[llvm] r194760 - Add addrspacecast instruction.

[Jeroen Ketema]

Hi Michele,

I'm currently on 195003.

I'm compiling the OpenCL code with:

lvm-build/Debug+Asserts/bin/clang -target nvptx--nvidiacl -emit-llvm -c  kernel.cl

(and similarly for the CUDA code)

Slightly simpler CUDA example:

__global__ void foo() {
  __shared__ int A[10];
  int* p = A;
  p[0] = 0;
}

Slightly simpler OpenCL example:

#define NULL ((void*)0)
__kernel void foo(int i, __global int *A)
{
  __global int *a;

  if (i == 0)
    a = A;
  else
    a = NULL;

  if (a != NULL)
    A[0] = 0;
}

Note the definition of NULL is as in libclc (llvm related project).

Best,

 Jeroen

On Nov 18, 2013, at 5:32 PM, Michele Scandale <michele.scandale at gmail.com> wrote:

> On 11/18/2013 03:13 PM, Jeroen Ketema wrote:
>> 
>> This commit seems to break compilation of certain CUDA and OpenCL kernels.
>> 
>> For example:
>> 
>> CUDA:
>> 
>> __global__ void foo() {
>>   __shared__ int A[10];
>>   int* p = A;
>>   p[threadIdx.x] = 0;
>> }
>> 
>> If clang/llvm is complied with debugging enabled I get:
>> 
>> Assertion failed: (CastInst::castIsValid(opc, C, Ty) && "Invalid constantexpr cast!"), function getCast, file llvm/lib/IR/Constants.cpp, line 1448.
>> 
>> OpenCL:
>> 
>> __kernel void foo(int i, __global int *A)
>> {
>>   __global int *a;
>> 
>>   if (i == 0)
>>     a = A;
>>   else
>>     a = NULL;
>> 
>>   if (a != NULL)
>>     A[get_global_id(0)] = get_global_id(0);
>> }
>> 
>> Again I get:
>> 
>> Assertion failed: (CastInst::castIsValid(opc, C, Ty) && "Invalid constantexpr cast!"), function getCast, file llvm/lib/IR/Constants.cpp, line 1448.
>> 
> 
> Hi Jeroen,
> 
> I'm not able to reproduce the failure, could you provide me the revision commit of clang and llvm you are using?
> 
> -Michele
> 
> 
>> Regards,
>> 
>>  Jeroen
>>