saiislam updated this revision to Diff 553179. saiislam marked 7 inline comments as done. saiislam added a comment.
Updated test case to check internalization of newly inserted global variable. Repository: rG LLVM Github Monorepo CHANGES SINCE LAST ACTION https://reviews.llvm.org/D139730/new/ https://reviews.llvm.org/D139730 Files: clang/lib/CodeGen/CGBuiltin.cpp clang/lib/CodeGen/CodeGenModule.cpp clang/lib/CodeGen/CodeGenModule.h clang/lib/CodeGen/TargetInfo.h clang/lib/CodeGen/Targets/AMDGPU.cpp clang/lib/Driver/ToolChain.cpp clang/lib/Driver/ToolChains/Clang.cpp clang/test/CodeGenCUDA/amdgpu-code-object-version-linking.cu clang/test/CodeGenCUDA/amdgpu-workgroup-size.cu clang/test/CodeGenOpenCL/opencl_types.cl clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp openmp/libomptarget/DeviceRTL/CMakeLists.txt openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp openmp/libomptarget/plugins-nextgen/amdgpu/utils/UtilitiesRTL.h
Index: openmp/libomptarget/plugins-nextgen/amdgpu/utils/UtilitiesRTL.h =================================================================== --- openmp/libomptarget/plugins-nextgen/amdgpu/utils/UtilitiesRTL.h +++ openmp/libomptarget/plugins-nextgen/amdgpu/utils/UtilitiesRTL.h @@ -25,6 +25,7 @@ #include "llvm/Support/MemoryBufferRef.h" #include "llvm/Support/YAMLTraits.h" +using namespace llvm::ELF; namespace llvm { namespace omp { @@ -34,17 +35,25 @@ // The implicit arguments of AMDGPU kernels. struct AMDGPUImplicitArgsTy { - uint64_t OffsetX; - uint64_t OffsetY; - uint64_t OffsetZ; - uint64_t HostcallPtr; - uint64_t Unused0; - uint64_t Unused1; - uint64_t Unused2; + uint32_t BlockCountX; + uint32_t BlockCountY; + uint32_t BlockCountZ; + uint16_t GroupSizeX; + uint16_t GroupSizeY; + uint16_t GroupSizeZ; + uint8_t Unused0[46]; // 46 byte offset. + uint16_t GridDims; + uint8_t Unused1[190]; // 190 byte offset. }; -static_assert(sizeof(AMDGPUImplicitArgsTy) == 56, - "Unexpected size of implicit arguments"); +enum IMPLICITARGS : uint32_t { + COV4_SIZE = 56, + COV5_SIZE = 256, +}; + +uint16_t getImplicitArgsSize(uint16_t Version) { + return Version < ELF::ELFABIVERSION_AMDGPU_HSA_V5 ? COV4_SIZE : COV5_SIZE; +} /// Parse a TargetID to get processor arch and feature map. /// Returns processor subarch. @@ -295,7 +304,8 @@ /// Reads the AMDGPU specific metadata from the ELF file and propagates the /// KernelInfoMap Error readAMDGPUMetaDataFromImage(MemoryBufferRef MemBuffer, - StringMap<KernelMetaDataTy> &KernelInfoMap) { + StringMap<KernelMetaDataTy> &KernelInfoMap, + uint16_t &ELFABIVersion) { Error Err = Error::success(); // Used later as out-parameter auto ELFOrError = object::ELF64LEFile::create(MemBuffer.getBuffer()); @@ -305,6 +315,12 @@ const object::ELF64LEFile ELFObj = ELFOrError.get(); ArrayRef<object::ELF64LE::Shdr> Sections = cantFail(ELFObj.sections()); KernelInfoReader Reader(KernelInfoMap); + + // Read the code object version from ELF image header + auto Header = ELFObj.getHeader(); + ELFABIVersion = (uint8_t)(Header.e_ident[ELF::EI_ABIVERSION]); + DP("ELFABIVERSION Version: %u\n", ELFABIVersion); + for (const auto &S : Sections) { if (S.sh_type != ELF::SHT_NOTE) continue; Index: openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp =================================================================== --- openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp +++ openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp @@ -381,6 +381,9 @@ /// Get the executable. hsa_executable_t getExecutable() const { return Executable; } + /// Get to Code Object Version of the ELF + uint16_t getELFABIVersion() const { return ELFABIVersion; } + /// Find an HSA device symbol by its name on the executable. Expected<hsa_executable_symbol_t> findDeviceSymbol(GenericDeviceTy &Device, StringRef SymbolName) const; @@ -401,6 +404,7 @@ hsa_executable_t Executable; hsa_code_object_t CodeObject; StringMap<utils::KernelMetaDataTy> KernelInfoMap; + uint16_t ELFABIVersion; }; /// Class implementing the AMDGPU kernel functionalities which derives from the @@ -408,8 +412,7 @@ struct AMDGPUKernelTy : public GenericKernelTy { /// Create an AMDGPU kernel with a name and an execution mode. AMDGPUKernelTy(const char *Name, OMPTgtExecModeFlags ExecutionMode) - : GenericKernelTy(Name, ExecutionMode), - ImplicitArgsSize(sizeof(utils::AMDGPUImplicitArgsTy)) {} + : GenericKernelTy(Name, ExecutionMode) {} /// Initialize the AMDGPU kernel. Error initImpl(GenericDeviceTy &Device, DeviceImageTy &Image) override { @@ -450,6 +453,9 @@ // TODO: Read the kernel descriptor for the max threads per block. May be // read from the image. + ImplicitArgsSize = utils::getImplicitArgsSize(AMDImage.getELFABIVersion()); + DP("ELFABIVersion: %d\n", AMDImage.getELFABIVersion()); + // Get additional kernel info read from image KernelInfo = AMDImage.getKernelInfo(getName()); if (!KernelInfo.has_value()) @@ -476,6 +482,10 @@ /// Get the HSA kernel object representing the kernel function. uint64_t getKernelObject() const { return KernelObject; } + /// Get the size of implicitargs based on the code object version + /// @return 56 for cov4 and 256 for cov5 + uint32_t getImplicitArgsSize() const { return ImplicitArgsSize; } + private: /// The kernel object to execute. uint64_t KernelObject; @@ -486,7 +496,7 @@ uint32_t PrivateSize; /// The size of implicit kernel arguments. - const uint32_t ImplicitArgsSize; + uint32_t ImplicitArgsSize; /// Additional Info for the AMD GPU Kernel std::optional<utils::KernelMetaDataTy> KernelInfo; @@ -2627,8 +2637,8 @@ if (Result) return Plugin::error("Loaded HSA executable does not validate"); - if (auto Err = - utils::readAMDGPUMetaDataFromImage(getMemoryBuffer(), KernelInfoMap)) + if (auto Err = utils::readAMDGPUMetaDataFromImage( + getMemoryBuffer(), KernelInfoMap, ELFABIVersion)) return Err; return Plugin::success(); @@ -2993,6 +3003,14 @@ if (GenericDevice.getRPCServer()) Stream->setRPCServer(GenericDevice.getRPCServer()); + if (getImplicitArgsSize() >= utils::COV5_SIZE) { + ImplArgs->BlockCountX = NumBlocks; + ImplArgs->GroupSizeX = NumThreads; + ImplArgs->GroupSizeY = 1; + ImplArgs->GroupSizeZ = 1; + ImplArgs->GridDims = 1; + } + // Push the kernel launch into the stream. return Stream->pushKernelLaunch(*this, AllArgs, NumThreads, NumBlocks, GroupSize, ArgsMemoryManager); Index: openmp/libomptarget/DeviceRTL/CMakeLists.txt =================================================================== --- openmp/libomptarget/DeviceRTL/CMakeLists.txt +++ openmp/libomptarget/DeviceRTL/CMakeLists.txt @@ -282,7 +282,7 @@ add_custom_target(omptarget.devicertl.amdgpu) foreach(gpu_arch ${LIBOMPTARGET_DEVICE_ARCHITECTURES}) if("${gpu_arch}" IN_LIST all_amdgpu_architectures) - compileDeviceRTLLibrary(${gpu_arch} amdgpu amdgcn-amd-amdhsa) + compileDeviceRTLLibrary(${gpu_arch} amdgpu amdgcn-amd-amdhsa -Xclang -mcode-object-version=none) elseif("${gpu_arch}" IN_LIST all_nvptx_architectures) compileDeviceRTLLibrary(${gpu_arch} nvptx nvptx64-nvidia-cuda --cuda-feature=+ptx61) else() Index: clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp =================================================================== --- clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -403,11 +403,18 @@ llvm::copy(LinkerArgs, std::back_inserter(CmdArgs)); } + // pass on -mllvm options to the clang + for (const opt::Arg *Arg : Args.filtered(OPT_mllvm)) { + CmdArgs.push_back("-mllvm"); + CmdArgs.push_back(Arg->getValue()); + } + if (Args.hasArg(OPT_debug)) CmdArgs.push_back("-g"); - if (SaveTemps) + if (SaveTemps) { CmdArgs.push_back("-save-temps"); + } if (Verbose) CmdArgs.push_back("-v"); Index: clang/test/CodeGenOpenCL/opencl_types.cl =================================================================== --- clang/test/CodeGenOpenCL/opencl_types.cl +++ clang/test/CodeGenOpenCL/opencl_types.cl @@ -1,5 +1,5 @@ -// RUN: %clang_cc1 -cl-std=CL2.0 %s -triple "spir-unknown-unknown" -emit-llvm -o - -O0 | FileCheck %s --check-prefixes=CHECK-COM,CHECK-SPIR -// RUN: %clang_cc1 -cl-std=CL2.0 %s -triple "amdgcn--amdhsa" -emit-llvm -o - -O0 | FileCheck %s --check-prefixes=CHECK-COM,CHECK-AMDGCN +// RUN: %clang_cc1 -cl-std=CL2.0 %s -triple "spir-unknown-unknown" -emit-llvm -o - -O0 | FileCheck %s --check-prefix=CHECK-SPIR +// RUN: %clang_cc1 -cl-std=CL2.0 %s -triple "amdgcn--amdhsa" -emit-llvm -o - -O0 | FileCheck %s --check-prefix=CHECK-AMDGCN #define CLK_ADDRESS_CLAMP_TO_EDGE 2 #define CLK_NORMALIZED_COORDS_TRUE 1 @@ -7,7 +7,6 @@ #define CLK_FILTER_LINEAR 0x20 constant sampler_t glb_smp = CLK_ADDRESS_CLAMP_TO_EDGE|CLK_NORMALIZED_COORDS_TRUE|CLK_FILTER_NEAREST; -// CHECK-COM-NOT: constant i32 void fnc1(image1d_t img) {} // CHECK-SPIR: @fnc1(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0) Index: clang/test/CodeGenCUDA/amdgpu-workgroup-size.cu =================================================================== --- clang/test/CodeGenCUDA/amdgpu-workgroup-size.cu +++ clang/test/CodeGenCUDA/amdgpu-workgroup-size.cu @@ -7,6 +7,10 @@ // RUN: -fcuda-is-device -mcode-object-version=5 -emit-llvm -o - -x hip %s \ // RUN: | FileCheck -check-prefix=COV5 %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \ +// RUN: -fcuda-is-device -mcode-object-version=none -emit-llvm -o - -x hip %s \ +// RUN: | FileCheck -check-prefix=COVNONE %s + #include "Inputs/cuda.h" // PRECOV5-LABEL: test_get_workgroup_size @@ -26,6 +30,36 @@ // COV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef // COV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 16 // COV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef + + +// COVNONE-LABEL: test_get_workgroup_size +// COVNONE: load i32, ptr addrspacecast (ptr addrspace(1) @llvm.amdgcn.abi.version to ptr), align {{.*}} +// COVNONE: [[ABI5_X:%.*]] = icmp sge i32 %{{.*}}, 500 +// COVNONE: call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() +// COVNONE: [[GEP_5_X:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 12 +// COVNONE: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr() +// COVNONE: [[GEP_4_X:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 4 +// COVNONE: select i1 [[ABI5_X]], ptr addrspace(4) [[GEP_5_X]], ptr addrspace(4) [[GEP_4_X]] +// COVNONE: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef + +// COVNONE: load i32, ptr addrspacecast (ptr addrspace(1) @llvm.amdgcn.abi.version to ptr), align {{.*}} +// COVNONE: [[ABI5_Y:%.*]] = icmp sge i32 %{{.*}}, 500 +// COVNONE: call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() +// COVNONE: [[GEP_5_Y:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 14 +// COVNONE: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr() +// COVNONE: [[GEP_4_Y:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 6 +// COVNONE: select i1 [[ABI5_Y]], ptr addrspace(4) [[GEP_5_Y]], ptr addrspace(4) [[GEP_4_Y]] +// COVNONE: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef + +// COVNONE: load i32, ptr addrspacecast (ptr addrspace(1) @llvm.amdgcn.abi.version to ptr), align {{.*}} +// COVNONE: [[ABI5_Z:%.*]] = icmp sge i32 %{{.*}}, 500 +// COVNONE: call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() +// COVNONE: [[GEP_5_Z:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 16 +// COVNONE: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr() +// COVNONE: [[GEP_4_Z:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 8 +// COVNONE: select i1 [[ABI5_Z]], ptr addrspace(4) [[GEP_5_Z]], ptr addrspace(4) [[GEP_4_Z]] +// COVNONE: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef + __device__ void test_get_workgroup_size(int d, int *out) { switch (d) { Index: clang/test/CodeGenCUDA/amdgpu-code-object-version-linking.cu =================================================================== --- /dev/null +++ clang/test/CodeGenCUDA/amdgpu-code-object-version-linking.cu @@ -0,0 +1,96 @@ +// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm-bc \ +// RUN: -mcode-object-version=4 -DUSER -x hip -o %t_4.bc %s + +// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm-bc \ +// RUN: -mcode-object-version=5 -DUSER -x hip -o %t_5.bc %s + +// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm-bc \ +// RUN: -mcode-object-version=none -DDEVICELIB -x hip -o %t_0.bc %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm -O3 \ +// RUN: %t_4.bc -mlink-builtin-bitcode %t_0.bc -o - |\ +// RUN: FileCheck -check-prefix=LINKED4 %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm -O3 \ +// RUN: %t_5.bc -mlink-builtin-bitcode %t_0.bc -o - |\ +// RUN: FileCheck -check-prefix=LINKED5 %s + +#include "Inputs/cuda.h" + +// LINKED4: @llvm.amdgcn.abi.version = weak_odr hidden local_unnamed_addr addrspace(4) constant i32 400 +// LINKED4-LABEL: bar +// LINKED4-NOT: load i32, ptr addrspacecast (ptr addrspace(4) @llvm.amdgcn.abi.version to ptr), align {{.*}} +// LINKED4-NOT: icmp sge i32 %{{.*}}, 500 +// LINKED4: call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() +// LINKED4: [[GEP_5_X:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 12 +// LINKED4: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr() +// LINKED4: [[GEP_4_X:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 4 +// LINKED4: select i1 false, ptr addrspace(4) [[GEP_5_X]], ptr addrspace(4) [[GEP_4_X]] +// LINKED4: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef + +// LINKED4-NOT: load i32, ptr addrspacecast (ptr addrspace(4) @llvm.amdgcn.abi.version to ptr), align {{.*}} +// LINKED4-NOT: icmp sge i32 %{{.*}}, 500 +// LINKED4: call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() +// LINKED4: [[GEP_5_Y:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 14 +// LINKED4: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr() +// LINKED4: [[GEP_4_Y:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 6 +// LINKED4: select i1 false, ptr addrspace(4) [[GEP_5_Y]], ptr addrspace(4) [[GEP_4_Y]] +// LINKED4: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef + +// LINKED4-NOT: load i32, ptr addrspacecast (ptr addrspace(4) @llvm.amdgcn.abi.version to ptr), align {{.*}} +// LINKED4-NOT: icmp sge i32 %{{.*}}, 500 +// LINKED4: call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() +// LINKED4: [[GEP_5_Z:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 16 +// LINKED4: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr() +// LINKED4: [[GEP_4_Z:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 8 +// LINKED4: select i1 false, ptr addrspace(4) [[GEP_5_Z]], ptr addrspace(4) [[GEP_4_Z]] +// LINKED4: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef +// LINKED4: "amdgpu_code_object_version", i32 400 + +// LINKED5: llvm.amdgcn.abi.version = weak_odr hidden local_unnamed_addr addrspace(4) constant i32 500 +// LINKED5-LABEL: bar +// LINKED5-NOT: load i32, ptr addrspacecast (ptr addrspace(4) @llvm.amdgcn.abi.version to ptr), align {{.*}} +// LINKED5-NOT: icmp sge i32 %{{.*}}, 500 +// LINKED5: call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() +// LINKED5: [[GEP_5_X:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 12 +// LINKED5: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr() +// LINKED5: [[GEP_4_X:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 4 +// LINKED5: select i1 true, ptr addrspace(4) [[GEP_5_X]], ptr addrspace(4) [[GEP_4_X]] +// LINKED5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef + +// LINKED5-NOT: load i32, ptr addrspacecast (ptr addrspace(4) @llvm.amdgcn.abi.version to ptr), align {{.*}} +// LINKED5-NOT: icmp sge i32 %{{.*}}, 500 +// LINKED5: call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() +// LINKED5: [[GEP_5_Y:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 14 +// LINKED5: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr() +// LINKED5: [[GEP_4_Y:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 6 +// LINKED5: select i1 true, ptr addrspace(4) [[GEP_5_Y]], ptr addrspace(4) [[GEP_4_Y]] +// LINKED5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef + +// LINKED5-NOT: load i32, ptr addrspacecast (ptr addrspace(4) @llvm.amdgcn.abi.version to ptr), align {{.*}} +// LINKED5-NOT: icmp sge i32 %{{.*}}, 500 +// LINKED5: call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() +// LINKED5: [[GEP_5_Z:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 16 +// LINKED5: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr() +// LINKED5: [[GEP_4_Z:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 8 +// LINKED5: select i1 true, ptr addrspace(4) [[GEP_5_Z]], ptr addrspace(4) [[GEP_4_Z]] +// LINKED5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef +// LINKED5: "amdgpu_code_object_version", i32 500 + +#ifdef DEVICELIB +__device__ void bar(int *x, int *y, int *z) +{ + *x = __builtin_amdgcn_workgroup_size_x(); + *y = __builtin_amdgcn_workgroup_size_y(); + *z = __builtin_amdgcn_workgroup_size_z(); +} +#endif + +#ifdef USER +__device__ void bar(int *x, int *y, int *z); +__device__ void foo() +{ + int *x, *y, *z; + bar(x, y, z); +} +#endif Index: clang/lib/Driver/ToolChains/Clang.cpp =================================================================== --- clang/lib/Driver/ToolChains/Clang.cpp +++ clang/lib/Driver/ToolChains/Clang.cpp @@ -8646,6 +8646,14 @@ CmdArgs.push_back("--device-debug"); } + // code-object-version=X needs to be passed to clang-linker-wrapper to ensure + // that it is used by lld. + if (const Arg *A = Args.getLastArg(options::OPT_mcode_object_version_EQ)) { + CmdArgs.push_back(Args.MakeArgString("-mllvm")); + CmdArgs.push_back(Args.MakeArgString( + Twine("--amdhsa-code-object-version=") + A->getValue())); + } + for (const auto &A : Args.getAllArgValues(options::OPT_Xcuda_ptxas)) CmdArgs.push_back(Args.MakeArgString("--ptxas-arg=" + A)); Index: clang/lib/Driver/ToolChain.cpp =================================================================== --- clang/lib/Driver/ToolChain.cpp +++ clang/lib/Driver/ToolChain.cpp @@ -1365,7 +1365,10 @@ // matches the current toolchain triple. If it is not present // at all, target and host share a toolchain. if (A->getOption().matches(options::OPT_m_Group)) { - if (SameTripleAsHost) + // Pass code object version to device toolchain + // to correctly set metadata in intermediate files. + if (SameTripleAsHost || + A->getOption().matches(options::OPT_mcode_object_version_EQ)) DAL->append(A); else Modified = true; Index: clang/lib/CodeGen/Targets/AMDGPU.cpp =================================================================== --- clang/lib/CodeGen/Targets/AMDGPU.cpp +++ clang/lib/CodeGen/Targets/AMDGPU.cpp @@ -8,6 +8,7 @@ #include "ABIInfoImpl.h" #include "TargetInfo.h" +#include "clang/Basic/TargetOptions.h" using namespace clang; using namespace clang::CodeGen; @@ -274,6 +275,8 @@ void setFunctionDeclAttributes(const FunctionDecl *FD, llvm::Function *F, CodeGenModule &CGM) const; + void emitTargetGlobals(CodeGen::CodeGenModule &CGM) const override; + void setTargetAttributes(const Decl *D, llvm::GlobalValue *GV, CodeGen::CodeGenModule &M) const override; unsigned getOpenCLKernelCallingConv() const override; @@ -354,6 +357,26 @@ } } +/// Emits control constants used to change per-architecture behaviour in the +/// AMDGPU ROCm device libraries. +void AMDGPUTargetCodeGenInfo::emitTargetGlobals( + CodeGen::CodeGenModule &CGM) const { + StringRef Name = "llvm.amdgcn.abi.version"; + if (CGM.getModule().getNamedGlobal(Name)) + return; + + auto *Type = llvm::IntegerType::getIntNTy(CGM.getModule().getContext(), 32); + llvm::Constant *COV = llvm::ConstantInt::get( + Type, CGM.getTarget().getTargetOpts().CodeObjectVersion); + + auto *GV = new llvm::GlobalVariable( + CGM.getModule(), Type, true, llvm::GlobalValue::WeakODRLinkage, COV, Name, + nullptr, llvm::GlobalValue::ThreadLocalMode::NotThreadLocal, + CGM.getContext().getTargetAddressSpace(LangAS::opencl_constant)); + GV->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Local); + GV->setVisibility(llvm::GlobalValue::VisibilityTypes::HiddenVisibility); +} + void AMDGPUTargetCodeGenInfo::setTargetAttributes( const Decl *D, llvm::GlobalValue *GV, CodeGen::CodeGenModule &M) const { if (requiresAMDGPUProtectedVisibility(D, GV)) { Index: clang/lib/CodeGen/TargetInfo.h =================================================================== --- clang/lib/CodeGen/TargetInfo.h +++ clang/lib/CodeGen/TargetInfo.h @@ -81,6 +81,9 @@ CodeGen::CodeGenModule &CGM, const llvm::MapVector<GlobalDecl, StringRef> &MangledDeclNames) const {} + /// Provides a convenient hook to handle extra target-specific globals. + virtual void emitTargetGlobals(CodeGen::CodeGenModule &CGM) const {} + /// Any further codegen related checks that need to be done on a function call /// in a target specific manner. virtual void checkFunctionCallABI(CodeGenModule &CGM, SourceLocation CallLoc, Index: clang/lib/CodeGen/CodeGenModule.h =================================================================== --- clang/lib/CodeGen/CodeGenModule.h +++ clang/lib/CodeGen/CodeGenModule.h @@ -1573,6 +1573,11 @@ void handleAMDGPUWavesPerEUAttr(llvm::Function *F, const AMDGPUWavesPerEUAttr *A); + llvm::Constant * + GetOrCreateLLVMGlobal(StringRef MangledName, llvm::Type *Ty, LangAS AddrSpace, + const VarDecl *D, + ForDefinition_t IsForDefinition = NotForDefinition); + private: llvm::Constant *GetOrCreateLLVMFunction( StringRef MangledName, llvm::Type *Ty, GlobalDecl D, bool ForVTable, @@ -1595,11 +1600,6 @@ void UpdateMultiVersionNames(GlobalDecl GD, const FunctionDecl *FD, StringRef &CurName); - llvm::Constant * - GetOrCreateLLVMGlobal(StringRef MangledName, llvm::Type *Ty, LangAS AddrSpace, - const VarDecl *D, - ForDefinition_t IsForDefinition = NotForDefinition); - bool GetCPUAndFeaturesAttributes(GlobalDecl GD, llvm::AttrBuilder &AttrBuilder, bool SetTargetFeatures = true); Index: clang/lib/CodeGen/CodeGenModule.cpp =================================================================== --- clang/lib/CodeGen/CodeGenModule.cpp +++ clang/lib/CodeGen/CodeGenModule.cpp @@ -1203,6 +1203,8 @@ getModule().addModuleFlag(llvm::Module::Error, "MaxTLSAlign", getContext().getTargetInfo().getMaxTLSAlign()); + getTargetCodeGenInfo().emitTargetGlobals(*this); + getTargetCodeGenInfo().emitTargetMetadata(*this, MangledDeclNames); EmitBackendOptionsMetadata(getCodeGenOpts()); Index: clang/lib/CodeGen/CGBuiltin.cpp =================================================================== --- clang/lib/CodeGen/CGBuiltin.cpp +++ clang/lib/CodeGen/CGBuiltin.cpp @@ -27,6 +27,7 @@ #include "clang/AST/OSLog.h" #include "clang/Basic/TargetBuiltins.h" #include "clang/Basic/TargetInfo.h" +#include "clang/Basic/TargetOptions.h" #include "clang/CodeGen/CGFunctionInfo.h" #include "clang/Frontend/FrontendDiagnostic.h" #include "llvm/ADT/APFloat.h" @@ -17097,24 +17098,58 @@ } // \p Index is 0, 1, and 2 for x, y, and z dimension, respectively. +/// Emit code based on Code Object ABI version. +/// COV_4 : Emit code to use dispatch ptr +/// COV_5 : Emit code to use implicitarg ptr +/// COV_NONE : Emit code to load a global variable "llvm.amdgcn.abi.version" +/// and use its value for COV_4 or COV_5 approach. It is used for +/// compiling device libraries in an ABI-agnostic way. +/// +/// Note: "llvm.amdgcn.abi.version" is supposed to be emitted and intialized by +/// clang during compilation of user code. Value *EmitAMDGPUWorkGroupSize(CodeGenFunction &CGF, unsigned Index) { - bool IsCOV_5 = CGF.getTarget().getTargetOpts().CodeObjectVersion == - clang::TargetOptions::COV_5; - Constant *Offset; - Value *DP; - if (IsCOV_5) { + llvm::LoadInst *LD; + + auto Cov = CGF.getTarget().getTargetOpts().CodeObjectVersion; + + if (Cov == clang::TargetOptions::COV_None) { + auto *ABIVersionC = CGF.CGM.GetOrCreateLLVMGlobal( + "llvm.amdgcn.abi.version", CGF.Int32Ty, LangAS::Default, nullptr, + CodeGen::NotForDefinition); + + Value *ABIVersion = CGF.Builder.CreateAlignedLoad(CGF.Int32Ty, ABIVersionC, + CGF.CGM.getIntAlign()); + + Value *IsCOV5 = CGF.Builder.CreateICmpSGE( + ABIVersion, + llvm::ConstantInt::get(CGF.Int32Ty, clang::TargetOptions::COV_5)); + // Indexing the implicit kernarg segment. - Offset = llvm::ConstantInt::get(CGF.Int32Ty, 12 + Index * 2); - DP = EmitAMDGPUImplicitArgPtr(CGF); - } else { + Value *ImplicitGEP = CGF.Builder.CreateConstGEP1_32( + CGF.Int8Ty, EmitAMDGPUImplicitArgPtr(CGF), 12 + Index * 2); + // Indexing the HSA kernel_dispatch_packet struct. - Offset = llvm::ConstantInt::get(CGF.Int32Ty, 4 + Index * 2); - DP = EmitAMDGPUDispatchPtr(CGF); + Value *DispatchGEP = CGF.Builder.CreateConstGEP1_32( + CGF.Int8Ty, EmitAMDGPUDispatchPtr(CGF), 4 + Index * 2); + + auto Result = CGF.Builder.CreateSelect(IsCOV5, ImplicitGEP, DispatchGEP); + LD = CGF.Builder.CreateLoad( + Address(Result, CGF.Int16Ty, CharUnits::fromQuantity(2))); + } else { + Value *GEP = nullptr; + if (Cov == clang::TargetOptions::COV_5) { + // Indexing the implicit kernarg segment. + GEP = CGF.Builder.CreateConstGEP1_32( + CGF.Int8Ty, EmitAMDGPUImplicitArgPtr(CGF), 12 + Index * 2); + } else { + // Indexing the HSA kernel_dispatch_packet struct. + GEP = CGF.Builder.CreateConstGEP1_32( + CGF.Int8Ty, EmitAMDGPUDispatchPtr(CGF), 4 + Index * 2); + } + LD = CGF.Builder.CreateLoad( + Address(GEP, CGF.Int16Ty, CharUnits::fromQuantity(2))); } - auto *GEP = CGF.Builder.CreateGEP(CGF.Int8Ty, DP, Offset); - auto *LD = CGF.Builder.CreateLoad( - Address(GEP, CGF.Int16Ty, CharUnits::fromQuantity(2))); llvm::MDBuilder MDHelper(CGF.getLLVMContext()); llvm::MDNode *RNode = MDHelper.createRange(APInt(16, 1), APInt(16, CGF.getTarget().getMaxOpenCLWorkGroupSize() + 1));
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits