diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/CMakeLists.txt b/clang/lib/CIR/Dialect/Transforms/TargetLowering/CMakeLists.txt index 21bfa30a111a..dab8dbbe5611 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/CMakeLists.txt +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/CMakeLists.txt @@ -13,6 +13,7 @@ add_clang_library(TargetLowering TargetInfo.cpp TargetLoweringInfo.cpp Targets/AArch64.cpp + Targets/NVPTX.cpp Targets/SPIR.cpp Targets/X86.cpp Targets/LoweringPrepareAArch64CXXABI.cpp diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp index 278091070763..e979e813336f 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp @@ -81,6 +81,8 @@ createTargetLoweringInfo(LowerModule &LM) { } case llvm::Triple::spirv64: return createSPIRVTargetLoweringInfo(LM); + case llvm::Triple::nvptx64: + return createNVPTXTargetLoweringInfo(LM); default: cir_cconv_unreachable("ABI NYI"); } diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetInfo.h b/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetInfo.h index 8184c4f0afc2..a03cf711babc 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetInfo.h +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetInfo.h @@ -30,6 +30,9 @@ createAArch64TargetLoweringInfo(LowerModule &CGM, cir::AArch64ABIKind AVXLevel); std::unique_ptr createSPIRVTargetLoweringInfo(LowerModule &CGM); +std::unique_ptr +createNVPTXTargetLoweringInfo(LowerModule &CGM); + } // namespace cir #endif // LLVM_CLANG_LIB_CIR_DIALECT_TRANSFORMS_TARGETLOWERING_TARGETINFO_H diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/NVPTX.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/NVPTX.cpp new file mode 100644 index 000000000000..64c13331d9ba --- /dev/null +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/NVPTX.cpp @@ -0,0 +1,71 @@ +//===- NVPTX.cpp - TargetInfo for NVPTX -----------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "ABIInfoImpl.h" +#include "LowerFunctionInfo.h" +#include "LowerTypes.h" +#include "TargetInfo.h" +#include "TargetLoweringInfo.h" +#include "clang/CIR/ABIArgInfo.h" +#include "clang/CIR/MissingFeatures.h" +#include "llvm/Support/ErrorHandling.h" + +using ABIArgInfo = cir::ABIArgInfo; +using MissingFeature = cir::MissingFeatures; + +namespace cir { + +//===----------------------------------------------------------------------===// +// NVPTX ABI Implementation +//===----------------------------------------------------------------------===// + +namespace { + +class NVPTXABIInfo : public ABIInfo { +public: + NVPTXABIInfo(LowerTypes <) : ABIInfo(lt) {} + +private: + void computeInfo(LowerFunctionInfo &fi) const override { + llvm_unreachable("NYI"); + } +}; + +class NVPTXTargetLoweringInfo : public TargetLoweringInfo { +public: + NVPTXTargetLoweringInfo(LowerTypes <) + : TargetLoweringInfo(std::make_unique(lt)) {} + + unsigned getTargetAddrSpaceFromCIRAddrSpace( + cir::AddressSpaceAttr addressSpaceAttr) const override { + using Kind = cir::AddressSpaceAttr::Kind; + switch (addressSpaceAttr.getValue()) { + case Kind::offload_private: + return 0; + case Kind::offload_local: + return 3; + case Kind::offload_global: + return 1; + case Kind::offload_constant: + return 2; + case Kind::offload_generic: + return 4; + default: + cir_cconv_unreachable("Unknown CIR address space for this target"); + } + } +}; + +} // namespace + +std::unique_ptr +createNVPTXTargetLoweringInfo(LowerModule &lowerModule) { + return std::make_unique(lowerModule.getTypes()); +} + +} // namespace cir diff --git a/clang/test/CIR/CodeGen/CUDA/simple.cu b/clang/test/CIR/CodeGen/CUDA/simple.cu index 51a1d3bb2f4b..905ad8048b05 100644 --- a/clang/test/CIR/CodeGen/CUDA/simple.cu +++ b/clang/test/CIR/CodeGen/CUDA/simple.cu @@ -32,6 +32,10 @@ __global__ void global_fn(int a) {} // CIR-HOST: cir.get_global @_Z24__device_stub__global_fni // CIR-HOST: cir.call @cudaLaunchKernel +// COM: LLVM-HOST: void @_Z24__device_stub__global_fni +// COM: LLVM-HOST: call i32 @__cudaPopCallConfiguration +// COM: LLVM-HOST: call i32 @cudaLaunchKernel(ptr @_Z24__device_stub__global_fni + int main() { global_fn<<<1, 1>>>(1); } @@ -46,3 +50,16 @@ int main() { // CIR-HOST: [[Arg:%[0-9]+]] = cir.const #cir.int<1> // CIR-HOST: cir.call @_Z24__device_stub__global_fni([[Arg]]) // CIR-HOST: } + +// COM: LLVM-HOST: define dso_local i32 @main +// COM: LLVM-HOST: alloca %struct.dim3 +// COM: LLVM-HOST: alloca %struct.dim3 +// COM: LLVM-HOST: call void @_ZN4dim3C1Ejjj +// COM: LLVM-HOST: call void @_ZN4dim3C1Ejjj +// COM: LLVM-HOST: [[LLVMConfigOK:%[0-9]+]] = call i32 @__cudaPushCallConfiguration +// COM: LLVM-HOST: br [[LLVMConfigOK]], label %[[Good:[0-9]+]], label [[Bad:[0-9]+]] +// COM: LLVM-HOST: [[Good]]: +// COM: LLVM-HOST: call void @_Z24__device_stub__global_fni +// COM: LLVM-HOST: br label [[Bad]] +// COM: LLVM-HOST: [[Bad]]: +// COM: LLVM-HOST: ret i32