Index: clang/include/clang/Basic/LangOptions.def =================================================================== --- clang/include/clang/Basic/LangOptions.def +++ clang/include/clang/Basic/LangOptions.def @@ -232,6 +232,7 @@ LANGOPT(GPURelocatableDeviceCode, 1, 0, "generate relocatable device code") LANGOPT(GPUAllowDeviceInit, 1, 0, "allowing device side global init functions for HIP") LANGOPT(GPUMaxThreadsPerBlock, 32, 256, "default max threads per block for kernel launch bounds for HIP") +LANGOPT(HIPCUID, 32, 0, "default compilation unit id for HIP") LANGOPT(SYCL , 1, 0, "SYCL") LANGOPT(SYCLIsDevice , 1, 0, "Generate code for SYCL device") Index: clang/include/clang/Driver/Action.h =================================================================== --- clang/include/clang/Driver/Action.h +++ clang/include/clang/Driver/Action.h @@ -213,14 +213,17 @@ class InputAction : public Action { const llvm::opt::Arg &Input; - + unsigned Id; virtual void anchor(); public: - InputAction(const llvm::opt::Arg &Input, types::ID Type); + InputAction(const llvm::opt::Arg &Input, types::ID Type, unsigned Id = 0); const llvm::opt::Arg &getInputArg() const { return Input; } + void setId(unsigned _Id) { Id = _Id; } + unsigned getId() const { return Id; } + static bool classof(const Action *A) { return A->getKind() == InputClass; } Index: clang/include/clang/Driver/CC1Options.td =================================================================== --- clang/include/clang/Driver/CC1Options.td +++ clang/include/clang/Driver/CC1Options.td @@ -891,6 +891,14 @@ HelpText<"Don't treat unattributed constexpr functions as __host__ __device__.">; //===----------------------------------------------------------------------===// +// HIP Options +//===----------------------------------------------------------------------===// + +def hip_cuid_EQ : Joined<["-"], "hip-cuid=">, + HelpText<"A unique id to identify a HIP compilation unit, which can be used to " + "identify a device-side static declaration in host compilation">; + +//===----------------------------------------------------------------------===// // OpenMP Options //===----------------------------------------------------------------------===// Index: clang/lib/AST/ASTContext.cpp =================================================================== --- clang/lib/AST/ASTContext.cpp +++ clang/lib/AST/ASTContext.cpp @@ -10055,12 +10055,19 @@ } else if (D->hasAttr()) { if (L == GVA_DiscardableODR) return GVA_StrongODR; - } else if (Context.getLangOpts().CUDA && Context.getLangOpts().CUDAIsDevice && - D->hasAttr()) { + } else if (Context.getLangOpts().CUDA && Context.getLangOpts().CUDAIsDevice) { // Device-side functions with __global__ attribute must always be // visible externally so they can be launched from host. - if (L == GVA_DiscardableODR || L == GVA_Internal) + if (D->hasAttr() && + (L == GVA_DiscardableODR || L == GVA_Internal)) return GVA_StrongODR; + // Externalize device side static file-scope variable for HIP. + if (Context.getLangOpts().HIP && Context.getLangOpts().HIPCUID && + (D->hasAttr() || D->hasAttr()) && + isa(D) && cast(D)->isFileVarDecl() && + cast(D)->getStorageClass() == SC_Static) { + return GVA_StrongExternal; + } } return L; } Index: clang/lib/CodeGen/CGCUDANV.cpp =================================================================== --- clang/lib/CodeGen/CGCUDANV.cpp +++ clang/lib/CodeGen/CGCUDANV.cpp @@ -238,6 +238,19 @@ DeviceSideName = std::string(Out.str()); } else DeviceSideName = std::string(ND->getIdentifier()->getName()); + + // Make unique name for device side static file-scope variable for HIP. + if (CGM.getLangOpts().HIP && CGM.getLangOpts().HIPCUID) { + if (const auto *VD = dyn_cast(ND)) { + if ((VD->hasAttr() || VD->hasAttr()) && + VD->isFileVarDecl() && VD->getStorageClass() == SC_Static) { + SmallString<256> Buffer; + llvm::raw_svector_ostream Out(Buffer); + Out << DeviceSideName << ".hip.static." << CGM.getLangOpts().HIPCUID; + DeviceSideName = std::string(Out.str()); + } + } + } return DeviceSideName; } Index: clang/lib/CodeGen/CodeGenModule.cpp =================================================================== --- clang/lib/CodeGen/CodeGenModule.cpp +++ clang/lib/CodeGen/CodeGenModule.cpp @@ -1087,6 +1087,15 @@ } } + // Make unique name for device side static file-scope variable for HIP. + if (CGM.getLangOpts().HIP && CGM.getLangOpts().HIPCUID) { + if (const auto *VD = dyn_cast(ND)) { + if ((VD->hasAttr() || VD->hasAttr()) && + VD->isFileVarDecl() && VD->getStorageClass() == SC_Static) { + Out << ".hip.static." << CGM.getLangOpts().HIPCUID; + } + } + } return std::string(Out.str()); } Index: clang/lib/Driver/Action.cpp =================================================================== --- clang/lib/Driver/Action.cpp +++ clang/lib/Driver/Action.cpp @@ -8,6 +8,7 @@ #include "clang/Driver/Action.h" #include "llvm/Support/ErrorHandling.h" +#include "llvm/Support/Process.h" #include #include @@ -163,8 +164,11 @@ void InputAction::anchor() {} -InputAction::InputAction(const Arg &_Input, types::ID _Type) - : Action(InputClass, _Type), Input(_Input) {} +InputAction::InputAction(const Arg &_Input, types::ID _Type, unsigned _Id) + : Action(InputClass, _Type), Input(_Input), Id(_Id) { + if (!Id) + Id = llvm::sys::Process::GetRandomNumber(); +} void BindArchAction::anchor() {} Index: clang/lib/Driver/Driver.cpp =================================================================== --- clang/lib/Driver/Driver.cpp +++ clang/lib/Driver/Driver.cpp @@ -2410,7 +2410,7 @@ : types::TY_CUDA_DEVICE; for (unsigned I = 0, E = GpuArchList.size(); I != E; ++I) { CudaDeviceActions.push_back( - C.MakeAction(IA->getInputArg(), Ty)); + C.MakeAction(IA->getInputArg(), Ty, IA->getId())); } return ABRT_Success; Index: clang/lib/Driver/ToolChains/Clang.cpp =================================================================== --- clang/lib/Driver/ToolChains/Clang.cpp +++ clang/lib/Driver/ToolChains/Clang.cpp @@ -6004,8 +6004,17 @@ CmdArgs.push_back("-fcuda-short-ptr"); } - if (IsHIP) + if (IsHIP) { + // Determine the original source input. + const Action *SourceAction = &JA; + while (SourceAction->getKind() != Action::InputClass) { + assert(!SourceAction->getInputs().empty() && "unexpected root action!"); + SourceAction = SourceAction->getInputs()[0]; + } + CmdArgs.push_back(Args.MakeArgString( + Twine("-hip-cuid=") + Twine(cast(SourceAction)->getId()))); CmdArgs.push_back("-fcuda-allow-variadic-functions"); + } // OpenMP offloading device jobs take the argument -fopenmp-host-ir-file-path // to specify the result of the compile phase on the host, so the meaningful Index: clang/lib/Frontend/CompilerInvocation.cpp =================================================================== --- clang/lib/Frontend/CompilerInvocation.cpp +++ clang/lib/Frontend/CompilerInvocation.cpp @@ -2579,12 +2579,16 @@ } Opts.HIPUseNewLaunchAPI = Args.hasArg(OPT_fhip_new_launch_api); Opts.HIPLambdaHostDevice = Args.hasArg(OPT_fhip_lambda_host_device); - if (Opts.HIP) + if (Opts.HIP) { Opts.GPUMaxThreadsPerBlock = getLastArgIntValue( Args, OPT_gpu_max_threads_per_block_EQ, Opts.GPUMaxThreadsPerBlock); - else if (Args.hasArg(OPT_gpu_max_threads_per_block_EQ)) - Diags.Report(diag::warn_ignored_hip_only_option) - << Args.getLastArg(OPT_gpu_max_threads_per_block_EQ)->getAsString(Args); + Opts.HIPCUID = getLastArgIntValue(Args, OPT_hip_cuid_EQ, Opts.HIPCUID); + } else { + for (auto Opt : {OPT_gpu_max_threads_per_block_EQ, OPT_hip_cuid_EQ}) + if (Args.hasArg(Opt)) + Diags.Report(diag::warn_ignored_hip_only_option) + << Args.getLastArg(Opt)->getAsString(Args); + } if (Opts.ObjC) { if (Arg *arg = Args.getLastArg(OPT_fobjc_runtime_EQ)) { Index: clang/test/CodeGenCUDA/static-device-var.cu =================================================================== --- /dev/null +++ clang/test/CodeGenCUDA/static-device-var.cu @@ -0,0 +1,57 @@ +// REQUIRES: x86-registered-target +// REQUIRES: amdgpu-registered-target + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \ +// RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefix=INT-DEV %s + +// RUN: %clang_cc1 -triple x86_64-gnu-linux \ +// RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefix=INT-HOST %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -hip-cuid=123 \ +// RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefix=EXT-DEV %s + +// RUN: %clang_cc1 -triple x86_64-gnu-linux -hip-cuid=123 \ +// RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefix=EXT-HOST %s + +#include "Inputs/cuda.h" + +// Test normal static device variables +// INT-DEV: @_ZL1x = internal addrspace(1) global i32 0 +// INT-HOST-DAG: @_ZL1x = internal global i32 undef +// INT-HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x\00" + +// Test externalized static device variables +// EXT-DEV: @_ZL1x.hip.static.123 = addrspace(1) externally_initialized global i32 0 +// EXT-HOST-DAG: @_ZL1x.hip.static.123 = internal global i32 undef +// EXT-HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x.hip.static.123\00" + +static __device__ int x; + +// Test normal static device variables +// INT-DEV: @_ZL1y = internal addrspace(4) global i32 0 +// INT-HOST-DAG: @_ZL1y = internal global i32 undef +// INT-HOST-DAG: @[[DEVNAMEY:[0-9]+]] = {{.*}}c"_ZL1y\00" + +// Test externalized static device variables +// EXT-DEV: @_ZL1y.hip.static.123 = addrspace(4) externally_initialized global i32 0 +// EXT-HOST-DAG: @_ZL1y.hip.static.123 = internal global i32 undef +// EXT-HOST-DAG: @[[DEVNAMEY:[0-9]+]] = {{.*}}c"_ZL1y.hip.static.123\00" + +static __constant__ int y; + +__global__ void kernel(int *a) { + a[0] = x; + a[1] = y; +} + +int* getDeviceSymbol(int *x); + +void foo() { + getDeviceSymbol(&x); + getDeviceSymbol(&y); +} + +// INT-HOST: __hipRegisterVar({{.*}}@_ZL1x{{.*}}@[[DEVNAMEX]] +// INT-HOST: __hipRegisterVar({{.*}}@_ZL1y{{.*}}@[[DEVNAMEY]] +// EXT-HOST: __hipRegisterVar({{.*}}@_ZL1x.hip.static.123{{.*}}@[[DEVNAMEX]] +// EXT-HOST: __hipRegisterVar({{.*}}@_ZL1y.hip.static.123{{.*}}@[[DEVNAMEY]] Index: clang/test/Driver/hip-cuid.hip =================================================================== --- /dev/null +++ clang/test/Driver/hip-cuid.hip @@ -0,0 +1,20 @@ +// REQUIRES: clang-driver +// REQUIRES: x86-registered-target +// REQUIRES: amdgpu-registered-target + +// RUN: %clang -### -x hip \ +// RUN: -target x86_64-unknown-linux-gnu \ +// RUN: --offload-arch=gfx900 \ +// RUN: --offload-arch=gfx906 \ +// RUN: -c -nogpulib \ +// RUN: %S/Inputs/hip_multiple_inputs/b.hip \ +// RUN: 2>&1 | FileCheck %s + +// CHECK: "{{.*}}clang{{.*}}" "-cc1" "-triple" "amdgcn-amd-amdhsa" +// CHECK-SAME: "-hip-cuid=[[CUID:[^"]+]]" + +// CHECK: "{{.*}}clang{{.*}}" "-cc1" "-triple" "amdgcn-amd-amdhsa" +// CHECK-SAME: "-hip-cuid=[[CUID]]" + +// CHECK: "{{.*}}clang{{.*}}" "-cc1" "-triple" "x86_64-unknown-linux-gnu" +// CHECK-SAME: "-hip-cuid=[[CUID]]"