Kindly up. -Michele
On 11/22/2013 11:07 AM, Michele Scandale wrote:
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 +}
_______________________________________________ cfe-commits mailing list [email protected] http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
