[llvm] r194760 - Add addrspacecast instruction.

Michele Scandale michele.scandale at gmail.com
Mon Nov 18 13:18:50 PST 2013


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




More information about the llvm-commits mailing list