[clang] [CIR][CUDA][NVPTX] Set ptx_kernel calling convention on CUDA kernels (PR #195382)

[David Rivera via cfe-commits]

https://github.com/RiverDave updated https://github.com/llvm/llvm-project/pull/195382

>From 86539463e003ee26dbed03d1d119f1e627e72ea4 Mon Sep 17 00:00:00 2001
From: David Rivera <davidriverg at gmail.com>
Date: Fri, 1 May 2026 19:19:43 -0400
Subject: [PATCH 1/2] [CIR][CUDA] Set ptx_kernel calling convention on CUDA
 kernels

---
 clang/lib/CIR/CodeGen/TargetInfo.cpp      | 18 ++++++++++
 clang/test/CIR/CodeGenCUDA/ptx-kernels.cu | 42 +++++++++++++++++++++++
 2 files changed, 60 insertions(+)
 create mode 100644 clang/test/CIR/CodeGenCUDA/ptx-kernels.cu

diff --git a/clang/lib/CIR/CodeGen/TargetInfo.cpp b/clang/lib/CIR/CodeGen/TargetInfo.cpp
index fc939cd9605ab..4390397754e3a 100644
--- a/clang/lib/CIR/CodeGen/TargetInfo.cpp
+++ b/clang/lib/CIR/CodeGen/TargetInfo.cpp
@@ -132,6 +132,24 @@ class NVPTXTargetCIRGenInfo : public TargetCIRGenInfo {
 public:
   NVPTXTargetCIRGenInfo(CIRGenTypes &cgt)
       : TargetCIRGenInfo(std::make_unique<NVPTXABIInfo>(cgt)) {}
+
+  void setTargetAttributes(const clang::Decl *decl, mlir::Operation *global,
+                           CIRGenModule &cgm) const override {
+    auto func = mlir::dyn_cast<cir::FuncOp>(global);
+    if (!func || func.isDeclaration())
+      return;
+
+    const auto *fd = dyn_cast_or_null<FunctionDecl>(decl);
+    if (!fd)
+      return;
+
+    if (cgm.getLangOpts().CUDA && fd->hasAttr<CUDAGlobalAttr>())
+      func.setCallingConv(cir::CallingConv::PTXKernel);
+
+    // TODO(CIR): NoInline on kernels, CUDALaunchBoundsAttr,
+    // CUDAGridConstantAttr param attrs, nvvm.annotations for
+    // surface/texture VarDecls.
+  }
 };
 } // namespace
 
diff --git a/clang/test/CIR/CodeGenCUDA/ptx-kernels.cu b/clang/test/CIR/CodeGenCUDA/ptx-kernels.cu
new file mode 100644
index 0000000000000..155e59638eac7
--- /dev/null
+++ b/clang/test/CIR/CodeGenCUDA/ptx-kernels.cu
@@ -0,0 +1,42 @@
+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda -fclangir \
+// RUN:            -fcuda-is-device -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR %s --input-file=%t.cir
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda -fclangir \
+// RUN:            -fcuda-is-device -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM %s --input-file=%t.ll
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda \
+// RUN:            -fcuda-is-device -emit-llvm %s -o %t.ogcg.ll
+// RUN: FileCheck --check-prefix=OGCG %s --input-file=%t.ogcg.ll
+
+#include "Inputs/cuda.h"
+
+// CIR: cir.func {{.*}} @device_function()
+// LLVM: define{{.*}} void @device_function
+// OGCG: define{{.*}} void @device_function
+extern "C"
+__device__ void device_function() {}
+
+// CIR: cir.func {{.*}} @global_function() cc(ptx_kernel)
+// LLVM: define{{.*}} ptx_kernel void @global_function
+// OGCG: define{{.*}} ptx_kernel void @global_function
+extern "C"
+__global__ void global_function() {
+  device_function();
+}
+
+template <typename T> __global__ void templated_kernel(T param) {}
+template __global__ void templated_kernel<int>(int);
+// CIR-DAG: cir.func {{.*}} @_Z16templated_kernelIiEvT_({{.*}}) cc(ptx_kernel)
+// LLVM-DAG: define{{.*}} ptx_kernel void @_Z16templated_kernelIiEvT_(
+// OGCG-DAG: define{{.*}} ptx_kernel void @_Z16templated_kernelIiEvT_(
+
+namespace {
+__global__ void anonymous_ns_kernel() {}
+// CIR-DAG: cir.func {{.*}} @_ZN12_GLOBAL__N_119anonymous_ns_kernelEv() cc(ptx_kernel)
+// LLVM-DAG: define{{.*}} ptx_kernel void @_ZN12_GLOBAL__N_119anonymous_ns_kernelEv(
+// OGCG-DAG: define{{.*}} ptx_kernel void @_ZN12_GLOBAL__N_119anonymous_ns_kernelEv(
+}

>From 627813ca996c6f23988c36da8781e1ab6c396c28 Mon Sep 17 00:00:00 2001
From: David Rivera <davidriverg at gmail.com>
Date: Fri, 1 May 2026 20:53:27 -0400
Subject: [PATCH 2/2] Fix cc mismatch

---
 clang/test/CIR/CodeGenCUDA/address-spaces.cu | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/clang/test/CIR/CodeGenCUDA/address-spaces.cu b/clang/test/CIR/CodeGenCUDA/address-spaces.cu
index cc1791a8f2244..2f235c8702899 100644
--- a/clang/test/CIR/CodeGenCUDA/address-spaces.cu
+++ b/clang/test/CIR/CodeGenCUDA/address-spaces.cu
@@ -86,7 +86,7 @@ __global__ void fn() {
 // CIR-DEVICE:   cir.store {{.*}}%[[VAL]], %[[J]] : !s32i, !cir.ptr<!s32i>
 // CIR-DEVICE:   cir.return
 
-// LLVM-DEVICE: define dso_local void @_Z2fnv()
+// LLVM-DEVICE: define dso_local ptx_kernel void @_Z2fnv()
 // LLVM-DEVICE:   %[[ALLOCA:.*]] = alloca i32, i64 1, align 4
 // LLVM-DEVICE:   store i32 0, ptr %[[ALLOCA]], align 4
 // LLVM-DEVICE:   %[[VAL:.*]] = load i32, ptr %[[ALLOCA]], align 4