20 Replies Latest reply on Feb 25, 2013 12:17 AM by himanshu.gautam

    OpenCL compiler crashes, example with 2 line kernel

    jaap_nhl_lwd

      I have problems compiling my some of my kernels for AMD GPUs, no problems with NVDIA OpenCL.

       

      Used platform OpenCL 1.2 AMD-APP (1084.4) Juniper HD5700 Win7-64.

       

      Minimal kernel extracted from bigger one as example:

       

      kernel void AMDBug (const uint xx, global uint *gg) {
         uint x = get_global_id(0) + xx;
         if (x < xx) *gg = 1;
      }

       

      Compiler crashed when compiling for GPU, no problem for CPU.

      Last message before crash:

      LLVM ERROR: Cannot select: 0xb1e670: i8 = setcc <lots of hex codes>

       

      Please your advice.

      Best regards, Jaap.

        • Re: OpenCL compiler crashes, example with 2 line kernel
          himanshu.gautam

          Will try this next week.

          Meanwhile - if it is not a hassle - can you post your little sample over here as a zip file or so.

          • Re: OpenCL compiler crashes, example with 2 line kernel
            nou

            copy paste this code into amd kernel analyzer cause this error.

            LLVM ERROR: Cannot select: 0x7f0c601926c0: i8 = setcc 0x7f0c60248ae0, 0x7f0c601927c0, 0x7f0c601924c0 [ID=12]
              0x7f0c60248ae0: i32 = AMDILISD::ADD 0x7f0c601927c0, 0x7f0c602485e0 [ID=11]
                0x7f0c601927c0: i32 = AMDILISD::VEXTRACT 0x7f0c602489e0, 0x7f0c602486e0 [ORD=2] [ID=10]
                  0x7f0c602489e0: v4i32,ch = llvm.AMDIL.get.global.id 0x7f0c60242160, 0x7f0c602488e0 [ORD=1] [ID=9]
                    0x7f0c602488e0: i32 = TargetConstant<2561> [ORD=1] [ID=2]
                  0x7f0c602486e0: i32 = TargetConstant<1> [ORD=2] [ID=7]
                0x7f0c602485e0: i32,ch = CopyFromReg 0x7f0c60242160, 0x7f0c602484e0 [ORD=3] [ID=8]
                  0x7f0c602484e0: i32 = Register %vreg0 [ORD=3] [ID=1]
              0x7f0c601927c0: i32 = AMDILISD::VEXTRACT 0x7f0c602489e0, 0x7f0c602486e0 [ORD=2] [ID=10]
                0x7f0c602489e0: v4i32,ch = llvm.AMDIL.get.global.id 0x7f0c60242160, 0x7f0c602488e0 [ORD=1] [ID=9]
                  0x7f0c602488e0: i32 = TargetConstant<2561> [ORD=1] [ID=2]
                0x7f0c602486e0: i32 = TargetConstant<1> [ORD=2] [ID=7]
            
            • Re: OpenCL compiler crashes, example with 2 line kernel
              drallan

              I have problems compiling my some of my kernels for AMD GPUs, no problems with NVDIA OpenCL.

              Used platform OpenCL 1.2 AMD-APP (1084.4) Juniper HD5700 Win7-64.

              Minimal kernel extracted from bigger one as example:

               

              kernel void AMDBug (const uint xx, global uint *gg) {
                 uint x = get_global_id(0) + xx;
                 if (x < xx) *gg = 1;
              }

               

              Compiler crashed when compiling for GPU, no problem for CPU.

              Last message before crash:

              LLVM ERROR: Cannot select: 0xb1e670: i8 = setcc <lots of hex codes>

               

              This is very familiar. I reported almost the same problem here in November starting about Catalyst 12.11.

              http://devgurus.amd.com/message/1285016#1285016

              In that case, it was a compiler error triggered by the expression  c = hi + (x < c),  which dumped the same kind of LLVM error message. The compiler erred on the conditional x < c, which is also in the current code. Most other conditionals like x <= c. worked fine.

               

              I would love to know what happens if you use  if (x <= xx) *gg = 1 instead.

              Or try if(x <= xx-1) to be logically the same.