r364381 - [HIP] Support attribute hip_pinned_shadow
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Tue Jun 25 20:47:37 PDT 2019
Author: yaxunl
Date: Tue Jun 25 20:47:37 2019
New Revision: 364381
URL: http://llvm.org/viewvc/llvm-project?rev=364381&view=rev
Log:
[HIP] Support attribute hip_pinned_shadow
This patch introduces support of hip_pinned_shadow variable for HIP.
A hip_pinned_shadow variable is a global variable with attribute hip_pinned_shadow.
It has external linkage on device side and has no initializer. It has internal
linkage on host side and has initializer or static constructor. It can be accessed
in both device code and host code.
This allows HIP runtime to implement support of HIP texture reference.
Differential Revision: https://reviews.llvm.org/D62738
Added:
cfe/trunk/test/AST/ast-dump-hip-pinned-shadow.cu
cfe/trunk/test/CodeGenCUDA/hip-pinned-shadow.cu
cfe/trunk/test/SemaCUDA/hip-pinned-shadow.cu
Modified:
cfe/trunk/include/clang/Basic/Attr.td
cfe/trunk/include/clang/Basic/AttrDocs.td
cfe/trunk/lib/CodeGen/CodeGenModule.cpp
cfe/trunk/lib/CodeGen/TargetInfo.cpp
cfe/trunk/lib/Driver/ToolChains/HIP.cpp
cfe/trunk/lib/Sema/SemaDeclAttr.cpp
cfe/trunk/test/Driver/hip-toolchain-no-rdc.hip
cfe/trunk/test/Driver/hip-toolchain-rdc.hip
cfe/trunk/test/Misc/pragma-attribute-supported-attributes-list.test
Modified: cfe/trunk/include/clang/Basic/Attr.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/Attr.td?rev=364381&r1=364380&r2=364381&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/Attr.td (original)
+++ cfe/trunk/include/clang/Basic/Attr.td Tue Jun 25 20:47:37 2019
@@ -295,6 +295,7 @@ class LangOpt<string name, code customCo
def MicrosoftExt : LangOpt<"MicrosoftExt">;
def Borland : LangOpt<"Borland">;
def CUDA : LangOpt<"CUDA">;
+def HIP : LangOpt<"HIP">;
def COnly : LangOpt<"COnly", "!LangOpts.CPlusPlus">;
def CPlusPlus : LangOpt<"CPlusPlus">;
def OpenCL : LangOpt<"OpenCL">;
@@ -957,6 +958,13 @@ def CUDADevice : InheritableAttr {
let Documentation = [Undocumented];
}
+def HIPPinnedShadow : InheritableAttr {
+ let Spellings = [GNU<"hip_pinned_shadow">, Declspec<"__hip_pinned_shadow__">];
+ let Subjects = SubjectList<[Var]>;
+ let LangOpts = [HIP];
+ let Documentation = [HIPPinnedShadowDocs];
+}
+
def CUDADeviceBuiltin : IgnoredAttr {
let Spellings = [GNU<"device_builtin">, Declspec<"__device_builtin__">];
let LangOpts = [CUDA];
Modified: cfe/trunk/include/clang/Basic/AttrDocs.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/AttrDocs.td?rev=364381&r1=364380&r2=364381&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/AttrDocs.td (original)
+++ cfe/trunk/include/clang/Basic/AttrDocs.td Tue Jun 25 20:47:37 2019
@@ -4183,3 +4183,15 @@ This attribute does not affect optimizat
``__attribute__((malloc))``.
}];
}
+
+def HIPPinnedShadowDocs : Documentation {
+ let Category = DocCatType;
+ let Content = [{
+The GNU style attribute __attribute__((hip_pinned_shadow)) or MSVC style attribute
+__declspec(hip_pinned_shadow) can be added to the definition of a global variable
+to indicate it is a HIP pinned shadow variable. A HIP pinned shadow variable can
+be accessed on both device side and host side. It has external linkage and is
+not initialized on device side. It has internal linkage and is initialized by
+the initializer on host side.
+ }];
+}
\ No newline at end of file
Modified: cfe/trunk/lib/CodeGen/CodeGenModule.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenModule.cpp?rev=364381&r1=364380&r2=364381&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CodeGenModule.cpp (original)
+++ cfe/trunk/lib/CodeGen/CodeGenModule.cpp Tue Jun 25 20:47:37 2019
@@ -2415,7 +2415,8 @@ void CodeGenModule::EmitGlobal(GlobalDec
if (!Global->hasAttr<CUDADeviceAttr>() &&
!Global->hasAttr<CUDAGlobalAttr>() &&
!Global->hasAttr<CUDAConstantAttr>() &&
- !Global->hasAttr<CUDASharedAttr>())
+ !Global->hasAttr<CUDASharedAttr>() &&
+ !(LangOpts.HIP && Global->hasAttr<HIPPinnedShadowAttr>()))
return;
} else {
// We need to emit host-side 'shadows' for all global
@@ -3781,7 +3782,12 @@ void CodeGenModule::EmitGlobalVarDefinit
!getLangOpts().CUDAIsDevice &&
(D->hasAttr<CUDAConstantAttr>() || D->hasAttr<CUDADeviceAttr>() ||
D->hasAttr<CUDASharedAttr>());
- if (getLangOpts().CUDA && (IsCUDASharedVar || IsCUDAShadowVar))
+ // HIP pinned shadow of initialized host-side global variables are also
+ // left undefined.
+ bool IsHIPPinnedShadowVar =
+ getLangOpts().CUDAIsDevice && D->hasAttr<HIPPinnedShadowAttr>();
+ if (getLangOpts().CUDA &&
+ (IsCUDASharedVar || IsCUDAShadowVar || IsHIPPinnedShadowVar))
Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy));
else if (!InitExpr) {
// This is a tentative definition; tentative definitions are
@@ -3892,7 +3898,8 @@ void CodeGenModule::EmitGlobalVarDefinit
// global variables become internal definitions. These have to
// be internal in order to prevent name conflicts with global
// host variables with the same name in a different TUs.
- if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()) {
+ if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>() ||
+ D->hasAttr<HIPPinnedShadowAttr>()) {
Linkage = llvm::GlobalValue::InternalLinkage;
// Shadow variables and their properties must be registered
@@ -3916,7 +3923,8 @@ void CodeGenModule::EmitGlobalVarDefinit
}
}
- GV->setInitializer(Init);
+ if (!IsHIPPinnedShadowVar)
+ GV->setInitializer(Init);
if (emitter) emitter->finalize(GV);
// If it is safe to mark the global 'constant', do so now.
Modified: cfe/trunk/lib/CodeGen/TargetInfo.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/TargetInfo.cpp?rev=364381&r1=364380&r2=364381&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/TargetInfo.cpp (original)
+++ cfe/trunk/lib/CodeGen/TargetInfo.cpp Tue Jun 25 20:47:37 2019
@@ -7874,12 +7874,24 @@ static bool requiresAMDGPUProtectedVisib
return D->hasAttr<OpenCLKernelAttr>() ||
(isa<FunctionDecl>(D) && D->hasAttr<CUDAGlobalAttr>()) ||
(isa<VarDecl>(D) &&
- (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()));
+ (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>() ||
+ D->hasAttr<HIPPinnedShadowAttr>()));
+}
+
+static bool requiresAMDGPUDefaultVisibility(const Decl *D,
+ llvm::GlobalValue *GV) {
+ if (GV->getVisibility() != llvm::GlobalValue::HiddenVisibility)
+ return false;
+
+ return isa<VarDecl>(D) && D->hasAttr<HIPPinnedShadowAttr>();
}
void AMDGPUTargetCodeGenInfo::setTargetAttributes(
const Decl *D, llvm::GlobalValue *GV, CodeGen::CodeGenModule &M) const {
- if (requiresAMDGPUProtectedVisibility(D, GV)) {
+ if (requiresAMDGPUDefaultVisibility(D, GV)) {
+ GV->setVisibility(llvm::GlobalValue::DefaultVisibility);
+ GV->setDSOLocal(false);
+ } else if (requiresAMDGPUProtectedVisibility(D, GV)) {
GV->setVisibility(llvm::GlobalValue::ProtectedVisibility);
GV->setDSOLocal(true);
}
Modified: cfe/trunk/lib/Driver/ToolChains/HIP.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Driver/ToolChains/HIP.cpp?rev=364381&r1=364380&r2=364381&view=diff
==============================================================================
--- cfe/trunk/lib/Driver/ToolChains/HIP.cpp (original)
+++ cfe/trunk/lib/Driver/ToolChains/HIP.cpp Tue Jun 25 20:47:37 2019
@@ -170,9 +170,8 @@ void AMDGCN::Linker::constructLldCommand
const char *InputFileName) const {
// Construct lld command.
// The output from ld.lld is an HSA code object file.
- ArgStringList LldArgs{"-flavor", "gnu", "--no-undefined",
- "-shared", "-o", Output.getFilename(),
- InputFileName};
+ ArgStringList LldArgs{
+ "-flavor", "gnu", "-shared", "-o", Output.getFilename(), InputFileName};
SmallString<128> LldPath(C.getDriver().Dir);
llvm::sys::path::append(LldPath, "lld");
const char *Lld = Args.MakeArgString(LldPath);
Modified: cfe/trunk/lib/Sema/SemaDeclAttr.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDeclAttr.cpp?rev=364381&r1=364380&r2=364381&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDeclAttr.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp Tue Jun 25 20:47:37 2019
@@ -6786,6 +6786,10 @@ static void ProcessDeclAttribute(Sema &S
case ParsedAttr::AT_CUDAHost:
handleSimpleAttributeWithExclusions<CUDAHostAttr, CUDAGlobalAttr>(S, D, AL);
break;
+ case ParsedAttr::AT_HIPPinnedShadow:
+ handleSimpleAttributeWithExclusions<HIPPinnedShadowAttr, CUDADeviceAttr,
+ CUDAConstantAttr>(S, D, AL);
+ break;
case ParsedAttr::AT_GNUInline:
handleGNUInlineAttr(S, D, AL);
break;
Added: cfe/trunk/test/AST/ast-dump-hip-pinned-shadow.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/AST/ast-dump-hip-pinned-shadow.cu?rev=364381&view=auto
==============================================================================
--- cfe/trunk/test/AST/ast-dump-hip-pinned-shadow.cu (added)
+++ cfe/trunk/test/AST/ast-dump-hip-pinned-shadow.cu Tue Jun 25 20:47:37 2019
@@ -0,0 +1,13 @@
+// RUN: %clang_cc1 -fcuda-is-device -ast-dump -ast-dump-filter tex -x hip %s | FileCheck -strict-whitespace %s
+// RUN: %clang_cc1 -ast-dump -ast-dump-filter tex -x hip %s | FileCheck -strict-whitespace %s
+struct textureReference {
+ int a;
+};
+
+// CHECK: HIPPinnedShadowAttr
+template <class T, int texType, int hipTextureReadMode>
+struct texture : public textureReference {
+texture() { a = 1; }
+};
+
+__attribute__((hip_pinned_shadow)) texture<float, 1, 1> tex;
Added: cfe/trunk/test/CodeGenCUDA/hip-pinned-shadow.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/hip-pinned-shadow.cu?rev=364381&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/hip-pinned-shadow.cu (added)
+++ cfe/trunk/test/CodeGenCUDA/hip-pinned-shadow.cu Tue Jun 25 20:47:37 2019
@@ -0,0 +1,23 @@
+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -std=c++11 -fvisibility hidden -fapply-global-visibility-to-externs \
+// RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefixes=HIPDEV %s
+// RUN: %clang_cc1 -triple x86_64 -std=c++11 \
+// RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefixes=HIPHOST %s
+
+struct textureReference {
+ int a;
+};
+
+template <class T, int texType, int hipTextureReadMode>
+struct texture : public textureReference {
+texture() { a = 1; }
+};
+
+__attribute__((hip_pinned_shadow)) texture<float, 2, 1> tex;
+// CUDADEV-NOT: @tex
+// CUDAHOST-NOT: call i32 @__hipRegisterVar{{.*}}@tex
+// HIPDEV: @tex = external addrspace(1) global %struct.texture
+// HIPDEV-NOT: declare{{.*}}void @_ZN7textureIfLi2ELi1EEC1Ev
+// HIPHOST: define{{.*}}@_ZN7textureIfLi2ELi1EEC1Ev
+// HIPHOST: call i32 @__hipRegisterVar{{.*}}@tex{{.*}}i32 0, i32 4, i32 0, i32 0)
Modified: cfe/trunk/test/Driver/hip-toolchain-no-rdc.hip
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/Driver/hip-toolchain-no-rdc.hip?rev=364381&r1=364380&r2=364381&view=diff
==============================================================================
--- cfe/trunk/test/Driver/hip-toolchain-no-rdc.hip (original)
+++ cfe/trunk/test/Driver/hip-toolchain-no-rdc.hip Tue Jun 25 20:47:37 2019
@@ -37,7 +37,7 @@
// CHECK-SAME: "-filetype=obj"
// CHECK-SAME: "-mcpu=gfx803" "-o" [[OBJ_DEV_A_803:".*-gfx803-.*o"]]
-// CHECK: [[LLD: ".*lld"]] "-flavor" "gnu" "--no-undefined" "-shared"
+// CHECK: [[LLD: ".*lld"]] "-flavor" "gnu" "-shared"
// CHECK-SAME: "-o" "[[IMG_DEV_A_803:.*out]]" [[OBJ_DEV_A_803]]
//
@@ -65,7 +65,7 @@
// CHECK-SAME: "-filetype=obj"
// CHECK-SAME: "-mcpu=gfx900" "-o" [[OBJ_DEV_A_900:".*-gfx900-.*o"]]
-// CHECK: [[LLD: ".*lld"]] "-flavor" "gnu" "--no-undefined" "-shared"
+// CHECK: [[LLD: ".*lld"]] "-flavor" "gnu" "-shared"
// CHECK-SAME: "-o" "[[IMG_DEV_A_900:.*out]]" [[OBJ_DEV_A_900]]
//
@@ -109,7 +109,7 @@
// CHECK-SAME: "-filetype=obj"
// CHECK-SAME: "-mcpu=gfx803" "-o" [[OBJ_DEV_B_803:".*-gfx803-.*o"]]
-// CHECK: [[LLD: ".*lld"]] "-flavor" "gnu" "--no-undefined" "-shared"
+// CHECK: [[LLD: ".*lld"]] "-flavor" "gnu" "-shared"
// CHECK-SAME: "-o" "[[IMG_DEV_B_803:.*out]]" [[OBJ_DEV_B_803]]
//
@@ -137,7 +137,7 @@
// CHECK-SAME: "-filetype=obj"
// CHECK-SAME: "-mcpu=gfx900" "-o" [[OBJ_DEV_B_900:".*-gfx900-.*o"]]
-// CHECK: [[LLD: ".*lld"]] "-flavor" "gnu" "--no-undefined" "-shared"
+// CHECK: [[LLD: ".*lld"]] "-flavor" "gnu" "-shared"
// CHECK-SAME: "-o" "[[IMG_DEV_B_900:.*out]]" [[OBJ_DEV_B_900]]
//
Modified: cfe/trunk/test/Driver/hip-toolchain-rdc.hip
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/Driver/hip-toolchain-rdc.hip?rev=364381&r1=364380&r2=364381&view=diff
==============================================================================
--- cfe/trunk/test/Driver/hip-toolchain-rdc.hip (original)
+++ cfe/trunk/test/Driver/hip-toolchain-rdc.hip Tue Jun 25 20:47:37 2019
@@ -43,7 +43,7 @@
// CHECK-SAME: "-filetype=obj"
// CHECK-SAME: "-mcpu=gfx803" "-o" [[OBJ_DEV1:".*-gfx803-.*o"]]
-// CHECK: [[LLD: ".*lld"]] "-flavor" "gnu" "--no-undefined" "-shared"
+// CHECK: [[LLD: ".*lld"]] "-flavor" "gnu" "-shared"
// CHECK-SAME: "-o" "[[IMG_DEV1:.*out]]" [[OBJ_DEV1]]
// CHECK: [[CLANG]] "-cc1" "-triple" "amdgcn-amd-amdhsa"
@@ -75,7 +75,7 @@
// CHECK-SAME: "-filetype=obj"
// CHECK-SAME: "-mcpu=gfx900" "-o" [[OBJ_DEV2:".*-gfx900-.*o"]]
-// CHECK: [[LLD]] "-flavor" "gnu" "--no-undefined" "-shared"
+// CHECK: [[LLD]] "-flavor" "gnu" "-shared"
// CHECK-SAME: "-o" "[[IMG_DEV2:.*out]]" [[OBJ_DEV2]]
// CHECK: [[CLANG]] "-cc1" "-triple" "x86_64-unknown-linux-gnu"
Modified: cfe/trunk/test/Misc/pragma-attribute-supported-attributes-list.test
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/Misc/pragma-attribute-supported-attributes-list.test?rev=364381&r1=364380&r2=364381&view=diff
==============================================================================
--- cfe/trunk/test/Misc/pragma-attribute-supported-attributes-list.test (original)
+++ cfe/trunk/test/Misc/pragma-attribute-supported-attributes-list.test Tue Jun 25 20:47:37 2019
@@ -53,6 +53,7 @@
// CHECK-NEXT: FlagEnum (SubjectMatchRule_enum)
// CHECK-NEXT: Flatten (SubjectMatchRule_function)
// CHECK-NEXT: GNUInline (SubjectMatchRule_function)
+// CHECK-NEXT: HIPPinnedShadow (SubjectMatchRule_variable)
// CHECK-NEXT: Hot (SubjectMatchRule_function)
// CHECK-NEXT: IBAction (SubjectMatchRule_objc_method_is_instance)
// CHECK-NEXT: IFunc (SubjectMatchRule_function)
Added: cfe/trunk/test/SemaCUDA/hip-pinned-shadow.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/hip-pinned-shadow.cu?rev=364381&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/hip-pinned-shadow.cu (added)
+++ cfe/trunk/test/SemaCUDA/hip-pinned-shadow.cu Tue Jun 25 20:47:37 2019
@@ -0,0 +1,25 @@
+// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -std=c++11 -fvisibility hidden -fapply-global-visibility-to-externs \
+// RUN: -emit-llvm -o - -x hip %s -fsyntax-only -verify
+// RUN: %clang_cc1 -triple x86_64 -std=c++11 \
+// RUN: -emit-llvm -o - -x hip %s -fsyntax-only -verify
+
+#define __device__ __attribute__((device))
+#define __constant__ __attribute__((constant))
+#define __hip_pinned_shadow__ __attribute((hip_pinned_shadow))
+
+struct textureReference {
+ int a;
+};
+
+template <class T, int texType, int hipTextureReadMode>
+struct texture : public textureReference {
+texture() { a = 1; }
+};
+
+__hip_pinned_shadow__ texture<float, 2, 1> tex;
+__device__ __hip_pinned_shadow__ texture<float, 2, 1> tex2; // expected-error{{'hip_pinned_shadow' and 'device' attributes are not compatible}}
+ // expected-error at -1{{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables}}
+ // expected-note at -2{{conflicting attribute is here}}
+__constant__ __hip_pinned_shadow__ texture<float, 2, 1> tex3; // expected-error{{'hip_pinned_shadow' and 'constant' attributes are not compatible}}
+ // expected-error at -1{{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables}}
+ // expected-note at -2{{conflicting attribute is here}}
More information about the cfe-commits
mailing list