[Beignet] [PATCH 1/2] Refine error check in SetKernelArg() and support NULL buffer argument
Zhigang Gong
zhigang.gong at linux.intel.com
Thu Jun 13 20:47:28 PDT 2013
This patch itself is good for me. As to the null pointer support on kernel side,
I have one concern here:
It seems that we don't do anything to prevent we allocate a valid 0 pointer to a
buffer. The relocate function call is as below:
for (i = 0; i < k->thread_n; ++i)
for (j = 0; j < gpgpu->binded_n; ++j) {
*(uint32_t*)(curbe + gpgpu->binded_offset[j]+i*k->cst_sz) = gpgpu->binded_buf[j]->offset;
drm_intel_bo_emit_reloc(gpgpu->curbe_b.bo,
gpgpu->binded_offset[j]+i*k->cst_sz,
gpgpu->binded_buf[j],
0,
I915_GEM_DOMAIN_RENDER,
I915_GEM_DOMAIN_RENDER);
}
I'm not very sure whether drm_intel_bo_emit_reloc will allocate a zero offset or not for a
given buffer object. If we want to treat NULL pointer as an invalid pointer on kernel
side, we must do something to prevent the above relocate function to relocate a buffer
to zero offset. Nanhai, is this concern valid from your point of view?
On Thu, Jun 13, 2013 at 02:00:25PM +0800, Ruiling Song wrote:
> 1. refine error check in clSetKernelArg() to follow spec.
> 2. add support NULL buffer as argument, so user could write like below:
> __kernel void func(__global int * p1, constant int* p2) {
> if(p1) {
> //do some thing if p1 is not NULL.
> } else {
> //do other things if p1 is NULL
> }
> }
> Then calling clSetKernelArg(k, 0, sizeof(cl_mem), NULL);
>
> Signed-off-by: Ruiling Song <ruiling.song at intel.com>
> ---
> backend/src/llvm/llvm_gen_backend.cpp | 5 ++-
> src/cl_command_queue.c | 4 +--
> src/cl_kernel.c | 60 ++++++++++++++++++++++++---------
> 3 files changed, 51 insertions(+), 18 deletions(-)
>
> diff --git a/backend/src/llvm/llvm_gen_backend.cpp b/backend/src/llvm/llvm_gen_backend.cpp
> index f579873..3a59da3 100644
> --- a/backend/src/llvm/llvm_gen_backend.cpp
> +++ b/backend/src/llvm/llvm_gen_backend.cpp
> @@ -693,7 +693,10 @@ namespace gbe
> return doIt(uint64_t(0));
> }
> }
> -
> + // NULL pointers
> + if(isa<ConstantPointerNull>(CPV)) {
> + return doIt(uint32_t(0));
> + }
> // Floats and doubles
> const Type::TypeID typeID = CPV->getType()->getTypeID();
> switch (typeID) {
> diff --git a/src/cl_command_queue.c b/src/cl_command_queue.c
> index a3987d8..5c7f7ae 100644
> --- a/src/cl_command_queue.c
> +++ b/src/cl_command_queue.c
> @@ -138,7 +138,7 @@ cl_command_queue_bind_surface(cl_command_queue queue, cl_kernel k)
> for (i = 0; i < k->arg_n; ++i) {
> uint32_t offset; // location of the address in the curbe
> arg_type = gbe_kernel_get_arg_type(k->opaque, i);
> - if (arg_type != GBE_ARG_GLOBAL_PTR)
> + if (arg_type != GBE_ARG_GLOBAL_PTR || !k->args[i].mem)
> continue;
> offset = gbe_kernel_get_curbe_offset(k->opaque, GBE_CURBE_KERNEL_ARGUMENT, i);
> cl_gpgpu_bind_buf(queue->gpgpu, k->args[i].mem->bo, offset, cc_llc_l3);
> @@ -154,7 +154,7 @@ LOCAL cl_int cl_command_queue_upload_constant_buffer(cl_kernel k,
> for(i = 0; i < k->arg_n; i++) {
> enum gbe_arg_type arg_type = gbe_kernel_get_arg_type(k->opaque, i);
>
> - if(arg_type == GBE_ARG_CONSTANT_PTR) {
> + if(arg_type == GBE_ARG_CONSTANT_PTR && k->args[i].mem) {
> uint32_t offset = gbe_kernel_get_curbe_offset(k->opaque, GBE_CURBE_EXTRA_ARGUMENT, i+GBE_CONSTANT_BUFFER);
> cl_mem mem = k->args[i].mem;
> cl_buffer_map(mem->bo, 1);
> diff --git a/src/cl_kernel.c b/src/cl_kernel.c
> index 851acfa..41e6a8a 100644
> --- a/src/cl_kernel.c
> +++ b/src/cl_kernel.c
> @@ -105,14 +105,42 @@ cl_kernel_set_arg(cl_kernel k, cl_uint index, size_t sz, const void *value)
> return CL_INVALID_ARG_INDEX;
> arg_type = gbe_kernel_get_arg_type(k->opaque, index);
> arg_sz = gbe_kernel_get_arg_size(k->opaque, index);
> +
> if (UNLIKELY(arg_type != GBE_ARG_LOCAL_PTR && arg_sz != sz))
> return CL_INVALID_ARG_SIZE;
>
> - /* Copy the structure or the value directly into the curbe */
> - if (arg_type == GBE_ARG_VALUE) {
> + if(UNLIKELY(arg_type == GBE_ARG_LOCAL_PTR && sz == 0))
> + return CL_INVALID_ARG_SIZE;
> + if(arg_type == GBE_ARG_VALUE) {
> + if(UNLIKELY(value == NULL))
> + return CL_INVALID_ARG_VALUE;
> + } else if(arg_type == GBE_ARG_LOCAL_PTR) {
> + if(UNLIKELY(value != NULL))
> + return CL_INVALID_ARG_VALUE;
> + } else if(arg_type == GBE_ARG_SAMPLER) {
> if (UNLIKELY(value == NULL))
> - return CL_INVALID_KERNEL_ARGS;
> + return CL_INVALID_ARG_VALUE;
> +
> + cl_sampler s = *(cl_sampler*)value;
> + if(s->magic != CL_MAGIC_SAMPLER_HEADER)
> + return CL_INVALID_SAMPLER;
> + } else {
> + // should be image, GLOBAL_PTR, CONSTANT_PTR
> + if (UNLIKELY(value == NULL && arg_type == GBE_ARG_IMAGE))
> + return CL_INVALID_ARG_VALUE;
> + if(value != NULL) {
> + mem = *(cl_mem*)value;
> + if (UNLIKELY(mem->magic != CL_MAGIC_MEM_HEADER))
> + return CL_INVALID_MEM_OBJECT;
> +
> + if (UNLIKELY((arg_type == GBE_ARG_IMAGE && !mem->is_image)
> + || (arg_type != GBE_ARG_IMAGE && mem->is_image)))
> + return CL_INVALID_ARG_VALUE;
> + }
> + }
>
> + /* Copy the structure or the value directly into the curbe */
> + if (arg_type == GBE_ARG_VALUE) {
> offset = gbe_kernel_get_curbe_offset(k->opaque, GBE_CURBE_KERNEL_ARGUMENT, index);
> assert(offset + sz <= k->curbe_sz);
> memcpy(k->curbe + offset, value, sz);
> @@ -124,8 +152,6 @@ cl_kernel_set_arg(cl_kernel k, cl_uint index, size_t sz, const void *value)
>
> /* For a local pointer just save the size */
> if (arg_type == GBE_ARG_LOCAL_PTR) {
> - if (UNLIKELY(value != NULL))
> - return CL_INVALID_KERNEL_ARGS;
> k->args[index].local_sz = sz;
> k->args[index].is_set = 1;
> k->args[index].mem = NULL;
> @@ -136,8 +162,6 @@ cl_kernel_set_arg(cl_kernel k, cl_uint index, size_t sz, const void *value)
> if (arg_type == GBE_ARG_SAMPLER) {
> cl_sampler sampler;
> memcpy(&sampler, value, sz);
> - if (UNLIKELY(sampler->magic != CL_MAGIC_SAMPLER_HEADER))
> - return CL_INVALID_KERNEL_ARGS;
> k->args[index].local_sz = 0;
> k->args[index].is_set = 1;
> k->args[index].mem = NULL;
> @@ -146,15 +170,21 @@ cl_kernel_set_arg(cl_kernel k, cl_uint index, size_t sz, const void *value)
> return CL_SUCCESS;
> }
>
> - /* Otherwise, we just need to check that this is a buffer */
> - if (UNLIKELY(value == NULL))
> - return CL_INVALID_KERNEL_ARGS;
> + if(value == NULL) {
> + /* for buffer object GLOBAL_PTR CONSTANT_PTR, it maybe NULL */
> + int32_t offset = gbe_kernel_get_curbe_offset(k->opaque, GBE_CURBE_KERNEL_ARGUMENT, index);
> + *((uint32_t *)(k->curbe + offset)) = 0;
> + assert(arg_type == GBE_ARG_GLOBAL_PTR || arg_type == GBE_ARG_CONSTANT_PTR);
> +
> + if (k->args[index].mem)
> + cl_mem_delete(k->args[index].mem);
> + k->args[index].mem = NULL;
> + k->args[index].is_set = 1;
> + k->args[index].local_sz = 0;
> + return CL_SUCCESS;
> + }
> +
> mem = *(cl_mem*) value;
> - if (UNLIKELY(mem->magic != CL_MAGIC_MEM_HEADER))
> - return CL_INVALID_ARG_VALUE;
> - if (UNLIKELY((arg_type == GBE_ARG_IMAGE && !mem->is_image)
> - || (arg_type != GBE_ARG_IMAGE && mem->is_image)))
> - return CL_INVALID_ARG_VALUE;
>
> if(arg_type == GBE_ARG_CONSTANT_PTR) {
> int32_t cbOffset;
> --
> 1.7.9.5
>
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet
More information about the Beignet
mailing list