diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -322,7 +322,6 @@ 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 @@ 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 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -4613,18 +4613,6 @@ }]; } -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.h b/clang/lib/CodeGen/CodeGenModule.h --- a/clang/lib/CodeGen/CodeGenModule.h +++ b/clang/lib/CodeGen/CodeGenModule.h @@ -1037,7 +1037,7 @@ 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/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1955,9 +1955,9 @@ } } -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 @@ !Global->hasAttr() && !Global->hasAttr() && !Global->hasAttr() && - !(LangOpts.HIP && Global->hasAttr()) && !Global->getType()->isCUDADeviceBuiltinSurfaceType() && !Global->getType()->isCUDADeviceBuiltinTextureType()) return; @@ -3928,10 +3927,8 @@ 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 @@ // 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 @@ } } - // 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/TargetInfo.cpp b/clang/lib/CodeGen/TargetInfo.cpp --- a/clang/lib/CodeGen/TargetInfo.cpp +++ b/clang/lib/CodeGen/TargetInfo.cpp @@ -8407,23 +8407,13 @@ (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 --- a/clang/lib/Driver/ToolChains/HIP.cpp +++ b/clang/lib/Driver/ToolChains/HIP.cpp @@ -192,8 +192,9 @@ 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 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -6930,10 +6930,6 @@ 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 --- 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 --- 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 --- 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 --- 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 --- 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}}