[Beignet] [PATCH V2 1/2] Refine error check in SetKernelArg() and support NULL buffer argument
Ruiling Song
ruiling.song at intel.com
Fri Jun 14 01:32:53 PDT 2013
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);
3. as there maybe small possibility that we get a 0 starting address for bo,
we add an assert() to make sure we do not get a 0 addressed bo.
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 ++++++++++++++++++++++++---------
src/intel/intel_gpgpu.c | 10 +++++-
4 files changed, 60 insertions(+), 19 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 895d6fc..b296dd7 100644
--- a/src/cl_command_queue.c
+++ b/src/cl_command_queue.c
@@ -137,7 +137,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);
@@ -153,7 +153,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;
diff --git a/src/intel/intel_gpgpu.c b/src/intel/intel_gpgpu.c
index 4a48d1d..19567dc 100644
--- a/src/intel/intel_gpgpu.c
+++ b/src/intel/intel_gpgpu.c
@@ -341,11 +341,19 @@ intel_gpgpu_batch_reset(intel_gpgpu_t *gpgpu, size_t sz)
{
intel_batchbuffer_reset(gpgpu->batch, sz);
}
-
+/* check we do not get a 0 starting address for binded buf */
+static void
+intel_gpgpu_check_binded_buf_address(intel_gpgpu_t *gpgpu)
+{
+ uint32_t i;
+ for (i = 0; i < gpgpu->binded_n; ++i)
+ assert(gpgpu->binded_buf[i]->offset != 0);
+}
static void
intel_gpgpu_flush(intel_gpgpu_t *gpgpu)
{
intel_batchbuffer_flush(gpgpu->batch);
+ intel_gpgpu_check_binded_buf_address(gpgpu);
}
static void
--
1.7.9.5
More information about the Beignet
mailing list