12 Replies Latest reply on Jun 4, 2014 12:50 AM by dipak

    Passing a NULL buffer to a kernel

    adm271828

      According to OpenCL 1.1 (§ 5.7.2) it is possible to pass a null buffer to a kernel: "If the argument is a buffer object, the arg_value pointer can be NULL or point to a NULL value in which case a NULL value will be used as the value for the argument declared as a pointer to __global or __constant memory in the kernel."

      However, I can't see it working. Here is what I get for various attempts (SDK 2.4).

      Regards,

         Antoine

       

      err = clSetKernelArg (k, 0, 0, 0); // -> err = CL_INVALID_ARG_VALUE cl_mem buf = 0; err = clSetKernelArg (k, 0, sizeof(cl_mem), &buf); // -> err = CL_INVALID_MEM_OBJECT err = clSetKernelArg (k, 0, 0, &buf); // -> err = CL_INVALID_ARG_SIZE

        • Passing a NULL buffer to a kernel
          himanshu.gautam

          Can you please post a test case showing the issue. How do you create the buffer object?

            • Passing a NULL buffer to a kernel
              adm271828

              Hi,

              the test case is the code already provided. It shall be clear I don't create a buffer, since I want to pass a NULL, and my understanding from the OpenCL spec (quoted in initial message) is that it is sufficient to pass a NULL arg_value, in which case the kernel will receive a null pointer.

              Do you suggest I shall create a buffer with zero size and the kernel would receive a null pointer? I didn't try, but if it is the case, it doesn't seem to conform to OpenCL, and it would waste a buffer ressource (in case clCreateBuffer succeeds with null size).

              Maybe I missunderstood the spec?

              Best regards,

                  Antoine

                • Passing a NULL buffer to a kernel
                  nou

                  IMHO clSetKernelArg(k, 0, sizeof(cl_mem), 0); is correct way to specify NULL pointer. (but return CL_INVALID_ARG_VALUE) OpenCL specification is little contradiction in this area.

                  CL_INVALID_ARG_VALUE if arg_value specified is NULL for an argument that is not declared with the __local qualifier or vice-versa.

                  so you can specify NULL pointer but you get error if do so

                    • Passing a NULL buffer to a kernel
                      tonyo_au

                      I do

                      size_t localMemBufferSize = xxx;

                      clSetKernelArg(k,0,localMemBufferSize,NULL);

                       

                      this works on my NVidia cards - I haven't tried it on my ATI yet.

                       

                      I have had problems with local memory like this

                       

                      clSetKernelArg(k,0,sizeof(cl_mem),p1);

                      clSetKernelArg(k,1,localMemBufferSize,NULL);

                      when p1 structure wasn't an acceptable multiple of 16 bytes. I haven't resolved this issue except by trial and error with filler fields

                        • Passing a NULL buffer to a kernel
                          adm271828

                          clSetKernelArg (k, 0, size, NULL), with size > 0 is for local memory, as explained int the spec. It works on my ATI card.

                          Once again, I want to pass a null pointer to global or constant memory (as described in the spec...). Seems to be broken (but not critical).

                            • Passing a NULL buffer to a kernel
                              himanshu.gautam

                              hi adm,

                              This appears to be a snack in the opencl spec. So it is better to report at khronos forums.

                              Meanwhile, i was interested if you were having some specific purpose for allocating a NULL buffer or were just experimenting.

                              Thanks

                                • Passing a NULL buffer to a kernel
                                  adm271828

                                  Hi Himanshu,

                                  I'm not just experimenting. I'd like to pass an optional output buffer to collect extra information from a kernel.

                                  On the kernel side, I'd have written: kernel k (__global int4* out) { ... if (out != 0) { ... write extra info here } ... }.

                                  Of course I can do it by having this optional buffer+code enabled/disabled conditionnaly at compile time. But since I will select each version at runtime in a non predictible pattern, it means I will have two kernels loaded at the same time, wasting resources (the extra code executed when out != 0 is about 10 % of the whole kernel).

                                  This is not critical however (having intrinsics that will help me not to generate 2 instructions when only a single one is necessary is much more critical... for instance)

                                  As for OpenCL spec, as an end-user, I can't say the spec is incorrect (if this is what you mean). Somebody wrote this sentence, that seems to make sense. If during approval process, it was decided to remove it and this was not done, then it's up to people that make products compliant to the spec to say it to the OpenCL committee. Especialy if they are part of the committee.

                                  Best regards,

                                  Antoine

                                    • Passing a NULL buffer to a kernel
                                      hduregger

                                      I agree, that part of the specification caught me aswell.

                                      Apart from creating a dummy 'NULL' buffer to bind when that kernel argument is not used, I also passed an int used as a boolean, to indicate whether the kernel should output data into the buffer or not. This way you don't have to create a second kernel.

                        • Re: Passing a NULL buffer to a kernel
                          ajk

                          Hello. This issue is still unresolved. clSetKernelArg fails when passing NULL as arg_value.

                          There was the same bug in Intel OpenCL: https://software.intel.com/en-us/forums/topic/281206 and it was resolved.

                            • Re: Passing a NULL buffer to a kernel
                              dipak

                              Hi,


                              It would be a great help if you can share a sample code capturing what you want to do and the error you're facing.We will try to reproduce the issue. Please also let us know your system details like SDK, Driver, OS (Window/Linux) (32/64), GPU etc.?


                              Thanks,

                                • Re: Passing a NULL buffer to a kernel
                                  ajk

                                  I've further studied this case and found that:

                                   

                                  suppose, we have the following kernel:

                                  __kernel void test(__global float *a, __global const float *b, __global const float *c) {

                                    int idx = get_global_id(0);

                                    a[idx] += b[idx] + (c ? c[idx] : 0);

                                  }

                                  then:

                                  clSetKernelArg(krn, 2, 0, NULL) fails on AMD, succeeds with correct results of the following kernel execution on Intel, succeeds on NVIDIA, but gives incorrect results (NVIDIA just ignores buffer assignment),

                                  while

                                  clSetKernelArg(krn, 2, sizeof(cl_mem), NULL) succeeds on all platforms with correct results, (cl_mem p_null = NULL, clSetKernelArg(krn, 2, sizeof(cl_mem), &p_null) also succeeds).

                                   

                                  So, I suppose behaviour of clSetKernelArg with arg_size=0 and arg_value=NULL is vendor specific (the documentation of clSetKernelArg doesn't specify should arg_size be checked or not when arg_value is NULL sadly).

                                    • Re: Passing a NULL buffer to a kernel
                                      dipak

                                      Hi,

                                      Yes, your assumption is correct. To specify a NULL buffer you should use: clSetKernelArg(krn, 2, sizeof(cl_mem), NULL).

                                      As per the clSetKernelArg API spec (clSetKernelArg) :

                                       

                                      ".....

                                      arg_size: 

                                           Specifies the size of the argument value. If the argument is a memory object, the size is the size of the buffer or image object type. ...

                                       

                                          Also, it generates error:

                                          CL_INVALID_ARG_SIZE if arg_size does not match the size of the data type for an argument that is not a memory object or if the argument is a memory object and arg_size != sizeof(cl_mem) or if arg_size is zero and the argument is declared with the __local qualifier or if the argument is a sampler and arg_size != sizeof(cl_sampler).

                                       

                                      ...."

                                       

                                          So according to API, you can assume that you have to specify the size of the argument type. There is no explicit mention about the arg_size in case of NULL. Some vendor may support 0 as arg_size value in case of NULL. But, for portability, better to use common one.

                                       

                                      Regards,