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]);
+}
+
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]);
+}
+
_______________________________________________
cfe-commits mailing list
[email protected]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
