[Beignet] [PATCH 1/2] Refine error check in SetKernelArg() and support NULL buffer argument
Zou, Nanhai
nanhai.zou at intel.com
Thu Jun 13 21:29:28 PDT 2013
Usually offset 0 was pin mapped by some pages used by kernel GPU driver at driver initial time.
So it is very unlikely the offset will be 0, you can put an assert here to make sure this will never happen.
Thanks
Zou Nanhai
-----Original Message-----
From: beignet-bounces+nanhai.zou=intel.com at lists.freedesktop.org [mailto:beignet-bounces+nanhai.zou=intel.com at lists.freedesktop.org] On Behalf Of Zhigang Gong
Sent: Friday, June 14, 2013 11:47 AM
To: Song, Ruiling
Cc: beignet at lists.freedesktop.org
Subject: Re: [Beignet] [PATCH 1/2] Refine error check in SetKernelArg() and support NULL buffer argument
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
_______________________________________________
Beignet mailing list
Beignet at lists.freedesktop.org
http://lists.freedesktop.org/mailman/listinfo/beignet
More information about the Beignet
mailing list