[CIR][CUDA] Skeleton of NVPTX target lowering info (#1358)

Added a skeleton of NVPTX target lowering info.

This enables lowering of `simple.cu` (as it hardly tests device side
functionalities), so a test of LLVM IR is also added onto it.

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

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");
   }

clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetInfo.h

@@ -30,6 +30,9 @@ createAArch64TargetLoweringInfo(LowerModule &CGM, cir::AArch64ABIKind AVXLevel);
 std::unique_ptr<TargetLoweringInfo>
 createSPIRVTargetLoweringInfo(LowerModule &CGM);
 
+std::unique_ptr<TargetLoweringInfo>
+createNVPTXTargetLoweringInfo(LowerModule &CGM);
+
 } // namespace cir
 
 #endif // LLVM_CLANG_LIB_CIR_DIALECT_TRANSFORMS_TARGETLOWERING_TARGETINFO_H

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 &lt) : ABIInfo(lt) {}
+
+private:
+  void computeInfo(LowerFunctionInfo &fi) const override {
+    llvm_unreachable("NYI");
+  }
+};
+
+class NVPTXTargetLoweringInfo : public TargetLoweringInfo {
+public:
+  NVPTXTargetLoweringInfo(LowerTypes &lt)
+      : TargetLoweringInfo(std::make_unique<NVPTXABIInfo>(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<TargetLoweringInfo>
+createNVPTXTargetLoweringInfo(LowerModule &lowerModule) {
+  return std::make_unique<NVPTXTargetLoweringInfo>(lowerModule.getTypes());
+}
+
+} // namespace cir

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