Version that explicitly does not support C++: address spaces within C++
should be better discussed and it seems that currently there is no interest.
Originally this patch wanted to fix the case of OpenCL and CUDA. The address
space extension on C is fixed as well. *IMHO for now this patch should be fine
to address the original problem.*
Trying to play with address spaces in C++ now will bring to codegen assertion
due the illegality of bitcast between pointers of different address spaces:
solve this problem requires a definition of what is legal or not, and a
semantic extension, for C++ w.r.t. different address spaces.
Hi pekka.jaaskelainen, rjmccall, rsmith,
http://llvm-reviews.chandlerc.com/D2241
CHANGE SINCE LAST DIFF
http://llvm-reviews.chandlerc.com/D2241?vs=5830&id=5973#toc
Files:
include/clang/AST/OperationKinds.h
lib/AST/Expr.cpp
lib/AST/ExprConstant.cpp
lib/CodeGen/CGExpr.cpp
lib/CodeGen/CGExprAgg.cpp
lib/CodeGen/CGExprComplex.cpp
lib/CodeGen/CGExprConstant.cpp
lib/CodeGen/CGExprScalar.cpp
lib/Edit/RewriteObjCFoundationAPI.cpp
lib/Sema/SemaExpr.cpp
lib/StaticAnalyzer/Core/ExprEngineC.cpp
test/CodeGenCUDA/address-space-conversion.cu
test/CodeGenOpenCL/address-space-conversion.cl
Index: include/clang/AST/OperationKinds.h
===================================================================
--- include/clang/AST/OperationKinds.h
+++ include/clang/AST/OperationKinds.h
@@ -295,7 +295,10 @@
CK_BuiltinFnToFnPtr,
// Convert a zero value for OpenCL event_t initialization.
- CK_ZeroToOCLEvent
+ CK_ZeroToOCLEvent,
+
+ // Convert a pointer to a different address space.
+ CK_AddressSpaceConversion
};
static const CastKind CK_Invalid = static_cast<CastKind>(-1);
Index: lib/AST/Expr.cpp
===================================================================
--- lib/AST/Expr.cpp
+++ lib/AST/Expr.cpp
@@ -1474,6 +1474,11 @@
assert(getSubExpr()->getType()->isFunctionType());
goto CheckNoBasePath;
+ case CK_AddressSpaceConversion:
+ assert(getType()->isPointerType());
+ assert(getSubExpr()->getType()->isPointerType());
+ assert(getType()->getPointeeType().getAddressSpace() !=
+ getSubExpr()->getType()->getPointeeType().getAddressSpace());
// These should not have an inheritance path.
case CK_Dynamic:
case CK_ToUnion:
@@ -1636,6 +1641,8 @@
return "BuiltinFnToFnPtr";
case CK_ZeroToOCLEvent:
return "ZeroToOCLEvent";
+ case CK_AddressSpaceConversion:
+ return "AddressSpaceConversion";
}
llvm_unreachable("Unhandled cast kind!");
Index: lib/AST/ExprConstant.cpp
===================================================================
--- lib/AST/ExprConstant.cpp
+++ lib/AST/ExprConstant.cpp
@@ -7120,6 +7120,7 @@
case CK_BuiltinFnToFnPtr:
case CK_ZeroToOCLEvent:
case CK_NonAtomicToAtomic:
+ case CK_AddressSpaceConversion:
llvm_unreachable("invalid cast kind for integral value");
case CK_BitCast:
@@ -7592,6 +7593,7 @@
case CK_BuiltinFnToFnPtr:
case CK_ZeroToOCLEvent:
case CK_NonAtomicToAtomic:
+ case CK_AddressSpaceConversion:
llvm_unreachable("invalid cast kind for complex value");
case CK_LValueToRValue:
Index: lib/CodeGen/CGExpr.cpp
===================================================================
--- lib/CodeGen/CGExpr.cpp
+++ lib/CodeGen/CGExpr.cpp
@@ -2744,6 +2744,7 @@
case CK_ARCReclaimReturnedObject:
case CK_ARCExtendBlockObject:
case CK_CopyAndAutoreleaseBlockObject:
+ case CK_AddressSpaceConversion:
return EmitUnsupportedLValue(E, "unexpected cast lvalue");
case CK_Dependent:
Index: lib/CodeGen/CGExprAgg.cpp
===================================================================
--- lib/CodeGen/CGExprAgg.cpp
+++ lib/CodeGen/CGExprAgg.cpp
@@ -713,6 +713,7 @@
case CK_CopyAndAutoreleaseBlockObject:
case CK_BuiltinFnToFnPtr:
case CK_ZeroToOCLEvent:
+ case CK_AddressSpaceConversion:
llvm_unreachable("cast kind invalid for aggregate types");
}
}
Index: lib/CodeGen/CGExprComplex.cpp
===================================================================
--- lib/CodeGen/CGExprComplex.cpp
+++ lib/CodeGen/CGExprComplex.cpp
@@ -475,6 +475,7 @@
case CK_CopyAndAutoreleaseBlockObject:
case CK_BuiltinFnToFnPtr:
case CK_ZeroToOCLEvent:
+ case CK_AddressSpaceConversion:
llvm_unreachable("invalid cast kind for complex value");
case CK_FloatingRealToComplex:
Index: lib/CodeGen/CGExprConstant.cpp
===================================================================
--- lib/CodeGen/CGExprConstant.cpp
+++ lib/CodeGen/CGExprConstant.cpp
@@ -633,6 +633,9 @@
return llvm::ConstantStruct::get(STy, Elts);
}
+ case CK_AddressSpaceConversion:
+ return llvm::ConstantExpr::getAddrSpaceCast(C, destType);
+
case CK_LValueToRValue:
case CK_AtomicToNonAtomic:
case CK_NonAtomicToAtomic:
@@ -1062,13 +1065,13 @@
if (!Offset->isNullValue()) {
llvm::Constant *Casted = llvm::ConstantExpr::getBitCast(C, Int8PtrTy);
Casted = llvm::ConstantExpr::getGetElementPtr(Casted, Offset);
- C = llvm::ConstantExpr::getBitCast(Casted, C->getType());
+ C = llvm::ConstantExpr::getPointerCast(Casted, C->getType());
}
// 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,7 +1299,18 @@
case CK_AnyPointerToBlockPointerCast:
case CK_BitCast: {
Value *Src = Visit(const_cast<Expr*>(E));
- return Builder.CreateBitCast(Src, ConvertType(DestTy));
+ llvm::Type *SrcTy = Src->getType();
+ llvm::Type *DstTy = ConvertType(DestTy);
+ if (SrcTy->isPtrOrPtrVectorTy() && DstTy->isPtrOrPtrVectorTy() &&
+ SrcTy->getPointerAddressSpace() != DstTy->getPointerAddressSpace()) {
+ llvm::Type *MidTy = CGF.CGM.getDataLayout().getIntPtrType(SrcTy);
+ return Builder.CreateIntToPtr(Builder.CreatePtrToInt(Src, MidTy), DstTy);
+ }
+ return Builder.CreateBitCast(Src, DstTy);
+ }
+ case CK_AddressSpaceConversion: {
+ Value *Src = Visit(const_cast<Expr*>(E));
+ return Builder.CreateAddrSpaceCast(Src, ConvertType(DestTy));
}
case CK_AtomicToNonAtomic:
case CK_NonAtomicToAtomic:
@@ -1360,7 +1371,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: lib/Edit/RewriteObjCFoundationAPI.cpp
===================================================================
--- lib/Edit/RewriteObjCFoundationAPI.cpp
+++ lib/Edit/RewriteObjCFoundationAPI.cpp
@@ -1033,6 +1033,7 @@
case CK_IntegralComplexToReal:
case CK_IntegralComplexToBoolean:
case CK_AtomicToNonAtomic:
+ case CK_AddressSpaceConversion:
needsCast = true;
break;
Index: lib/Sema/SemaExpr.cpp
===================================================================
--- lib/Sema/SemaExpr.cpp
+++ lib/Sema/SemaExpr.cpp
@@ -4877,8 +4877,13 @@
case Type::STK_BlockPointer:
case Type::STK_ObjCObjectPointer:
switch (DestTy->getScalarTypeKind()) {
- case Type::STK_CPointer:
+ case Type::STK_CPointer: {
+ unsigned SrcAS = SrcTy->getPointeeType().getAddressSpace();
+ unsigned DestAS = DestTy->getPointeeType().getAddressSpace();
+ if (SrcAS != DestAS)
+ return CK_AddressSpaceConversion;
return CK_BitCast;
+ }
case Type::STK_BlockPointer:
return (SrcKind == Type::STK_BlockPointer
? CK_BitCast : CK_AnyPointerToBlockPointerCast);
@@ -7834,10 +7839,14 @@
diagnoseDistinctPointerComparison(*this, Loc, LHS, RHS, /*isError*/false);
}
if (LCanPointeeTy != RCanPointeeTy) {
+ unsigned AddrSpaceL = LCanPointeeTy.getAddressSpace();
+ unsigned AddrSpaceR = RCanPointeeTy.getAddressSpace();
+ CastKind Kind = AddrSpaceL != AddrSpaceR ? CK_AddressSpaceConversion
+ : CK_BitCast;
if (LHSIsNull && !RHSIsNull)
- LHS = ImpCastExprToType(LHS.take(), RHSType, CK_BitCast);
+ LHS = ImpCastExprToType(LHS.take(), RHSType, Kind);
else
- RHS = ImpCastExprToType(RHS.take(), LHSType, CK_BitCast);
+ RHS = ImpCastExprToType(RHS.take(), LHSType, Kind);
}
return ResultTy;
}
Index: lib/StaticAnalyzer/Core/ExprEngineC.cpp
===================================================================
--- lib/StaticAnalyzer/Core/ExprEngineC.cpp
+++ lib/StaticAnalyzer/Core/ExprEngineC.cpp
@@ -286,6 +286,7 @@
case CK_Dependent:
case CK_ArrayToPointerDecay:
case CK_BitCast:
+ case CK_AddressSpaceConversion:
case CK_IntegralCast:
case CK_NullToPointer:
case CK_IntegralToPointer:
Index: test/CodeGenCUDA/address-space-conversion.cu
===================================================================
--- test/CodeGenCUDA/address-space-conversion.cu
+++ test/CodeGenCUDA/address-space-conversion.cu
@@ -0,0 +1,48 @@
+// 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: @_Z27explicit_address_space_castPi
+ __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() {
+ // CHECK: @_Z23pointer_as_array_accessv
+ __shared__ int A[10];
+ int* p = A + 1;
+ p[x] = 0;
+ // CHECK: addrspacecast
+}
+
+__device__ int* callee(int* p) {
+ // CHECK: @_Z6calleePi
+ return p;
+}
+
+__global__ void caller() {
+ // CHECK: @_Z6callerv
+ __shared__ int A[10];
+ __shared__ int* p;
+ p = A;
+ // CHECK: addrspacecast
+
+ ((int*)A)[x] = 42;
+ // CHECK: addrspacecast
+ ((int*)A)[0] = 15;
+ // CHECK: addrspacecast
+
+ 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
+}
_______________________________________________
cfe-commits mailing list
cfe-commits@cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
