3 Replies Latest reply on Dec 5, 2014 1:50 AM by dipak

    clBuildProgram() Dereference Bug -- Best place to submit a Bug report?

    alaric

      Hello,

       

      I've been trying to build an OpenCL application with pyopencl, but lately I've been hitting a segfault when I call program.build(), which is a wrapper for clBuildProgram().  Oddly, the problem only happens when I run my CL code on my 270X GPU.  When I create a context on my 1090T CPU, the OpenCL program runs fine.  Additionally, other OpenCL programs were running just fine on my GPU.  This led me to believe that the segfault must be code dependent, and I proceeded to comment out my kernels one at a time to discover which kernels were causing problems.  The only three kernels that failed were each referencing and dereferencing pointers to an array (for style and clarity).  Here's an example:

       

      //Expected Grid Size: num_inputs x num_outputs x 1

      __kernel void nn_update_mat_trans_finish(

          __global nn_desc* net,

          __global nn_desc* delta_net,

          __global data_t* mat_trans_gradient,

          __global data_t* mat_trans

      ) {

          uint it_in = get_global_id(0);

          uint it_out = get_global_id(1);

       

          data_t* w = &mat_trans[net->num_inputs * it_out + it_in];

          data_t* dw = &mat_trans_gradient[delta_net->num_inputs * it_out + it_in];

       

          *w += *dw;

      }

       

      Unfortunately, it appears that the dereference operator is broken on somewhere in my GPU's OpenCL stack.  The code also fails to build when I try something along the lines of

      data_t* w = mat_trans + net->num_inputs * it_out + it_in;

       

      When I comment out the "*w += *dw;" line, then the code builds fine on my GPU. 

       

      My question is, where would be the proper place for an OpenCL newbie like myself to submit a bug report about this issue?  I'm not sure if the fault lies in FGLRX (A proprietary linux graphics driver), or somewhere else in the OpenCL stack.

       

      Thanks

        • Re: clBuildProgram() Dereference Bug -- Best place to submit a Bug report?
          alaric

          P.S., data_t is a typedef for float.  Additionally, for those who find this via Google, this, oddly, does compile:

           

          //Expected Grid Size: num_inputs x num_outputs x 1

          __kernel void nn_update_mat_trans_finish(

              __global nn_desc* net,

              __global data_t* mat_trans_gradient,

              __global data_t* mat_trans,

              __global nn_desc* delta_net

          ) {

              uint it_in = get_global_id(0);

              uint it_out = get_global_id(1);

           

              *(mat_trans + net->num_inputs * it_out + it_in) += mat_trans_gradient[delta_net->num_inputs * it_out + it_in];

          }

           

          While this does not compile:

          //Expected Grid Size: num_inputs x num_outputs x 1

          __kernel void nn_update_mat_trans_finish(

              __global nn_desc* net,

              __global float* mat_trans_gradient,

              __global float* mat_trans,

              __global nn_desc* delta_net

          ) {

              uint it_in = get_global_id(0);

              uint it_out = get_global_id(1);

           

              float* w = (mat_trans + net->num_inputs * it_out + it_in);

           

              *w += mat_trans_gradient[delta_net->num_inputs * it_out + it_in];

          }

            • Re: clBuildProgram() Dereference Bug -- Best place to submit a Bug report?
              andylewis

              That's really nice. I like to use it.

              • Re: clBuildProgram() Dereference Bug -- Best place to submit a Bug report?
                dipak

                Hi,

                My question is, where would be the proper place for an OpenCL newbie like myself to submit a bug report about this issue?

                 

                Done!...You've already reported a bug/issue... In this forum, you can share your OpenCL related problems/issues as well as can post a bug report (for attachment use advance editor mode).  Whenever report a issue/bug, please mention your set-up details such as OS, GPU, CPU, driver etc. In some cases, we suggest the users to provide a sample test-case that manifests the particular issue/problem, so that we can reproduce it and, if required, forward it to concerned team.

                 

                The only three kernels that failed were each referencing and dereferencing pointers to an array (for style and clarity).  Here's an example:

                 

                 

                The above kernel code seems incomplete [say declaration of "nn_desc" missing]. Can you please post the full kernel code?

                Did you try to build/compile the kernel using CodeXL? If not, please try and share your observation. Please mention your driver version and build options.

                 

                BTW, one suggestion:

                Modify the following line as below:

                float* w = (mat_trans + net->num_inputs * it_out + it_in);


                BY:


                __global float* w = (mat_trans + net->num_inputs * it_out + it_in);


                 

                Regards