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

clang/lib/CIR/CodeGen/Targets/AMDGPU.cpp

diff --git a/clang/lib/CIR/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CIR/CodeGen/Targets/AMDGPU.cpp
index 566bca1..7f1d903 100644
--- a/clang/lib/CIR/CodeGen/Targets/AMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/Targets/AMDGPU.cpp
@@ -240,7 +240,8 @@
     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);

clang/test/CIR/CodeGenHIP/amdgpu-attrs.hip

diff --git a/clang/test/CIR/CodeGenHIP/amdgpu-attrs.hip b/clang/test/CIR/CodeGenHIP/amdgpu-attrs.hip
index 6845373..636cb52 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