Index: cfe/trunk/include/clang/Basic/Attr.td =================================================================== --- cfe/trunk/include/clang/Basic/Attr.td +++ cfe/trunk/include/clang/Basic/Attr.td @@ -295,6 +295,7 @@ 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 @@ 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]; Index: cfe/trunk/include/clang/Basic/AttrDocs.td =================================================================== --- cfe/trunk/include/clang/Basic/AttrDocs.td +++ cfe/trunk/include/clang/Basic/AttrDocs.td @@ -4183,3 +4183,15 @@ ``__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 Index: cfe/trunk/lib/CodeGen/CodeGenModule.cpp =================================================================== --- cfe/trunk/lib/CodeGen/CodeGenModule.cpp +++ cfe/trunk/lib/CodeGen/CodeGenModule.cpp @@ -2415,7 +2415,8 @@ if (!Global->hasAttr() && !Global->hasAttr() && !Global->hasAttr() && - !Global->hasAttr()) + !Global->hasAttr() && + !(LangOpts.HIP && Global->hasAttr())) return; } else { // We need to emit host-side 'shadows' for all global @@ -3781,7 +3782,12 @@ !getLangOpts().CUDAIsDevice && (D->hasAttr() || D->hasAttr() || D->hasAttr()); - if (getLangOpts().CUDA && (IsCUDASharedVar || IsCUDAShadowVar)) + // HIP pinned shadow of initialized host-side global variables are also + // left undefined. + bool IsHIPPinnedShadowVar = + getLangOpts().CUDAIsDevice && D->hasAttr(); + 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 @@ // 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()) { + if (D->hasAttr() || D->hasAttr() || + D->hasAttr()) { Linkage = llvm::GlobalValue::InternalLinkage; // Shadow variables and their properties must be registered @@ -3916,7 +3923,8 @@ } } - 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. Index: cfe/trunk/lib/CodeGen/TargetInfo.cpp =================================================================== --- cfe/trunk/lib/CodeGen/TargetInfo.cpp +++ cfe/trunk/lib/CodeGen/TargetInfo.cpp @@ -7874,12 +7874,24 @@ return D->hasAttr() || (isa(D) && D->hasAttr()) || (isa(D) && - (D->hasAttr() || D->hasAttr())); + (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(); } 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); } Index: cfe/trunk/lib/Driver/ToolChains/HIP.cpp =================================================================== --- cfe/trunk/lib/Driver/ToolChains/HIP.cpp +++ cfe/trunk/lib/Driver/ToolChains/HIP.cpp @@ -170,9 +170,8 @@ 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); Index: cfe/trunk/lib/Sema/SemaDeclAttr.cpp =================================================================== --- cfe/trunk/lib/Sema/SemaDeclAttr.cpp +++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp @@ -6786,6 +6786,10 @@ case ParsedAttr::AT_CUDAHost: handleSimpleAttributeWithExclusions(S, D, AL); break; + case ParsedAttr::AT_HIPPinnedShadow: + handleSimpleAttributeWithExclusions(S, D, AL); + break; case ParsedAttr::AT_GNUInline: handleGNUInlineAttr(S, D, AL); break; Index: cfe/trunk/test/AST/ast-dump-hip-pinned-shadow.cu =================================================================== --- cfe/trunk/test/AST/ast-dump-hip-pinned-shadow.cu +++ cfe/trunk/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 +struct texture : public textureReference { +texture() { a = 1; } +}; + +__attribute__((hip_pinned_shadow)) texture tex; Index: cfe/trunk/test/CodeGenCUDA/hip-pinned-shadow.cu =================================================================== --- cfe/trunk/test/CodeGenCUDA/hip-pinned-shadow.cu +++ cfe/trunk/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 +struct texture : public textureReference { +texture() { a = 1; } +}; + +__attribute__((hip_pinned_shadow)) texture 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) Index: cfe/trunk/test/Driver/hip-toolchain-no-rdc.hip =================================================================== --- cfe/trunk/test/Driver/hip-toolchain-no-rdc.hip +++ cfe/trunk/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]] // Index: cfe/trunk/test/Driver/hip-toolchain-rdc.hip =================================================================== --- cfe/trunk/test/Driver/hip-toolchain-rdc.hip +++ cfe/trunk/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" Index: cfe/trunk/test/Misc/pragma-attribute-supported-attributes-list.test =================================================================== --- cfe/trunk/test/Misc/pragma-attribute-supported-attributes-list.test +++ cfe/trunk/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) Index: cfe/trunk/test/SemaCUDA/hip-pinned-shadow.cu =================================================================== --- cfe/trunk/test/SemaCUDA/hip-pinned-shadow.cu +++ cfe/trunk/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 +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}}