This revision was automatically updated to reflect the committed changes.
Closed by commit rGc97be2c37785: [hip] Remove `hip_pinned_shadow`. (authored by
hliao).
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D77583/new/
https://reviews.llvm.org/D77583
Files:
clang/include/clang/Basic/Attr.td
clang/include/clang/Basic/AttrDocs.td
clang/lib/CodeGen/CodeGenModule.cpp
clang/lib/CodeGen/CodeGenModule.h
clang/lib/CodeGen/TargetInfo.cpp
clang/lib/Driver/ToolChains/HIP.cpp
clang/lib/Sema/SemaDeclAttr.cpp
clang/test/AST/ast-dump-hip-pinned-shadow.cu
clang/test/Driver/hip-toolchain-no-rdc.hip
clang/test/Driver/hip-toolchain-rdc.hip
clang/test/Misc/pragma-attribute-supported-attributes-list.test
clang/test/SemaCUDA/hip-pinned-shadow.cu
Index: clang/test/SemaCUDA/hip-pinned-shadow.cu
===================================================================
--- 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 <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@-1{{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables}}
- // expected-note@-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@-1{{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables}}
- // expected-note@-2{{conflicting attribute is here}}
Index: clang/test/Misc/pragma-attribute-supported-attributes-list.test
===================================================================
--- clang/test/Misc/pragma-attribute-supported-attributes-list.test
+++ 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)
Index: clang/test/Driver/hip-toolchain-rdc.hip
===================================================================
--- clang/test/Driver/hip-toolchain-rdc.hip
+++ 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"
Index: clang/test/Driver/hip-toolchain-no-rdc.hip
===================================================================
--- clang/test/Driver/hip-toolchain-no-rdc.hip
+++ 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]]
//
Index: clang/test/AST/ast-dump-hip-pinned-shadow.cu
===================================================================
--- 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 <class T, int texType, int hipTextureReadMode>
-struct texture : public textureReference {
-texture() { a = 1; }
-};
-
-__attribute__((hip_pinned_shadow)) texture<float, 1, 1> tex;
Index: clang/lib/Sema/SemaDeclAttr.cpp
===================================================================
--- clang/lib/Sema/SemaDeclAttr.cpp
+++ clang/lib/Sema/SemaDeclAttr.cpp
@@ -6930,10 +6930,6 @@
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_CUDADeviceBuiltinSurfaceType:
handleSimpleAttributeWithExclusions<CUDADeviceBuiltinSurfaceTypeAttr,
CUDADeviceBuiltinTextureTypeAttr>(S, D,
Index: clang/lib/Driver/ToolChains/HIP.cpp
===================================================================
--- clang/lib/Driver/ToolChains/HIP.cpp
+++ 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<Command>(JA, *this, Lld, LldArgs, Inputs));
}
Index: clang/lib/CodeGen/TargetInfo.cpp
===================================================================
--- clang/lib/CodeGen/TargetInfo.cpp
+++ clang/lib/CodeGen/TargetInfo.cpp
@@ -8407,23 +8407,13 @@
(isa<FunctionDecl>(D) && D->hasAttr<CUDAGlobalAttr>()) ||
(isa<VarDecl>(D) &&
(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>();
+ cast<VarDecl>(D)->getType()->isCUDADeviceBuiltinSurfaceType() ||
+ cast<VarDecl>(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);
}
Index: clang/lib/CodeGen/CodeGenModule.h
===================================================================
--- clang/lib/CodeGen/CodeGenModule.h
+++ 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);
Index: clang/lib/CodeGen/CodeGenModule.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenModule.cpp
+++ 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<CUDAGlobalAttr>() &&
!Global->hasAttr<CUDAConstantAttr>() &&
!Global->hasAttr<CUDASharedAttr>() &&
- !(LangOpts.HIP && Global->hasAttr<HIPPinnedShadowAttr>()) &&
!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<HIPPinnedShadowAttr>();
- if (getLangOpts().CUDA && (IsCUDASharedVar || IsCUDAShadowVar ||
- IsCUDADeviceShadowVar || IsHIPPinnedShadowVar))
+ if (getLangOpts().CUDA &&
+ (IsCUDASharedVar || IsCUDAShadowVar || IsCUDADeviceShadowVar))
Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy));
else if (D->hasAttr<LoaderUninitializedAttr>())
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<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>() ||
- D->hasAttr<HIPPinnedShadowAttr>()) {
+ if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()) {
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);
Index: clang/include/clang/Basic/AttrDocs.td
===================================================================
--- clang/include/clang/Basic/AttrDocs.td
+++ 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 = [{
Index: clang/include/clang/Basic/Attr.td
===================================================================
--- clang/include/clang/Basic/Attr.td
+++ 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];
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
