summaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
Diffstat
-rw-r--r--clang/include/clang/Basic/Attr.td8
-rw-r--r--clang/include/clang/Basic/AttrDocs.td12
-rw-r--r--clang/lib/CodeGen/CodeGenModule.cpp16
-rw-r--r--clang/lib/CodeGen/TargetInfo.cpp16
-rw-r--r--clang/lib/Driver/ToolChains/HIP.cpp5
-rw-r--r--clang/lib/Sema/SemaDeclAttr.cpp4
-rw-r--r--clang/test/AST/ast-dump-hip-pinned-shadow.cu13
-rw-r--r--clang/test/CodeGenCUDA/hip-pinned-shadow.cu23
-rw-r--r--clang/test/Driver/hip-toolchain-no-rdc.hip8
-rw-r--r--clang/test/Driver/hip-toolchain-rdc.hip4
-rw-r--r--clang/test/Misc/pragma-attribute-supported-attributes-list.test1
-rw-r--r--clang/test/SemaCUDA/hip-pinned-shadow.cu25
12 files changed, 120 insertions, 15 deletions
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 5dd9cfc9582..93913b043ec 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -295,6 +295,7 @@ class LangOpt<string name, code customCode = [{}]> {
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/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index 97e50f92360..fac6116057d 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/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/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index 87a1d45abf4..b0e3b0bb98d 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -2415,7 +2415,8 @@ void CodeGenModule::EmitGlobal(GlobalDecl GD) {
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::EmitGlobalVarDefinition(const VarDecl *D,
!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::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<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::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/clang/lib/CodeGen/TargetInfo.cpp b/clang/lib/CodeGen/TargetInfo.cpp
index 8563a18b2ec..22f70a08b0f 100644
--- a/clang/lib/CodeGen/TargetInfo.cpp
+++ b/clang/lib/CodeGen/TargetInfo.cpp
@@ -7874,12 +7874,24 @@ static bool requiresAMDGPUProtectedVisibility(const Decl *D,
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);
}
diff --git a/clang/lib/Driver/ToolChains/HIP.cpp b/clang/lib/Driver/ToolChains/HIP.cpp
index a60485ab03b..2ec97e798fd 100644
--- a/clang/lib/Driver/ToolChains/HIP.cpp
+++ b/clang/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/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 8670f74d956..5a1712ff13a 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -6786,6 +6786,10 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
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;
diff --git a/clang/test/AST/ast-dump-hip-pinned-shadow.cu b/clang/test/AST/ast-dump-hip-pinned-shadow.cu
new file mode 100644
index 00000000000..53d7c8fbed8
--- /dev/null
+++ b/clang/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 <class T, int texType, int hipTextureReadMode>
+struct texture : public textureReference {
+texture() { a = 1; }
+};
+
+__attribute__((hip_pinned_shadow)) texture<float, 1, 1> tex;
diff --git a/clang/test/CodeGenCUDA/hip-pinned-shadow.cu b/clang/test/CodeGenCUDA/hip-pinned-shadow.cu
new file mode 100644
index 00000000000..75798f7e1de
--- /dev/null
+++ b/clang/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 <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)
diff --git a/clang/test/Driver/hip-toolchain-no-rdc.hip b/clang/test/Driver/hip-toolchain-no-rdc.hip
index 229839db6c8..540b9328605 100644
--- a/clang/test/Driver/hip-toolchain-no-rdc.hip
+++ b/clang/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/clang/test/Driver/hip-toolchain-rdc.hip b/clang/test/Driver/hip-toolchain-rdc.hip
index 055efe69adb..15ac5f19312 100644
--- a/clang/test/Driver/hip-toolchain-rdc.hip
+++ b/clang/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/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
index aa71e0bc8f5..fc9d86efe70 100644
--- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test
+++ b/clang/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/clang/test/SemaCUDA/hip-pinned-shadow.cu b/clang/test/SemaCUDA/hip-pinned-shadow.cu
new file mode 100644
index 00000000000..c58f7097af6
--- /dev/null
+++ b/clang/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 <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}}
OpenPOWER on IntegriCloud