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

[Michele Scandale]

  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
+}

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
+}

_______________________________________________
cfe-commits mailing list
cfe-commits@cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits