16 Replies Latest reply on Jun 17, 2014 4:35 AM by pinform

    Strange behavior in kernel - Losing variables during kernel run

    dasmurkel

      Hi,

       

      I've come across some very strange kernel behavior when running a certain kernel on my Radeon Device. Interestingly, the exact same code performs fine using the CPU-Device, either with the AMD-APP SDK or the Intel SDK (I have a Core i5).

       

      The kernel looks as follows:

       

      kernel void
      incident_bitmask (const uint size,
                        const uint intersectionBufferPerRay,
                        const global uint* intersectionCount,
                        const global intersection* intersections,
                        const global ray* rays,
                        const global transmitter* sourceTransmitters,
                        const global wall* walls,
                        const global building* buildings,
                        float rx_planeHeight,
                        global uint* bitmask,
                        global incidentPoint* incidents)
      {
        const uint id = get_global_id (0);
        if (id >= size)
          {
            return;
          }

       

      // ... Some stuff ...
        printf ("ID: %i, intersectionBufPerRay: %i\n", id, intersectionBufferPerRay);
        const uint resultBufferSize = 2 * intersectionBufferPerRay + 1;

       

      // ... Lots of more stuff, branches, loopsand function calls ...


        printf ("resultBufferSize for ID %i: %i\n", id, resultBufferSize);

      }

       

      I run this kernel with a global size of 6 and obviously the "size" parameter is also set to 6. This is my output:

       

      ID: 0, intersectionBufPerRay: 100
      resultBufferSize for ID 0: 201
      ID: 1, intersectionBufPerRay: 100
      resultBufferSize for ID 1: 201
      ID: 2, intersectionBufPerRay: 100
      resultBufferSize for ID 2: 201
      ID: 3, intersectionBufPerRay: 100
      resultBufferSize for ID 0: 201
      ID: 4, intersectionBufPerRay: 100
      resultBufferSize for ID 0: 201
      ID: 5, intersectionBufPerRay: 100
      resultBufferSize for ID 0: 201
      ID: 6, intersectionBufPerRay: 100
      resultBufferSize for ID 0: 201

       

       

      Please note that I get the IDs printed out correctly for the first prinf but that I get only 0-2 for the second printf. The rest of the ids (which were set to the correct global_id at the beginning of the kernel and are const, so there is no way I changed them) is set to zero.

       

      Another thing is, that this only happens sometimes and hat the margin for which the values are correct changes. Sometimes everything is fine, sometimes I get 0-3 correct and the rest is wrong, sometimes the "intersectionBufferPerRay" also prints zero.

       

      Environment: Ubuntu 12.10, Catalyst 13.4, Radeon HD 7850

       

      Any help is greatly appreciated!

        • Re: Strange behavior in kernel - Losing variables during kernel run
          himanshu.gautam

          Hi,

           

          Just with the above its difficult to say anything..

          There are some points which are actualy need to check

          1. you are saying the global size which you set is 6. then actually you should get the id range from 0 to 5 not 6... So i have doubt that your host program is having some problem.

          2. You are mentioning that in the second run the results are not coming properly which means that are you executing the program for the second time or you are just calling the kernel in some loop.

          So to understand all these please post the sample code.

            • Re: Strange behavior in kernel - Losing variables during kernel run
              dasmurkel

              Thank you for your answer.

               

              The kernel together with is helper functions is very long and since I couldn't find any environment for code in the forum editor, I cropped it.

               

              As to your points:

              1. Global range and "size" are, of course, the same. But it should not matter anyway.

              2. I meant reruns of the whole program.

               

              I did, however, find the problem and you are correct, it couldn't have been found with my short code snippet. At some point, I declared some functions inline. That messed up the kernel. Once I removed the "inline" keyword, the results were correct (again).

               

              Then I looked for "inline" in the standard and didn't find it mentioned anywhere, neither as allowed, nor as restricted from the C99 subset. So I guess I feel like there should be a compiler warning if "inline" is used but not supported. If it is supported - did I just find my first compiler bug?

               

              If requested, I can still post the entire code.

                • Re: Strange behavior in kernel - Losing variables during kernel run
                  himanshu.gautam

                  Hi,

                   

                  could you please post your code. Because i have written one sample code with inline function and its working fine for me.

                  Check the below code... inline function F1 just receiving argument and returning the same...

                  inline int F1(unsigned int val)
                  {
                      return val;
                  }

                  __kernel void templateKernel(__global  unsigned int * output,
                                               __global  unsigned int * input,
                                               const     unsigned int multiplier)

                      uint tid = get_global_id(0);
                      if(tid==0)printf(" tid:%d", tid);
                      output[tid] = input[tid] * multiplier;
                      output[tid] =   F1(output[tid]);

                  }