[PATCH] Fix for CUDA codegen due to missing generation of implicit address space conversions

[Michele Scandale]

Hi jholewinski, rjmccall,

Here a small patch that fixes the CUDA code generator handling implicit address space conversions. In CUDA address spaces at sema level are attributes for the variable declaration, making my previous patch on casts not effective. In such cases we must generate some addrspacecast.

The specific test cases were provided by Jeroen Ketama.

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

Files:
  lib/CodeGen/CGCall.cpp
  lib/CodeGen/CGExpr.cpp
  lib/CodeGen/CodeGenFunction.cpp
  lib/CodeGen/CodeGenFunction.h
  test/CodeGenCUDA/address-spaces-implicit-conversion.cu

Index: lib/CodeGen/CGCall.cpp
===================================================================
--- lib/CodeGen/CGCall.cpp
+++ lib/CodeGen/CGCall.cpp
@@ -2430,7 +2430,7 @@
         // can happen due to trivial type mismatches.
         if (IRArgNo < IRFuncTy->getNumParams() &&
             V->getType() != IRFuncTy->getParamType(IRArgNo))
-          V = Builder.CreateBitCast(V, IRFuncTy->getParamType(IRArgNo));
+          V = EmitImplicitTypeCoercion(V, IRFuncTy->getParamType(IRArgNo));
         Args.push_back(V);
         
         checkArgMatches(V, IRArgNo, IRFuncTy);
Index: lib/CodeGen/CGExpr.cpp
===================================================================
--- lib/CodeGen/CGExpr.cpp
+++ lib/CodeGen/CGExpr.cpp
@@ -1432,6 +1432,13 @@
   }
 
   assert(Src.isScalar() && "Can't emit an agg store with this method");
+
+  // Handle CUDA implicit address space conversions.
+  if (getLangOpts().CUDA)
+    if (Src.getScalarVal()->getType()->isPtrOrPtrVectorTy())
+      Src = RValue::get(EmitBitCastOrAddrSpaceCast(Src.getScalarVal(),
+                                                   ConvertType(Dst.getType())));
+
   EmitStoreOfScalar(Src.getScalarVal(), Dst, isInit);
 }
 
Index: lib/CodeGen/CodeGenFunction.cpp
===================================================================
--- lib/CodeGen/CodeGenFunction.cpp
+++ lib/CodeGen/CodeGenFunction.cpp
@@ -1499,3 +1499,21 @@
 }
 
 CodeGenFunction::CGCapturedStmtInfo::~CGCapturedStmtInfo() { }
+
+llvm::Value *CodeGenFunction::EmitBitCastOrAddrSpaceCast(llvm::Value *V,
+                                                         llvm::Type *Ty) {
+  if (V->getType()->isPtrOrPtrVectorTy() && Ty->isPtrOrPtrVectorTy() &&
+      V->getType()->getPointerAddressSpace() != Ty->getPointerAddressSpace())
+    return Builder.CreateAddrSpaceCast(V, Ty);
+
+  return Builder.CreateBitCast(V, Ty);
+}
+
+llvm::Value *CodeGenFunction::EmitImplicitTypeCoercion(llvm::Value *V,
+                                                       llvm::Type *Ty) {
+  // Handle implicit pointer address space conversion for CUDA.
+  if (getLangOpts().CUDA)
+    return EmitBitCastOrAddrSpaceCast(V, Ty);
+
+  return Builder.CreateBitCast(V, Ty);
+}
Index: lib/CodeGen/CodeGenFunction.h
===================================================================
--- lib/CodeGen/CodeGenFunction.h
+++ lib/CodeGen/CodeGenFunction.h
@@ -2369,6 +2369,9 @@
 
   RValue EmitAtomicExpr(AtomicExpr *E, llvm::Value *Dest = 0);
 
+  llvm::Value *EmitImplicitTypeCoercion(llvm::Value *V, llvm::Type *Ty);
+  llvm::Value *EmitBitCastOrAddrSpaceCast(llvm::Value *V, llvm::Type *Ty);
+
   //===--------------------------------------------------------------------===//
   //                         Annotations Emission
   //===--------------------------------------------------------------------===//
Index: test/CodeGenCUDA/address-spaces-implicit-conversion.cu
===================================================================
--- test/CodeGenCUDA/address-spaces-implicit-conversion.cu
+++ test/CodeGenCUDA/address-spaces-implicit-conversion.cu
@@ -0,0 +1,41 @@
+// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple nvptx-unknown-unknown | FileCheck %s
+
+#include "../SemaCUDA/cuda.h"
+
+
+__device__ void bar(int &x) {
+  // CHECK: bar
+  x = 0;
+}
+
+int x;
+
+__global__ void implicit_as_cast_from_shared_scalar() {
+  // CHECK: implicit_as_cast_from_shared_scalar
+  __shared__ int b;
+  int &y = b;
+  // CHECK: addrspacecast
+  bar(y);
+}
+
+__global__ void implicit_as_cast_from_shared_array() {
+  // CHECK: implicit_as_cast_from_shared_array
+  __shared__ int b[1024];
+  // CHECK: addrspacecast
+  bar(b[x]);
+}
+
+__global__ void implicit_as_cast_from_device_scalar() {
+  // CHECK: implicit_as_cast_from_device_scalar
+  __device__ int b;
+  int &y = b;
+  bar(y);
+}
+
+__global__ void implicit_as_cast_from_device_array() {
+  // CHECK: implicit_as_cast_from_device_array
+  __device__ int b[1024];
+  // CHECK: addrspacecast
+  bar(b[x]);
+}
+
-------------- next part --------------
A non-text attachment was scrubbed...
Name: D2462.1.patch
Type: text/x-patch
Size: 3998 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20131223/b62988ba/attachment.bin>