3 Replies Latest reply on Sep 3, 2010 9:41 AM by davibu

    [BUG] Compiler problem with switch() statement and variable declaration

    davibu

      I have experience this problem with a large complex kernel and I try to exemplify the problem here with an example.

      The AMD OpenCL C compiler compiles the following code:

      __kernel void test(
              __global int *v,
              ) {
          int gid = get_global_id(0);

          switch(gid) {
              case 0:
                  int a = 0; // <=== Problem here, missing { ... }
                  gid = a + 1;
                  break;
          }

          v[gid] = gid;

      }

      by firing only a warning (i.e. "warning: a declaration cannot have a label"). First of all, the compiler should refuse to compile the above code because it has an error. The right code is:

      __kernel void test(
              __global int *v,
              ) {
          int gid = get_global_id(0);

          switch(gid) {
              case 0:

      { // <=== Addition
                  int a = 0;
                  gid = a + 1;

                  break;

      } // <=== Addition
          }

          v[gid] = gid;

      }

      Second, the compiler seems to generate some buggy code like if multiple variables were sharing the same register at the same time. This lead indeed to some very wired behaviour of the code.

      The problem was found on Linux Ubuntu 10.04 with SDK 2.2 and Catalyst 10.8.

      Both NVIDIA OpenCL and Apple OpenCL show the right behaviour: they refuse to compile the code with an error.

       

        • [BUG] Compiler problem with switch() statement and variable declaration
          himanshu.gautam

          hi davibu,

          i tried the problem you specified on vista X64.It indeed show only a warning in the first case and do compile.But i do not think declaring a label inside a case

          causes any accuracy issue.May be AMD doesn't consider the problem critical enough to be reported as an error.A warning is definitely noticeable.

          • [BUG] Compiler problem with switch() statement and variable declaration
            genaganna

             

            Originally posted by: davibu I have experience this problem with a large complex kernel and I try to exemplify the problem here with an example.

             

            The AMD OpenCL C compiler compiles the following code:

             

            __kernel void test(         __global int *v,         ) {     int gid = get_global_id(0);     switch(gid) {         case 0:             int a = 0; // <=== Problem here, missing { ... }             gid = a + 1;             break;     }

             

                v[gid] = gid;

             

            }

             

            by firing only a warning (i.e. "warning: a declaration cannot have a label"). First of all, the compiler should refuse to compile the above code because it has an error. The right code is:

             

            __kernel void test(         __global int *v,         ) {     int gid = get_global_id(0);     switch(gid) {         case 0:

             

            { // <=== Addition             int a = 0;             gid = a + 1;

             

                        break;

             

            } // <=== Addition     }

             

                v[gid] = gid;

             

            }

             

            Second, the compiler seems to generate some buggy code like if multiple variables were sharing the same register at the same time. This lead indeed to some very wired behaviour of the code.

             

            The problem was found on Linux Ubuntu 10.04 with SDK 2.2 and Catalyst 10.8.

             

            Both NVIDIA OpenCL and Apple OpenCL show the right behaviour: they refuse to compile the code with an error.

             

             



            Thanks for reporting first issue. Upcoming releases will give appropriate error for this.

            Could you please tell me where buggy code gererated(You second case)? Please give kernel code which shows this issue.