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

[Michele Scandale]

  This new version of the patch add another missing case (thanks to Jeroen Ketema).

  About the doubt pointed out by David, I think that the correctness should be checked at sematic level: if there is the semantic correctness then the codegen change should be correct too.

  Thanks in advance for your contributions.

Hi pekka.jaaskelainen, rjmccall, rsmith,

http://llvm-reviews.chandlerc.com/D2241

CHANGE SINCE LAST DIFF
  http://llvm-reviews.chandlerc.com/D2241?vs=5726&id=5784#toc

Files:
  lib/CodeGen/CGExprConstant.cpp
  lib/CodeGen/CGExprScalar.cpp
  test/CodeGenCUDA/address-space-conversion.cu
  test/CodeGenOpenCL/address-space-conversion.cl

Index: lib/CodeGen/CGExprConstant.cpp
===================================================================
--- lib/CodeGen/CGExprConstant.cpp
+++ lib/CodeGen/CGExprConstant.cpp
@@ -1068,7 +1068,7 @@
       // Convert to the appropriate type; this could be an lvalue for
       // an integer.
       if (isa<llvm::PointerType>(DestTy))
-        return llvm::ConstantExpr::getBitCast(C, DestTy);
+        return llvm::ConstantExpr::getPointerCast(C, DestTy);
 
       return llvm::ConstantExpr::getPtrToInt(C, DestTy);
     } else {
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,40 @@
+// 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;
+
+// CHECK: @_ZZ6callervE1p = internal addrspace(3) global i32* addrspacecast
+
+__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
+}
+
+__device__ int* callee(int* p) {
+	// CHECK: callee
+  return p;
+}
+
+__global__ void caller() {
+	// CHECK: @_Z6callerv
+  __shared__ int A[10];
+  __shared__ int* p = A;
+  int *np = callee(p);
+	A[2] = 5;
+	np[0] = 2;
+}
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
+}
-------------- next part --------------
A non-text attachment was scrubbed...
Name: D2241.4.patch
Type: text/x-patch
Size: 3092 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20131126/260fc167/attachment.bin>