From: Eli Bendersky Date: Mon, 24 Mar 2014 22:05:38 +0000 (+0000) Subject: Proper handling of static local variables with address space qualifiers. X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=6e06873ac6b644d94907eb274133f5f30e0bce91;p=clang Proper handling of static local variables with address space qualifiers. Similar to the implementation for globals in r157167. Patch by Jingyue Wu. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@204677 91177308-0d34-0410-b5e6-96231b3b80d8 --- diff --git a/lib/CodeGen/CGDecl.cpp b/lib/CodeGen/CGDecl.cpp index e16845c680..430b56caa0 100644 --- a/lib/CodeGen/CGDecl.cpp +++ b/lib/CodeGen/CGDecl.cpp @@ -183,7 +183,7 @@ static std::string GetStaticDeclName(CodeGenFunction &CGF, const VarDecl &D, return ContextName + Separator + D.getNameAsString(); } -llvm::GlobalVariable * +llvm::Constant * CodeGenFunction::CreateStaticVarDecl(const VarDecl &D, const char *Separator, llvm::GlobalValue::LinkageTypes Linkage) { @@ -212,6 +212,13 @@ CodeGenFunction::CreateStaticVarDecl(const VarDecl &D, if (D.getTLSKind()) CGM.setTLSMode(GV, D); + // Make sure the result is of the correct type. + unsigned ExpectedAddrSpace = CGM.getContext().getTargetAddressSpace(Ty); + if (AddrSpace != ExpectedAddrSpace) { + llvm::PointerType *PTy = llvm::PointerType::get(LTy, ExpectedAddrSpace); + return llvm::ConstantExpr::getAddrSpaceCast(GV, PTy); + } + return GV; } @@ -298,12 +305,8 @@ void CodeGenFunction::EmitStaticVarDecl(const VarDecl &D, llvm::Constant *addr = CGM.getStaticLocalDeclAddress(&D); - llvm::GlobalVariable *var; - if (addr) { - var = cast(addr->stripPointerCasts()); - } else { - addr = var = CreateStaticVarDecl(D, ".", Linkage); - } + if (!addr) + addr = CreateStaticVarDecl(D, ".", Linkage); // Store into LocalDeclMap before generating initializer to handle // circular references. @@ -319,6 +322,8 @@ void CodeGenFunction::EmitStaticVarDecl(const VarDecl &D, // Save the type in case adding the initializer forces a type change. llvm::Type *expectedType = addr->getType(); + llvm::GlobalVariable *var = + cast(addr->stripPointerCasts()); // If this value has an initializer, emit it. if (D.getInit()) var = AddInitializerToStaticVarDecl(D, var); @@ -339,7 +344,8 @@ void CodeGenFunction::EmitStaticVarDecl(const VarDecl &D, // // FIXME: It is really dangerous to store this in the map; if anyone // RAUW's the GV uses of this constant will be invalid. - llvm::Constant *castedAddr = llvm::ConstantExpr::getBitCast(var, expectedType); + llvm::Constant *castedAddr = + llvm::ConstantExpr::getPointerBitCastOrAddrSpaceCast(var, expectedType); DMEntry = castedAddr; CGM.setStaticLocalDeclAddress(&D, castedAddr); diff --git a/lib/CodeGen/CodeGenFunction.h b/lib/CodeGen/CodeGenFunction.h index bbc4fae992..c583a30548 100644 --- a/lib/CodeGen/CodeGenFunction.h +++ b/lib/CodeGen/CodeGenFunction.h @@ -2336,9 +2336,9 @@ public: /// CreateStaticVarDecl - Create a zero-initialized LLVM global for /// a static local variable. - llvm::GlobalVariable *CreateStaticVarDecl(const VarDecl &D, - const char *Separator, - llvm::GlobalValue::LinkageTypes Linkage); + llvm::Constant *CreateStaticVarDecl(const VarDecl &D, + const char *Separator, + llvm::GlobalValue::LinkageTypes Linkage); /// AddInitializerToStaticVarDecl - Add the initializer for 'D' to the /// global variable that has already been created for it. If the initializer diff --git a/test/CodeGenCUDA/address-spaces.cu b/test/CodeGenCUDA/address-spaces.cu index 04344526f4..a28886f4b1 100644 --- a/test/CodeGenCUDA/address-spaces.cu +++ b/test/CodeGenCUDA/address-spaces.cu @@ -1,5 +1,8 @@ // RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple nvptx-unknown-unknown | FileCheck %s +// Verifies Clang emits correct address spaces and addrspacecast instructions +// for CUDA code. + #include "../SemaCUDA/cuda.h" // CHECK: @i = addrspace(1) global @@ -11,6 +14,18 @@ __constant__ int j; // CHECK: @k = addrspace(3) global __shared__ int k; +struct MyStruct { + int data1; + int data2; +}; + +// CHECK: @_ZZ5func0vE1a = internal addrspace(3) global %struct.MyStruct zeroinitializer +// CHECK: @_ZZ5func1vE1a = internal addrspace(3) global float 0.000000e+00 +// CHECK: @_ZZ5func2vE1a = internal addrspace(3) global [256 x float] zeroinitializer +// CHECK: @_ZZ5func3vE1a = internal addrspace(3) global float 0.000000e+00 +// CHECK: @_ZZ5func4vE1a = internal addrspace(3) global float 0.000000e+00 +// CHECK: @b = addrspace(3) global float 0.000000e+00 + __device__ void foo() { // CHECK: load i32* addrspacecast (i32 addrspace(1)* @i to i32*) i++; @@ -22,15 +37,66 @@ __device__ void foo() { k++; static int li; - // CHECK: load i32 addrspace(1)* @_ZZ3foovE2li + // CHECK: load i32* addrspacecast (i32 addrspace(1)* @_ZZ3foovE2li to i32*) li++; __constant__ int lj; - // CHECK: load i32 addrspace(4)* @_ZZ3foovE2lj + // CHECK: load i32* addrspacecast (i32 addrspace(4)* @_ZZ3foovE2lj to i32*) lj++; __shared__ int lk; - // CHECK: load i32 addrspace(3)* @_ZZ3foovE2lk + // CHECK: load i32* addrspacecast (i32 addrspace(3)* @_ZZ3foovE2lk to i32*) lk++; } +__device__ void func0() { + __shared__ MyStruct a; + MyStruct *ap = &a; // composite type + ap->data1 = 1; + ap->data2 = 2; +} +// CHECK: define void @_Z5func0v() +// CHECK: store %struct.MyStruct* addrspacecast (%struct.MyStruct addrspace(3)* @_ZZ5func0vE1a to %struct.MyStruct*), %struct.MyStruct** %ap + +__device__ void callee(float *ap) { + *ap = 1.0f; +} + +__device__ void func1() { + __shared__ float a; + callee(&a); // implicit cast from parameters +} +// CHECK: define void @_Z5func1v() +// CHECK: call void @_Z6calleePf(float* addrspacecast (float addrspace(3)* @_ZZ5func1vE1a to float*)) + +__device__ void func2() { + __shared__ float a[256]; + float *ap = &a[128]; // implicit cast from a decayed array + *ap = 1.0f; +} +// CHECK: define void @_Z5func2v() +// CHECK: store float* getelementptr inbounds ([256 x float]* addrspacecast ([256 x float] addrspace(3)* @_ZZ5func2vE1a to [256 x float]*), i32 0, i32 128), float** %ap + +__device__ void func3() { + __shared__ float a; + float *ap = reinterpret_cast(&a); // explicit cast + *ap = 1.0f; +} +// CHECK: define void @_Z5func3v() +// CHECK: store float* addrspacecast (float addrspace(3)* @_ZZ5func3vE1a to float*), float** %ap + +__device__ void func4() { + __shared__ float a; + float *ap = (float *)&a; // explicit c-style cast + *ap = 1.0f; +} +// CHECK: define void @_Z5func4v() +// CHECK: store float* addrspacecast (float addrspace(3)* @_ZZ5func4vE1a to float*), float** %ap + +__shared__ float b; + +__device__ float *func5() { + return &b; // implicit cast from a return value +} +// CHECK: define float* @_Z5func5v() +// CHECK: ret float* addrspacecast (float addrspace(3)* @b to float*)