From: Michael Liao Date: Thu, 26 Mar 2020 15:21:45 +0000 (-0400) Subject: [hip] Remove `hip_pinned_shadow`. X-Git-Tag: llvmorg-12-init~9806 X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=c97be2c377852fad7eb38409aae5692fa417e49b;p=platform%2Fupstream%2Fllvm.git [hip] Remove `hip_pinned_shadow`. Summary: - Use `device_builtin_surface` and `device_builtin_texture` for surface/texture reference support. So far, both the host and device use the same reference type, which could be revised later when interface/implementation is stablized. Reviewers: yaxunl Subscribers: cfe-commits Tags: #clang Differential Revision: https://reviews.llvm.org/D77583 --- diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index f55ce2c..c586f9b 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -322,7 +322,6 @@ class LangOpt { def MicrosoftExt : LangOpt<"MicrosoftExt">; def Borland : LangOpt<"Borland">; def CUDA : LangOpt<"CUDA">; -def HIP : LangOpt<"HIP">; def SYCL : LangOpt<"SYCLIsDevice">; def COnly : LangOpt<"", "!LangOpts.CPlusPlus">; def CPlusPlus : LangOpt<"CPlusPlus">; @@ -1052,13 +1051,6 @@ 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/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index fb1c82a..36561c04 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -4613,18 +4613,6 @@ only call one function. }]; } -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. - }]; -} - def CUDADeviceBuiltinSurfaceTypeDocs : Documentation { let Category = DocCatType; let Content = [{ diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 1645a9e..8b7d52b 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1955,9 +1955,9 @@ void CodeGenModule::SetFunctionAttributes(GlobalDecl GD, llvm::Function *F, } } -void CodeGenModule::addUsedGlobal(llvm::GlobalValue *GV, bool SkipCheck) { - assert(SkipCheck || (!GV->isDeclaration() && - "Only globals with definition can force usage.")); +void CodeGenModule::addUsedGlobal(llvm::GlobalValue *GV) { + assert(!GV->isDeclaration() && + "Only globals with definition can force usage."); LLVMUsed.emplace_back(GV); } @@ -2520,7 +2520,6 @@ void CodeGenModule::EmitGlobal(GlobalDecl GD) { !Global->hasAttr() && !Global->hasAttr() && !Global->hasAttr() && - !(LangOpts.HIP && Global->hasAttr()) && !Global->getType()->isCUDADeviceBuiltinSurfaceType() && !Global->getType()->isCUDADeviceBuiltinTextureType()) return; @@ -3928,10 +3927,8 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D, D->getType()->isCUDADeviceBuiltinTextureType()); // HIP pinned shadow of initialized host-side global variables are also // left undefined. - bool IsHIPPinnedShadowVar = - getLangOpts().CUDAIsDevice && D->hasAttr(); - if (getLangOpts().CUDA && (IsCUDASharedVar || IsCUDAShadowVar || - IsCUDADeviceShadowVar || IsHIPPinnedShadowVar)) + if (getLangOpts().CUDA && + (IsCUDASharedVar || IsCUDAShadowVar || IsCUDADeviceShadowVar)) Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy)); else if (D->hasAttr()) Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy)); @@ -4039,8 +4036,7 @@ 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() || - D->hasAttr()) { + if (D->hasAttr() || D->hasAttr()) { Linkage = llvm::GlobalValue::InternalLinkage; // Shadow variables and their properties must be registered with CUDA // runtime. Skip Extern global variables, which will be registered in @@ -4087,15 +4083,7 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D, } } - // HIPPinnedShadowVar should remain in the final code object irrespective of - // whether it is used or not within the code. Add it to used list, so that - // it will not get eliminated when it is unused. Also, it is an extern var - // within device code, and it should *not* get initialized within device code. - if (IsHIPPinnedShadowVar) - addUsedGlobal(GV, /*SkipCheck=*/true); - else - GV->setInitializer(Init); - + GV->setInitializer(Init); if (emitter) emitter->finalize(GV); diff --git a/clang/lib/CodeGen/CodeGenModule.h b/clang/lib/CodeGen/CodeGenModule.h index dc2c61b..a84c5bd 100644 --- a/clang/lib/CodeGen/CodeGenModule.h +++ b/clang/lib/CodeGen/CodeGenModule.h @@ -1037,7 +1037,7 @@ public: void MaybeHandleStaticInExternC(const SomeDecl *D, llvm::GlobalValue *GV); /// Add a global to a list to be added to the llvm.used metadata. - void addUsedGlobal(llvm::GlobalValue *GV, bool SkipCheck = false); + void addUsedGlobal(llvm::GlobalValue *GV); /// Add a global to a list to be added to the llvm.compiler.used metadata. void addCompilerUsedGlobal(llvm::GlobalValue *GV); diff --git a/clang/lib/CodeGen/TargetInfo.cpp b/clang/lib/CodeGen/TargetInfo.cpp index e2825a3..019c533 100644 --- a/clang/lib/CodeGen/TargetInfo.cpp +++ b/clang/lib/CodeGen/TargetInfo.cpp @@ -8407,23 +8407,13 @@ static bool requiresAMDGPUProtectedVisibility(const Decl *D, (isa(D) && D->hasAttr()) || (isa(D) && (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(); + cast(D)->getType()->isCUDADeviceBuiltinSurfaceType() || + cast(D)->getType()->isCUDADeviceBuiltinTextureType())); } void AMDGPUTargetCodeGenInfo::setTargetAttributes( const Decl *D, llvm::GlobalValue *GV, CodeGen::CodeGenModule &M) const { - if (requiresAMDGPUDefaultVisibility(D, GV)) { - GV->setVisibility(llvm::GlobalValue::DefaultVisibility); - GV->setDSOLocal(false); - } else if (requiresAMDGPUProtectedVisibility(D, GV)) { + if (requiresAMDGPUProtectedVisibility(D, GV)) { GV->setVisibility(llvm::GlobalValue::ProtectedVisibility); GV->setDSOLocal(true); } diff --git a/clang/lib/Driver/ToolChains/HIP.cpp b/clang/lib/Driver/ToolChains/HIP.cpp index e4ace81..8676f8b 100644 --- a/clang/lib/Driver/ToolChains/HIP.cpp +++ b/clang/lib/Driver/ToolChains/HIP.cpp @@ -192,8 +192,9 @@ 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", "-shared", "-o", Output.getFilename(), InputFileName}; + ArgStringList LldArgs{"-flavor", "gnu", "--no-undefined", + "-shared", "-o", Output.getFilename(), + InputFileName}; const char *Lld = Args.MakeArgString(getToolChain().GetProgramPath("lld")); C.addCommand(std::make_unique(JA, *this, Lld, LldArgs, Inputs)); } diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 5aacb2f..66f5d04 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -6930,10 +6930,6 @@ 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_CUDADeviceBuiltinSurfaceType: handleSimpleAttributeWithExclusions(S, D, diff --git a/clang/test/AST/ast-dump-hip-pinned-shadow.cu b/clang/test/AST/ast-dump-hip-pinned-shadow.cu deleted file mode 100644 index 53d7c8f..0000000 --- a/clang/test/AST/ast-dump-hip-pinned-shadow.cu +++ /dev/null @@ -1,13 +0,0 @@ -// 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/clang/test/Driver/hip-toolchain-no-rdc.hip b/clang/test/Driver/hip-toolchain-no-rdc.hip index 4371334..fd5cf1b 100644 --- a/clang/test/Driver/hip-toolchain-no-rdc.hip +++ b/clang/test/Driver/hip-toolchain-no-rdc.hip @@ -38,7 +38,7 @@ // CHECK-SAME: "-filetype=obj" // CHECK-SAME: "-o" [[OBJ_DEV_A_803:".*-gfx803-.*o"]] -// CHECK: [[LLD: ".*lld.*"]] "-flavor" "gnu" "-shared" +// CHECK: [[LLD: ".*lld.*"]] "-flavor" "gnu" "--no-undefined" "-shared" // CHECK-SAME: "-o" "[[IMG_DEV_A_803:.*out]]" [[OBJ_DEV_A_803]] // @@ -67,7 +67,7 @@ // CHECK-SAME: "-filetype=obj" // CHECK-SAME: "-o" [[OBJ_DEV_A_900:".*-gfx900-.*o"]] -// CHECK: [[LLD]] "-flavor" "gnu" "-shared" +// CHECK: [[LLD]] "-flavor" "gnu" "--no-undefined" "-shared" // CHECK-SAME: "-o" "[[IMG_DEV_A_900:.*out]]" [[OBJ_DEV_A_900]] // @@ -112,7 +112,7 @@ // CHECK-SAME: "-filetype=obj" // CHECK-SAME: "-o" [[OBJ_DEV_B_803:".*-gfx803-.*o"]] -// CHECK: [[LLD]] "-flavor" "gnu" "-shared" +// CHECK: [[LLD]] "-flavor" "gnu" "--no-undefined" "-shared" // CHECK-SAME: "-o" "[[IMG_DEV_B_803:.*out]]" [[OBJ_DEV_B_803]] // @@ -141,7 +141,7 @@ // CHECK-SAME: "-filetype=obj" // CHECK-SAME: "-o" [[OBJ_DEV_B_900:".*-gfx900-.*o"]] -// CHECK: [[LLD]] "-flavor" "gnu" "-shared" +// CHECK: [[LLD]] "-flavor" "gnu" "--no-undefined" "-shared" // CHECK-SAME: "-o" "[[IMG_DEV_B_900:.*out]]" [[OBJ_DEV_B_900]] // diff --git a/clang/test/Driver/hip-toolchain-rdc.hip b/clang/test/Driver/hip-toolchain-rdc.hip index 203784f..0880d0f 100644 --- a/clang/test/Driver/hip-toolchain-rdc.hip +++ b/clang/test/Driver/hip-toolchain-rdc.hip @@ -44,7 +44,7 @@ // CHECK-SAME: "-filetype=obj" // CHECK-SAME: "-o" [[OBJ_DEV1:".*-gfx803-.*o"]] -// CHECK: [[LLD: ".*lld.*"]] "-flavor" "gnu" "-shared" +// CHECK: [[LLD: ".*lld.*"]] "-flavor" "gnu" "--no-undefined" "-shared" // CHECK-SAME: "-o" "[[IMG_DEV1:.*out]]" [[OBJ_DEV1]] // CHECK: [[CLANG]] "-cc1" "-triple" "amdgcn-amd-amdhsa" @@ -77,7 +77,7 @@ // CHECK-SAME: "-filetype=obj" // CHECK-SAME: "-o" [[OBJ_DEV2:".*-gfx900-.*o"]] -// CHECK: [[LLD]] "-flavor" "gnu" "-shared" +// CHECK: [[LLD]] "-flavor" "gnu" "--no-undefined" "-shared" // CHECK-SAME: "-o" "[[IMG_DEV2:.*out]]" [[OBJ_DEV2]] // CHECK: [[CLANG]] "-cc1" "-triple" "x86_64-unknown-linux-gnu" diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test index d3705cf..e475070 100644 --- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test +++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test @@ -61,7 +61,6 @@ // 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/clang/test/SemaCUDA/hip-pinned-shadow.cu b/clang/test/SemaCUDA/hip-pinned-shadow.cu deleted file mode 100644 index c58f709..0000000 --- a/clang/test/SemaCUDA/hip-pinned-shadow.cu +++ /dev/null @@ -1,25 +0,0 @@ -// 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}}