[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

diff --git a/include/clang/Basic/Attr.td b/include/clang/Basic/Attr.td
index 5dd9cfc95827cc2df6664b4e366b789d3b314b6b..93913b043ec8ac1204fdfda8f88319d4519df00b 100644 (file)
--- a/include/clang/Basic/Attr.td
+++ b/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/include/clang/Basic/AttrDocs.td b/include/clang/Basic/AttrDocs.td
index 97e50f9236032b4f4788ec8dc36870f79e90abb6..fac6116057dc2103c7dc90105ef51bbbd005b1c6 100644 (file)
--- 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 87a1d45abf46e8be542dcbb29cbf00ceaeb08383..b0e3b0bb98de80784130c5c00c9526568715e305 100644 (file)
--- a/lib/CodeGen/CodeGenModule.cpp
+++ b/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/lib/CodeGen/TargetInfo.cpp b/lib/CodeGen/TargetInfo.cpp
index 8563a18b2ecd944c5b542a8bded155fb0b3f5e98..22f70a08b0f93d52c59f2759c6af36ee8d6911f9 100644 (file)
--- a/lib/CodeGen/TargetInfo.cpp
+++ b/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/lib/Driver/ToolChains/HIP.cpp b/lib/Driver/ToolChains/HIP.cpp
index a60485ab03b7e91d9ee9e09874e720cde6458fcf..2ec97e798fd0509493269d5e047df2adb8808b39 100644 (file)
--- 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 8670f74d9560154e668a814d6c8a338aaea3e93c..5a1712ff13a5d60211908ba8b4497e629fbb26f3 100644 (file)
--- 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<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/test/AST/ast-dump-hip-pinned-shadow.cu b/test/AST/ast-dump-hip-pinned-shadow.cu
new file mode 100644 (file)
index 0000000..53d7c8f
--- /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 <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/test/CodeGenCUDA/hip-pinned-shadow.cu b/test/CodeGenCUDA/hip-pinned-shadow.cu
new file mode 100644 (file)
index 0000000..75798f7
--- /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 <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/test/Driver/hip-toolchain-no-rdc.hip b/test/Driver/hip-toolchain-no-rdc.hip
index 229839db6c85db3371b861d57c6695e8cfd02059..540b9328605388e22ed2541a9e36d88607d2f4ec 100644 (file)
--- 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 055efe69adb746e7e3798e79dc9140725750abd5..15ac5f19312b23d041fb6ea4885465c40ad8790a 100644 (file)
--- 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 aa71e0bc8f58eef6ff8c11c4e6ea20f4ba58b63a..fc9d86efe705b4d043d01a8947f35f2ce45dd7b4 100644 (file)
--- 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 (file)
index 0000000..c58f709
--- /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 <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}}