From 8d068d6ad41e3e39a0ce34148f968eadd48b53bb Mon Sep 17 00:00:00 2001 From: Yaxun Liu Date: Wed, 26 Jun 2019 03:47:37 +0000 Subject: [PATCH] [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 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@364381 91177308-0d34-0410-b5e6-96231b3b80d8 --- include/clang/Basic/Attr.td | 8 ++++++ include/clang/Basic/AttrDocs.td | 12 +++++++++ lib/CodeGen/CodeGenModule.cpp | 16 +++++++++--- lib/CodeGen/TargetInfo.cpp | 16 ++++++++++-- lib/Driver/ToolChains/HIP.cpp | 5 ++-- lib/Sema/SemaDeclAttr.cpp | 4 +++ test/AST/ast-dump-hip-pinned-shadow.cu | 13 ++++++++++ test/CodeGenCUDA/hip-pinned-shadow.cu | 23 +++++++++++++++++ test/Driver/hip-toolchain-no-rdc.hip | 8 +++--- test/Driver/hip-toolchain-rdc.hip | 4 +-- ...a-attribute-supported-attributes-list.test | 1 + test/SemaCUDA/hip-pinned-shadow.cu | 25 +++++++++++++++++++ 12 files changed, 120 insertions(+), 15 deletions(-) create mode 100644 test/AST/ast-dump-hip-pinned-shadow.cu create mode 100644 test/CodeGenCUDA/hip-pinned-shadow.cu create mode 100644 test/SemaCUDA/hip-pinned-shadow.cu diff --git a/include/clang/Basic/Attr.td b/include/clang/Basic/Attr.td index 5dd9cfc958..93913b043e 100644 --- a/include/clang/Basic/Attr.td +++ b/include/clang/Basic/Attr.td @@ -295,6 +295,7 @@ class LangOpt { 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]; diff --git a/include/clang/Basic/AttrDocs.td b/include/clang/Basic/AttrDocs.td index 97e50f9236..fac6116057 100644 --- a/include/clang/Basic/AttrDocs.td +++ b/include/clang/Basic/AttrDocs.td @@ -4183,3 +4183,15 @@ This attribute does not affect optimizations in any way, unlike GCC's ``__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 diff --git a/lib/CodeGen/CodeGenModule.cpp b/lib/CodeGen/CodeGenModule.cpp index 87a1d45abf..b0e3b0bb98 100644 --- a/lib/CodeGen/CodeGenModule.cpp +++ b/lib/CodeGen/CodeGenModule.cpp @@ -2415,7 +2415,8 @@ void CodeGenModule::EmitGlobal(GlobalDecl GD) { if (!Global->hasAttr() && !Global->hasAttr() && !Global->hasAttr() && - !Global->hasAttr()) + !Global->hasAttr() && + !(LangOpts.HIP && Global->hasAttr())) return; } else { // We need to emit host-side 'shadows' for all global @@ -3781,7 +3782,12 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D, !getLangOpts().CUDAIsDevice && (D->hasAttr() || D->hasAttr() || D->hasAttr()); - if (getLangOpts().CUDA && (IsCUDASharedVar || IsCUDAShadowVar)) + // HIP pinned shadow of initialized host-side global variables are also + // left undefined. + bool IsHIPPinnedShadowVar = + getLangOpts().CUDAIsDevice && D->hasAttr(); + 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::EmitGlobalVarDefinition(const VarDecl *D, // 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() || D->hasAttr()) { + if (D->hasAttr() || D->hasAttr() || + D->hasAttr()) { Linkage = llvm::GlobalValue::InternalLinkage; // Shadow variables and their properties must be registered @@ -3916,7 +3923,8 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D, } } - 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. diff --git a/lib/CodeGen/TargetInfo.cpp b/lib/CodeGen/TargetInfo.cpp index 8563a18b2e..22f70a08b0 100644 --- a/lib/CodeGen/TargetInfo.cpp +++ b/lib/CodeGen/TargetInfo.cpp @@ -7874,12 +7874,24 @@ static bool requiresAMDGPUProtectedVisibility(const Decl *D, return D->hasAttr() || (isa(D) && D->hasAttr()) || (isa(D) && - (D->hasAttr() || D->hasAttr())); + (D->hasAttr() || D->hasAttr() || + D->hasAttr())); +} + +static bool requiresAMDGPUDefaultVisibility(const Decl *D, + llvm::GlobalValue *GV) { + if (GV->getVisibility() != llvm::GlobalValue::HiddenVisibility) + return false; + + return isa(D) && D->hasAttr(); } 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); } diff --git a/lib/Driver/ToolChains/HIP.cpp b/lib/Driver/ToolChains/HIP.cpp index a60485ab03..2ec97e798f 100644 --- a/lib/Driver/ToolChains/HIP.cpp +++ b/lib/Driver/ToolChains/HIP.cpp @@ -170,9 +170,8 @@ void AMDGCN::Linker::constructLldCommand(Compilation &C, const JobAction &JA, 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); diff --git a/lib/Sema/SemaDeclAttr.cpp b/lib/Sema/SemaDeclAttr.cpp index 8670f74d95..5a1712ff13 100644 --- a/lib/Sema/SemaDeclAttr.cpp +++ b/lib/Sema/SemaDeclAttr.cpp @@ -6786,6 +6786,10 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, case ParsedAttr::AT_CUDAHost: handleSimpleAttributeWithExclusions(S, D, AL); break; + case ParsedAttr::AT_HIPPinnedShadow: + handleSimpleAttributeWithExclusions(S, D, AL); + break; case ParsedAttr::AT_GNUInline: handleGNUInlineAttr(S, D, AL); break; diff --git a/test/AST/ast-dump-hip-pinned-shadow.cu b/test/AST/ast-dump-hip-pinned-shadow.cu new file mode 100644 index 0000000000..53d7c8fbed --- /dev/null +++ b/test/AST/ast-dump-hip-pinned-shadow.cu @@ -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 +struct texture : public textureReference { +texture() { a = 1; } +}; + +__attribute__((hip_pinned_shadow)) texture tex; diff --git a/test/CodeGenCUDA/hip-pinned-shadow.cu b/test/CodeGenCUDA/hip-pinned-shadow.cu new file mode 100644 index 0000000000..75798f7e1d --- /dev/null +++ b/test/CodeGenCUDA/hip-pinned-shadow.cu @@ -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 +struct texture : public textureReference { +texture() { a = 1; } +}; + +__attribute__((hip_pinned_shadow)) texture 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) diff --git a/test/Driver/hip-toolchain-no-rdc.hip b/test/Driver/hip-toolchain-no-rdc.hip index 229839db6c..540b932860 100644 --- a/test/Driver/hip-toolchain-no-rdc.hip +++ b/test/Driver/hip-toolchain-no-rdc.hip @@ -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]] // diff --git a/test/Driver/hip-toolchain-rdc.hip b/test/Driver/hip-toolchain-rdc.hip index 055efe69ad..15ac5f1931 100644 --- a/test/Driver/hip-toolchain-rdc.hip +++ b/test/Driver/hip-toolchain-rdc.hip @@ -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" diff --git a/test/Misc/pragma-attribute-supported-attributes-list.test b/test/Misc/pragma-attribute-supported-attributes-list.test index aa71e0bc8f..fc9d86efe7 100644 --- a/test/Misc/pragma-attribute-supported-attributes-list.test +++ b/test/Misc/pragma-attribute-supported-attributes-list.test @@ -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) diff --git a/test/SemaCUDA/hip-pinned-shadow.cu b/test/SemaCUDA/hip-pinned-shadow.cu new file mode 100644 index 0000000000..c58f7097af --- /dev/null +++ b/test/SemaCUDA/hip-pinned-shadow.cu @@ -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 +struct texture : public textureReference { +texture() { a = 1; } +}; + +__hip_pinned_shadow__ texture tex; +__device__ __hip_pinned_shadow__ texture tex2; // expected-error{{'hip_pinned_shadow' and 'device' attributes are not compatible}} + // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables}} + // expected-note@-2{{conflicting attribute is here}} +__constant__ __hip_pinned_shadow__ texture tex3; // expected-error{{'hip_pinned_shadow' and 'constant' attributes are not compatible}} + // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables}} + // expected-note@-2{{conflicting attribute is here}} -- 2.40.0