[llvm] r194760 - Add addrspacecast instruction.

[Michele Scandale]

On 11/18/2013 10:18 PM, Michele Scandale wrote:
> On 11/18/2013 07:29 PM, Jeroen Ketema wrote:
>>
>> 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).
>>
> 
> Thanks, it was the definition of NULL that I was using the prevent me to
> reproduce the problem. From a quick look, is a problem related to clang: now
> bitcast changed semantic and so clang must be updated in order to match the new
> semantic. A quick workaround seems to solve the problem, but I think that a
> better solution would be to introduce a new CastKind in clang in order to track
> address space conversion during semantic analysis, in order to drive nicely the
> code generation.
> Till tomorrow I wont able to work on this.
> 
> -Michele
> 

Here I prepared a possible fix: http://llvm-reviews.chandlerc.com/D2241

-Michele