forked from OSchip/llvm-project
[CUDA] Create offloading entries when using the new driver
The changes made in D123460 generalized the code generation for OpenMP's offloading entries. We can use the same scheme to register globals for CUDA code. This patch adds the code generation to create these offloading entries when compiling using the new offloading driver mode. The offloading entries are simple structs that contain the information necessary to register the global. The struct used is as follows: ``` Type struct __tgt_offload_entry { void *addr; // Pointer to the offload entry info. // (function or global) char *name; // Name of the function or global. size_t size; // Size of the entry info (0 if it a function). int32_t flags; int32_t reserved; }; ``` Currently CUDA handles RDC code generation by deferring the registration of globals in the current TU to a callback function containing the modules ID. Later all the module IDs will be used to register all of the globals at once. Rather than mimic this, offloading entries allow us to mimic the way OpenMP registers globals. That is, we create a simple global struct for each device global to be registered. These are placed at a special section `cuda_offloading_entires`. Because this section is a valid C-identifier, the linker will profide a `__start` and `__stop` pointer that we can use to iterate and register all globals at runtime. the registration requires a flag variable to indicate which registration function to use. I have assigned the flags somewhat arbitrarily, but these use the following values. Kernel: 0 Variable: 0 Managed: 1 Surface: 2 Texture: 3 Depends on D120272 Reviewed By: tra Differential Revision: https://reviews.llvm.org/D123471
This commit is contained in:
parent
c7ba568f40
commit
0035f7154c
|
@ -266,6 +266,7 @@ LANGOPT(GPUAllowDeviceInit, 1, 0, "allowing device side global init functions fo
|
||||||
LANGOPT(GPUMaxThreadsPerBlock, 32, 1024, "default max threads per block for kernel launch bounds for HIP")
|
LANGOPT(GPUMaxThreadsPerBlock, 32, 1024, "default max threads per block for kernel launch bounds for HIP")
|
||||||
LANGOPT(GPUDeferDiag, 1, 0, "defer host/device related diagnostic messages for CUDA/HIP")
|
LANGOPT(GPUDeferDiag, 1, 0, "defer host/device related diagnostic messages for CUDA/HIP")
|
||||||
LANGOPT(GPUExcludeWrongSideOverloads, 1, 0, "always exclude wrong side overloads in overloading resolution for CUDA/HIP")
|
LANGOPT(GPUExcludeWrongSideOverloads, 1, 0, "always exclude wrong side overloads in overloading resolution for CUDA/HIP")
|
||||||
|
LANGOPT(OffloadingNewDriver, 1, 0, "use the new driver for generating offloading code.")
|
||||||
|
|
||||||
LANGOPT(SYCLIsDevice , 1, 0, "Generate code for SYCL device")
|
LANGOPT(SYCLIsDevice , 1, 0, "Generate code for SYCL device")
|
||||||
LANGOPT(SYCLIsHost , 1, 0, "SYCL host compilation")
|
LANGOPT(SYCLIsHost , 1, 0, "SYCL host compilation")
|
||||||
|
|
|
@ -2526,9 +2526,9 @@ defm openmp_optimistic_collapse : BoolFOption<"openmp-optimistic-collapse",
|
||||||
PosFlag<SetTrue, [CC1Option]>, NegFlag<SetFalse>, BothFlags<[NoArgumentUnused, HelpHidden]>>;
|
PosFlag<SetTrue, [CC1Option]>, NegFlag<SetFalse>, BothFlags<[NoArgumentUnused, HelpHidden]>>;
|
||||||
def static_openmp: Flag<["-"], "static-openmp">,
|
def static_openmp: Flag<["-"], "static-openmp">,
|
||||||
HelpText<"Use the static host OpenMP runtime while linking.">;
|
HelpText<"Use the static host OpenMP runtime while linking.">;
|
||||||
def offload_new_driver : Flag<["--"], "offload-new-driver">, Flags<[CC1Option]>, Group<Action_Group>,
|
def offload_new_driver : Flag<["--"], "offload-new-driver">, Flags<[CC1Option]>, Group<f_Group>,
|
||||||
HelpText<"Use the new driver for offloading compilation.">;
|
MarshallingInfoFlag<LangOpts<"OffloadingNewDriver">>, HelpText<"Use the new driver for offloading compilation.">;
|
||||||
def no_offload_new_driver : Flag<["--"], "no-offload-new-driver">, Flags<[CC1Option]>, Group<Action_Group>,
|
def no_offload_new_driver : Flag<["--"], "no-offload-new-driver">, Flags<[CC1Option]>, Group<f_Group>,
|
||||||
HelpText<"Don't Use the new driver for offloading compilation.">;
|
HelpText<"Don't Use the new driver for offloading compilation.">;
|
||||||
def offload_device_only : Flag<["--"], "offload-device-only">,
|
def offload_device_only : Flag<["--"], "offload-device-only">,
|
||||||
HelpText<"Only compile for the offloading device.">;
|
HelpText<"Only compile for the offloading device.">;
|
||||||
|
@ -2543,7 +2543,7 @@ def cuda_host_only : Flag<["--"], "cuda-host-only">, Alias<offload_host_only>,
|
||||||
def cuda_compile_host_device : Flag<["--"], "cuda-compile-host-device">, Alias<offload_host_device>,
|
def cuda_compile_host_device : Flag<["--"], "cuda-compile-host-device">, Alias<offload_host_device>,
|
||||||
HelpText<"Compile CUDA code for both host and device (default). Has no "
|
HelpText<"Compile CUDA code for both host and device (default). Has no "
|
||||||
"effect on non-CUDA compilations.">;
|
"effect on non-CUDA compilations.">;
|
||||||
def fopenmp_new_driver : Flag<["-"], "fopenmp-new-driver">, Flags<[CC1Option]>, Group<Action_Group>,
|
def fopenmp_new_driver : Flag<["-"], "fopenmp-new-driver">, Flags<[CC1Option]>, Group<f_Group>,
|
||||||
HelpText<"Use the new driver for OpenMP offloading.">;
|
HelpText<"Use the new driver for OpenMP offloading.">;
|
||||||
def fno_openmp_new_driver : Flag<["-"], "fno-openmp-new-driver">, Flags<[CC1Option]>, Group<Action_Group>,
|
def fno_openmp_new_driver : Flag<["-"], "fno-openmp-new-driver">, Flags<[CC1Option]>, Group<Action_Group>,
|
||||||
Alias<no_offload_new_driver>, HelpText<"Don't use the new driver for OpenMP offloading.">;
|
Alias<no_offload_new_driver>, HelpText<"Don't use the new driver for OpenMP offloading.">;
|
||||||
|
|
|
@ -157,6 +157,8 @@ private:
|
||||||
llvm::Function *makeModuleDtorFunction();
|
llvm::Function *makeModuleDtorFunction();
|
||||||
/// Transform managed variables for device compilation.
|
/// Transform managed variables for device compilation.
|
||||||
void transformManagedVars();
|
void transformManagedVars();
|
||||||
|
/// Create offloading entries to register globals in RDC mode.
|
||||||
|
void createOffloadingEntries();
|
||||||
|
|
||||||
public:
|
public:
|
||||||
CGNVCUDARuntime(CodeGenModule &CGM);
|
CGNVCUDARuntime(CodeGenModule &CGM);
|
||||||
|
@ -210,7 +212,8 @@ static std::unique_ptr<MangleContext> InitDeviceMC(CodeGenModule &CGM) {
|
||||||
CGNVCUDARuntime::CGNVCUDARuntime(CodeGenModule &CGM)
|
CGNVCUDARuntime::CGNVCUDARuntime(CodeGenModule &CGM)
|
||||||
: CGCUDARuntime(CGM), Context(CGM.getLLVMContext()),
|
: CGCUDARuntime(CGM), Context(CGM.getLLVMContext()),
|
||||||
TheModule(CGM.getModule()),
|
TheModule(CGM.getModule()),
|
||||||
RelocatableDeviceCode(CGM.getLangOpts().GPURelocatableDeviceCode),
|
RelocatableDeviceCode(CGM.getLangOpts().GPURelocatableDeviceCode ||
|
||||||
|
CGM.getLangOpts().OffloadingNewDriver),
|
||||||
DeviceMC(InitDeviceMC(CGM)) {
|
DeviceMC(InitDeviceMC(CGM)) {
|
||||||
CodeGen::CodeGenTypes &Types = CGM.getTypes();
|
CodeGen::CodeGenTypes &Types = CGM.getTypes();
|
||||||
ASTContext &Ctx = CGM.getContext();
|
ASTContext &Ctx = CGM.getContext();
|
||||||
|
@ -1107,6 +1110,40 @@ void CGNVCUDARuntime::transformManagedVars() {
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
// Creates offloading entries for all the kernels and globals that must be
|
||||||
|
// registered. The linker will provide a pointer to this section so we can
|
||||||
|
// register the symbols with the linked device image.
|
||||||
|
void CGNVCUDARuntime::createOffloadingEntries() {
|
||||||
|
llvm::OpenMPIRBuilder OMPBuilder(CGM.getModule());
|
||||||
|
OMPBuilder.initialize();
|
||||||
|
|
||||||
|
StringRef Section = "cuda_offloading_entries";
|
||||||
|
for (KernelInfo &I : EmittedKernels)
|
||||||
|
OMPBuilder.emitOffloadingEntry(KernelHandles[I.Kernel],
|
||||||
|
getDeviceSideName(cast<NamedDecl>(I.D)), 0,
|
||||||
|
DeviceVarFlags::OffloadGlobalEntry, Section);
|
||||||
|
|
||||||
|
for (VarInfo &I : DeviceVars) {
|
||||||
|
uint64_t VarSize =
|
||||||
|
CGM.getDataLayout().getTypeAllocSize(I.Var->getValueType());
|
||||||
|
if (I.Flags.getKind() == DeviceVarFlags::Variable) {
|
||||||
|
OMPBuilder.emitOffloadingEntry(
|
||||||
|
I.Var, getDeviceSideName(I.D), VarSize,
|
||||||
|
I.Flags.isManaged() ? DeviceVarFlags::OffloadGlobalManagedEntry
|
||||||
|
: DeviceVarFlags::OffloadGlobalEntry,
|
||||||
|
Section);
|
||||||
|
} else if (I.Flags.getKind() == DeviceVarFlags::Surface) {
|
||||||
|
OMPBuilder.emitOffloadingEntry(I.Var, getDeviceSideName(I.D), VarSize,
|
||||||
|
DeviceVarFlags::OffloadGlobalSurfaceEntry,
|
||||||
|
Section);
|
||||||
|
} else if (I.Flags.getKind() == DeviceVarFlags::Texture) {
|
||||||
|
OMPBuilder.emitOffloadingEntry(I.Var, getDeviceSideName(I.D), VarSize,
|
||||||
|
DeviceVarFlags::OffloadGlobalTextureEntry,
|
||||||
|
Section);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
// Returns module constructor to be added.
|
// Returns module constructor to be added.
|
||||||
llvm::Function *CGNVCUDARuntime::finalizeModule() {
|
llvm::Function *CGNVCUDARuntime::finalizeModule() {
|
||||||
if (CGM.getLangOpts().CUDAIsDevice) {
|
if (CGM.getLangOpts().CUDAIsDevice) {
|
||||||
|
@ -1135,7 +1172,11 @@ llvm::Function *CGNVCUDARuntime::finalizeModule() {
|
||||||
}
|
}
|
||||||
return nullptr;
|
return nullptr;
|
||||||
}
|
}
|
||||||
|
if (!(CGM.getLangOpts().OffloadingNewDriver && RelocatableDeviceCode))
|
||||||
return makeModuleCtorFunction();
|
return makeModuleCtorFunction();
|
||||||
|
|
||||||
|
createOffloadingEntries();
|
||||||
|
return nullptr;
|
||||||
}
|
}
|
||||||
|
|
||||||
llvm::GlobalValue *CGNVCUDARuntime::getKernelHandle(llvm::Function *F,
|
llvm::GlobalValue *CGNVCUDARuntime::getKernelHandle(llvm::Function *F,
|
||||||
|
|
|
@ -52,6 +52,19 @@ public:
|
||||||
Texture, // Builtin texture
|
Texture, // Builtin texture
|
||||||
};
|
};
|
||||||
|
|
||||||
|
/// The kind flag for an offloading entry.
|
||||||
|
enum OffloadEntryKindFlag : uint32_t {
|
||||||
|
/// Mark the entry as a global entry. This indicates the presense of a
|
||||||
|
/// kernel if the size size field is zero and a variable otherwise.
|
||||||
|
OffloadGlobalEntry = 0x0,
|
||||||
|
/// Mark the entry as a managed global variable.
|
||||||
|
OffloadGlobalManagedEntry = 0x1,
|
||||||
|
/// Mark the entry as a surface variable.
|
||||||
|
OffloadGlobalSurfaceEntry = 0x2,
|
||||||
|
/// Mark the entry as a texture variable.
|
||||||
|
OffloadGlobalTextureEntry = 0x3,
|
||||||
|
};
|
||||||
|
|
||||||
private:
|
private:
|
||||||
unsigned Kind : 2;
|
unsigned Kind : 2;
|
||||||
unsigned Extern : 1;
|
unsigned Extern : 1;
|
||||||
|
|
|
@ -6086,6 +6086,10 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
|
||||||
options::OPT_fno_openmp_extensions);
|
options::OPT_fno_openmp_extensions);
|
||||||
}
|
}
|
||||||
|
|
||||||
|
// Forward the new driver to change offloading code generation.
|
||||||
|
if (Args.hasArg(options::OPT_offload_new_driver))
|
||||||
|
CmdArgs.push_back("--offload-new-driver");
|
||||||
|
|
||||||
SanitizeArgs.addArgs(TC, Args, CmdArgs, InputType);
|
SanitizeArgs.addArgs(TC, Args, CmdArgs, InputType);
|
||||||
|
|
||||||
const XRayArgs &XRay = TC.getXRayArgs();
|
const XRayArgs &XRay = TC.getXRayArgs();
|
||||||
|
|
|
@ -0,0 +1,33 @@
|
||||||
|
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals
|
||||||
|
// RUN: %clang_cc1 -std=c++11 -triple x86_64-unknown-linux-gnu \
|
||||||
|
// RUN: --offload-new-driver -emit-llvm -o - -x cuda %s | FileCheck \
|
||||||
|
// RUN: --check-prefix=HOST %s
|
||||||
|
|
||||||
|
#include "Inputs/cuda.h"
|
||||||
|
|
||||||
|
//.
|
||||||
|
// HOST: @x = internal global i32 undef, align 4
|
||||||
|
// HOST: @.omp_offloading.entry_name = internal unnamed_addr constant [8 x i8] c"_Z3foov\00"
|
||||||
|
// HOST: @.omp_offloading.entry._Z3foov = weak constant %struct.__tgt_offload_entry { ptr @_Z18__device_stub__foov, ptr @.omp_offloading.entry_name, i64 0, i32 0, i32 0 }, section "cuda_offloading_entries", align 1
|
||||||
|
// HOST: @.omp_offloading.entry_name.1 = internal unnamed_addr constant [8 x i8] c"_Z3barv\00"
|
||||||
|
// HOST: @.omp_offloading.entry._Z3barv = weak constant %struct.__tgt_offload_entry { ptr @_Z18__device_stub__barv, ptr @.omp_offloading.entry_name.1, i64 0, i32 0, i32 0 }, section "cuda_offloading_entries", align 1
|
||||||
|
// HOST: @.omp_offloading.entry_name.2 = internal unnamed_addr constant [2 x i8] c"x\00"
|
||||||
|
// HOST: @.omp_offloading.entry.x = weak constant %struct.__tgt_offload_entry { ptr @x, ptr @.omp_offloading.entry_name.2, i64 4, i32 0, i32 0 }, section "cuda_offloading_entries", align 1
|
||||||
|
//.
|
||||||
|
// HOST-LABEL: @_Z18__device_stub__foov(
|
||||||
|
// HOST-NEXT: entry:
|
||||||
|
// HOST-NEXT: [[TMP0:%.*]] = call i32 @cudaLaunch(ptr @_Z18__device_stub__foov)
|
||||||
|
// HOST-NEXT: br label [[SETUP_END:%.*]]
|
||||||
|
// HOST: setup.end:
|
||||||
|
// HOST-NEXT: ret void
|
||||||
|
//
|
||||||
|
__global__ void foo() {}
|
||||||
|
// HOST-LABEL: @_Z18__device_stub__barv(
|
||||||
|
// HOST-NEXT: entry:
|
||||||
|
// HOST-NEXT: [[TMP0:%.*]] = call i32 @cudaLaunch(ptr @_Z18__device_stub__barv)
|
||||||
|
// HOST-NEXT: br label [[SETUP_END:%.*]]
|
||||||
|
// HOST: setup.end:
|
||||||
|
// HOST-NEXT: ret void
|
||||||
|
//
|
||||||
|
__global__ void bar() {}
|
||||||
|
__device__ int x = 1;
|
Loading…
Reference in New Issue