[PATCH] Support for address space casting in order to map these on the new 'addrspacecast' IR instruction

[Michele Scandale]

Kindly up.

-Michele

On 11/22/2013 11:07 AM, Michele Scandale wrote:
>    Changed OpenCL testcase, the previous version was illegal w.r.t. OpenCL 1.2 standard (casting pointer to another address space is illegal).
>    Fixed case of ArrayToPointerDecay and added CUDA test cases based those suggested by Jeroen Ketema (the user that notifed the problem to llvm-commits list).
>
>    This version of the patch just modify the CodeGen part as the previous.
>
> Hi rjmccall, rsmith,
>
> http://llvm-reviews.chandlerc.com/D2241
>
> CHANGE SINCE LAST DIFF
>    http://llvm-reviews.chandlerc.com/D2241?vs=5723&id=5726#toc
>
> Files:
>    lib/CodeGen/CGExprScalar.cpp
>    test/CodeGenCUDA/address-space-conversion.cu
>    test/CodeGenOpenCL/address-space-conversion.cl
>
> Index: lib/CodeGen/CGExprScalar.cpp
> ===================================================================
> --- lib/CodeGen/CGExprScalar.cpp
> +++ lib/CodeGen/CGExprScalar.cpp
> @@ -1299,6 +1299,8 @@
>     case CK_AnyPointerToBlockPointerCast:
>     case CK_BitCast: {
>       Value *Src = Visit(const_cast<Expr*>(E));
> +    if (E->getType()->isPointerType() && DestTy->isPointerType())
> +      return Builder.CreatePointerCast(Src, ConvertType(DestTy));
>       return Builder.CreateBitCast(Src, ConvertType(DestTy));
>     }
>     case CK_AtomicToNonAtomic:
> @@ -1360,7 +1362,7 @@
>
>       // Make sure the array decay ends up being the right type.  This matters if
>       // the array type was of an incomplete type.
> -    return CGF.Builder.CreateBitCast(V, ConvertType(CE->getType()));
> +    return CGF.Builder.CreatePointerCast(V, ConvertType(CE->getType()));
>     }
>     case CK_FunctionToPointerDecay:
>       return EmitLValue(E).getAddress();
> Index: test/CodeGenCUDA/address-space-conversion.cu
> ===================================================================
> --- test/CodeGenCUDA/address-space-conversion.cu
> +++ test/CodeGenCUDA/address-space-conversion.cu
> @@ -0,0 +1,24 @@
> +// RUN: %clang_cc1 %s -triple nvptx-- -fcuda-is-device -emit-llvm -o - | FileCheck %s
> +
> +#include "../SemaCUDA/cuda.h"
> +
> +#define N 32
> +
> +extern __shared__ int x;
> +
> +__global__ void explicit_address_space_cast(int* p) {
> +	// CHECK: explicit_address_space_cast
> +   __shared__ unsigned char x[N];
> +
> +   for (unsigned int i=0; i<(N/4); i++) {
> +     ((unsigned int *)x)[i] = 0;
> +		// CHECK: addrspacecast
> +   }
> +}
> +
> +__global__ void pointer_as_array_access() {
> +   __shared__ int A[10];
> +   int* p = A + 1;
> +   p[x] = 0;
> +	 // CHECK: addrspacecast
> +}
> Index: test/CodeGenOpenCL/address-space-conversion.cl
> ===================================================================
> --- test/CodeGenOpenCL/address-space-conversion.cl
> +++ test/CodeGenOpenCL/address-space-conversion.cl
> @@ -0,0 +1,14 @@
> +// RUN: %clang_cc1 %s -ffake-address-space-map -emit-llvm -o - | FileCheck %s
> +
> +#define NULL ((void*)0)
> +
> +void null_pointer_implicit_conversion(int i, __global int *A) {
> +	// CHECK: null_pointer_implicit_conversion
> +	__global int *b;
> +
> +	b = i > 42 ? A : NULL;
> +
> +	if (b != NULL)
> +	  A[0] = b[5];
> +	// CHECK: null
> +}
>