[LLVMdev] PTX backend fatal error

[Alberto Magni]

Hi everybody,

I am testing the PTX backend using the OpenCL NVIDIA SDK benchmarks.
Compiling the Histogram64.cl program I get a several backend errors.

I isolated one of them in the following kernel program:

__kernel void kernel_function(__global int *input) {
    __local char localArray[16];
    for(unsigned int index = 0; index < 16; ++index)
      localArray[index] = 0;
    input[0] = localArray[get_local_id(0)];
}

fatal error: error in backend: Cannot select:
      0x5810cc0: i32,ch = load 0x57fa148,
      0x5810ac0, 0x58105c0<LD1[%arrayidx1], sext
      from i8> [ID=9]
  0x5810ac0: i32 = add 0x58109c0, 0x5813640 [ORD=113] [ID=8]
    0x58109c0: i32 = PTXISD::COPY_ADDRESS 0x5813540 [ID=7]
      0x5813540: i32 = TargetGlobalAddress<[16 x i8] addrspace(4)*
@kernel_function.localArray> 0 [ID=4]
    0x5813640: i32,ch = load 0x57fa148, 0x5810dc0,
0x58105c0<LD4[%retval.i]> [ORD=110] [ID=5]
      0x5810dc0: i32 = FrameIndex<0> [ORD=110] [ID=1]
      0x58105c0: i32 = undef [ORD=110] [ID=2]
  0x58105c0: i32 = undef [ORD=110] [ID=2]

The command I am using is:

clang kernels/fatal_error_test.cl -O0 -include ocldef.h -include
builtin_functions_ptx.cl
                                                   -D__x86_64__
-ccc-host-triple ptx32 -Xclang
                                                   -target-feature
-Xclang +ptx23 -Xclang
                                                   -target-feature
-Xclang +compute20

Any ideas ?

Best regards

Alberto