7 Replies Latest reply on Mar 3, 2013 5:26 AM by himanshu.gautam

    printf() misbehave?

    Raistmer

      Hello

       

      I want to check correctness of one of kernels so trying to print debug info from it:

       

       

      __kernel void GPU_fetch_array_kernel_cl(__global float* src, __global int* offsets, __global int* f_int, __global float4* dest,

                          const uint offset, const uint stride){

                uint threadID=get_global_id(0)+offset;

                uint j=get_global_id(1);

                int k=0;

                int l=0;

                float4 acc=(float4)0.f;

                int n_per=f_int[threadID];

                for(k=0;k<n_per;k++){

                                    l=offsets[k*stride+threadID];

                                    l+=(4*j);//R: index to data array computed

                                    acc.x+=src[l];

                                    acc.y+=src[l+1];

                                    acc.z+=src[l+2];

                                     acc.w+=src[l+3];

      #if 1

                printf("l=%d\n",l);

      #endif

                }

                dest[threadID*480+j]=acc;

      }

       

       

      to ensure that printf() allowed I used next line near beginning of CL file:

       

       

      #pragma OPENCL EXTENSION cl_amd_printf : enable

       

       

      but then no debug info from kernel printed and app crashes on this buffer read:

       

       

      err=clEnqueueReadBuffer(cq, gpu_results,CL_TRUE,0,sizeof(cl_uint),&CPU_result,0, NULL, NULL);OCL_LOG_ERR("clEnqueueReadBuffer->CPU_result");

       

       

      When I comment out extension enabling, that is:

       

       

      //#pragma OPENCL EXTENSION cl_amd_printf : enable

       

       

      debug info from kernel starts printing (!). And no compilation error on printf line.

      Also, no crash on later buffer reading.

      Why so? Why "legal" extension enabling leads to crash while illegal usage of printf from kernel works just OK ?

        • Re: printf() misbehave?
          himanshu.gautam

          1. Which driver version?

          2. Which platform? (Windows, linux)

          3. Distribution Name (win7, win8, ubuntu, suse etc..)

          4. Bitness (32 or 64-bit?)

          5. Are you running the code on CPU or GPU target?

          6. clinfo details of your platform

            • Re: printf() misbehave?
              Raistmer

              1.

                Driver version:


              CAL 1.4.1646 (VM)
                Version:



              OpenCL 1.2 AMD-APP (938.2)

               

              2-4: Windows7 x64 Home Premium SP1

              5: GPU target

              6: clinfo.txt attached. In short, cl_amd_printf listed withing extensions.

               

              BTW, one more observation. App prints its own "hello" string into stdout before any kernel call.

              But it appears AFTER dump from kernel in case where kernel dump works w/o crash. Not FIFO stdout ?

               

              .

                • Re: printf() misbehave?
                  himanshu.gautam

                  You seem to have a very old version of driver. Can you please upgrade to Catalyst 13.1?

                  (or)

                  I can check with Cat 13.1 within this week - provided you upload a ZIP file with a repro case on it.

                    • Re: printf() misbehave?
                      Raistmer

                      It was some joke?

                      In just prev my thread we discuss inability of 13.1 to compile my kernels at all. Do you really think I will install 13.1 after that ??? Moreover, just now I read your own experiment that confirmed that 13.1 fails with kernels compilation.

                      "Too old" driver at least WORKS. I can't say the same about more recent ones.

                      Printf extension was included long before this particular driver release. And it should work there just as in any another driver with extension support. Do you have some info  what driver after I use had some fix (listed in its release notes) for this extension? If not, why bother with upgrade even if 13.1 would be normal one ???

                      What I learnt in hard way dealing with AMD drivers that one should be very careful when do any driver change. Too big chances that working config will be lost forever (OS reinstall/unroll or smth like this will be required). Change working AMD driver is the last thing to try.

                        • Re: printf() misbehave?
                          himanshu.gautam

                          Hi Raistmer,

                          I do not see any problems in using cl_amd_printf extension with 13.1 driver. I can check once with your driver if you can give me its version (12.10 again?). As far as other bug that you reported, it is being looked into, and I will update the relevent thread when I get some update.

                          Anyways OpenCL 1.2 specification does not require you to enable that extension anymore for printf. Printf is available by default.

                          1 of 1 people found this helpful
                            • Re: printf() misbehave?
                              Raistmer

                              Actually, I don't know exactly what "Catalyst driver" it is. Described results were recived on C-60 based Acer Aspire One netbook with driver from manufacturer (Acer). AMD VISION Engine is quite useless tool for now, it provides no info about driver version at all. And GPU-Z shows display driver version 8.920.0.0 (but also it states that device name is Ontario while it's Loveland).

                              What I can say is only what clinfo reports:

                               

                                Driver version:


                              CAL 1.4.1646 (VM)
                                Version:



                              OpenCL 1.2 AMD-APP (938.2)

                               

                              I will try to do the same kernel modifications on another PC with Catalyst 12.8 installed and report how it will react.

                              Cause I already found workaround and can use printf from GPU for debugging it's not show-stopper issue, but still better to remove it than to leave it as it is.