[clang] 857baf3 - [CIR][AMDGPU] Set amdgpu_kernel calling convention on HIP kernels (#195381)

[via cfe-commits]

Author: David Rivera
Date: 2026-05-04T08:29:44-04:00
New Revision: 857baf37de587a94b0527870aa7c72b69e4ffc44

URL: https://github.com/llvm/llvm-project/commit/857baf37de587a94b0527870aa7c72b69e4ffc44
DIFF: https://github.com/llvm/llvm-project/commit/857baf37de587a94b0527870aa7c72b69e4ffc44.diff

LOG: [CIR][AMDGPU] Set amdgpu_kernel calling convention on HIP kernels (#195381)

Added: 
    

Modified: 
    clang/lib/CIR/CodeGen/Targets/AMDGPU.cpp
    clang/test/CIR/CodeGenHIP/amdgpu-attrs.hip

Removed: 
    


################################################################################
diff  --git a/clang/lib/CIR/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CIR/CodeGen/Targets/AMDGPU.cpp
index 566bca14c4d59..7f1d903853660 100644
--- a/clang/lib/CIR/CodeGen/Targets/AMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/Targets/AMDGPU.cpp
@@ -240,7 +240,8 @@ void clang::CIRGen::setAMDGPUTargetFunctionAttributes(const Decl *decl,
     const bool isHIPKernel =
         cgm.getLangOpts().HIP && fd->hasAttr<CUDAGlobalAttr>();
 
-    // TODO(CIR): Set amdgpu_kernel calling convention for HIP kernels.
+    if (isHIPKernel)
+      func.setCallingConv(cir::CallingConv::AMDGPUKernel);
 
     handleAMDGPUFlatWorkGroupSizeAttr(fd, func, cgm, builder, isOpenCLKernel,
                                       isHIPKernel);

diff  --git a/clang/test/CIR/CodeGenHIP/amdgpu-attrs.hip b/clang/test/CIR/CodeGenHIP/amdgpu-attrs.hip
index 68453731b6f39..636cb529cafb7 100644
--- a/clang/test/CIR/CodeGenHIP/amdgpu-attrs.hip
+++ b/clang/test/CIR/CodeGenHIP/amdgpu-attrs.hip
@@ -25,8 +25,8 @@
 // Test that AMDGPU-specific attributes are generated for HIP kernels
 
 // Test: Default attributes for simple kernel
-// CIR: cir.func{{.*}} @_Z13kernel_simplev(){{.*}}"cir.amdgpu-flat-work-group-size" = "1,1024"
-// LLVM: define{{.*}} void @_Z13kernel_simplev(){{.*}} #[[SIMPLE_ATTR:[0-9]+]]
+// CIR: cir.func{{.*}} @_Z13kernel_simplev() cc(amdgpu_kernel){{.*}}"cir.amdgpu-flat-work-group-size" = "1,1024"
+// LLVM: define{{.*}} amdgpu_kernel void @_Z13kernel_simplev(){{.*}} #[[SIMPLE_ATTR:[0-9]+]]
 __global__ void kernel_simple() {}
 
 // Test: Explicit flat work group size attribute