add info

2 files changed, 17 insertions, 0 deletions

diff --git a/include/cl_kernel.h b/include/cl_kernel.h
index a75eb171..e61b6d98 100644
--- a/include/cl_kernel.h
+++ b/include/cl_kernel.h
@@ -78,6 +78,7 @@ typedef struct _cl_kernel {
   uint32_t arg_num;                 /* Number of arguments */
   cl_kernel_workgroup_info wg_info; /* WorkGroup info for each device. */
   pthread_mutex_t lock;             /* The lock to protect this struct. */
+  cl_bool in_enqueue;               /* Kernel are enqueued, we do not want it to be changed. */
   void* pdata;
 } _cl_kernel;
 
diff --git a/libclapi/cl_kernel.c b/libclapi/cl_kernel.c
index ecccb1f6..c8b6e191 100644
--- a/libclapi/cl_kernel.c
+++ b/libclapi/cl_kernel.c
@@ -210,6 +210,7 @@ LOCAL void cl_release_kernel(cl_kernel k)
   if (cl_ref_dec(&k->ref_n) > 1)
     return;
 
+  assert(k->in_enqueue == CL_FALSE);
   build_status = CL_CALLOC(k->program->ctx->device_num, sizeof(int32_t));
   assert(build_status != NULL);
   CL_MUTEX_LOCK(&k->program->lock);
@@ -923,6 +924,14 @@ clSetKernelArg(cl_kernel     kernel,
   cl_int err = CL_SUCCESS;
   CHECK_KERNEL(kernel);
 
+  CL_MUTEX_LOCK(&kernel->lock);
+  if (kernel->in_enqueue == CL_TRUE) {
+    CL_MUTEX_UNLOCK(&kernel->lock);
+    err = CL_OUT_OF_RESOURCES;
+    goto error;
+  }
+  CL_MUTEX_UNLOCK(&kernel->lock);
+
   err = cl_kernel_set_arg(kernel, arg_index, arg_size, arg_value);
 error:
   return err;
@@ -1044,9 +1053,16 @@ clEnqueueNDRangeKernel(cl_command_queue  command_queue,
     }
   }
 
+  CL_MUTEX_LOCK(&kernel->lock);
+  kernel->in_enqueue = CL_TRUE;
+  CL_MUTEX_UNLOCK(&kernel->lock);
   /* Do device specific checks are enqueue the kernel */
   err = cl_enqueue_ND_range(command_queue, kernel, work_dim, fixed_global_off, fixed_global_sz,
                             fixed_local_sz, num_events_in_wait_list, event_wait_list, event);
+  CL_MUTEX_LOCK(&kernel->lock);
+  kernel->in_enqueue = CL_FALSE;
+  CL_MUTEX_UNLOCK(&kernel->lock);
+
   if(err != CL_SUCCESS)
     goto error;