[CIR][AMDGPU] Add AMDGPU target support for AST to CIR codegen

clang/include/clang/CIR/Dialect/IR/CIROps.td

@@ -4003,7 +4003,8 @@ def CIR_CallingConv : CIR_I32EnumAttr<"CallingConv", "calling convention", [
   I32EnumAttrCase<"SpirKernel", 2, "spir_kernel">,
   I32EnumAttrCase<"SpirFunction", 3, "spir_function">,
   I32EnumAttrCase<"OpenCLKernel", 4, "opencl_kernel">,
-  I32EnumAttrCase<"PTXKernel", 5, "ptx_kernel">
+  I32EnumAttrCase<"PTXKernel", 5, "ptx_kernel">,
+  I32EnumAttrCase<"AMDGPUKernel", 6, "amdgpu_kernel">
 ]>;
 
 def CIR_OptionalPriorityAttr : OptionalAttr<

clang/lib/CIR/CodeGen/CIRGenCall.cpp

@@ -1393,10 +1393,10 @@ CIRGenTypes::arrangeFunctionDeclaration(const FunctionDecl *FD) {
     if (MD->isInstance())
       return arrangeCXXMethodDeclaration(MD);
 
-  auto FTy = FD->getType()->getCanonicalTypeUnqualified();
+  CanQualType FTy = FD->getType()->getCanonicalTypeUnqualified();
 
   assert(isa<FunctionType>(FTy));
-  // TODO: setCUDAKernelCallingConvention
+  setCUDAKernelCallingConvention(FTy, CGM, FD);
 
   // When declaring a function without a prototype, always use a non-variadic
   // type.

clang/lib/CIR/CodeGen/TargetInfo.cpp

@@ -371,15 +371,21 @@ class AMDGPUABIInfo : public ABIInfo {
 
   cir::ABIArgInfo classifyReturnType(QualType retTy) const;
   cir::ABIArgInfo classifyArgumentType(QualType ty) const;
+  cir::ABIArgInfo classifyKernelArgumentType(QualType ty) const;
 };
 
 class AMDGPUTargetCIRGenInfo : public TargetCIRGenInfo {
 public:
   AMDGPUTargetCIRGenInfo(CIRGenTypes &cgt)
       : TargetCIRGenInfo(std::make_unique<AMDGPUABIInfo>(cgt)) {}
 
+  cir::CallingConv getOpenCLKernelCallingConv() const override {
+    return cir::CallingConv::AMDGPUKernel;
+  }
+
   void setCUDAKernelCallingConvention(const FunctionType *&ft) const override {
-    llvm_unreachable("NYI");
+    ft = getABIInfo().getContext().adjustFunctionType(
+        ft, ft->getExtInfo().withCallingConv(CC_DeviceKernel));
   }
 };
 
@@ -519,13 +525,44 @@ cir::ABIArgInfo NVPTXABIInfo::classifyArgumentType(QualType ty) const {
   llvm_unreachable("not yet implemented");
 }
 
-// Skeleton only. Implement when used in TargetLower stage.
 cir::ABIArgInfo AMDGPUABIInfo::classifyReturnType(QualType retTy) const {
-  llvm_unreachable("not yet implemented");
+  if (retTy->isVoidType())
+    return cir::ABIArgInfo::getIgnore();
+
+  if (!isAggregateTypeForABI(retTy)) {
+    return (isPromotableIntegerTypeForABI(retTy)
+                ? cir::ABIArgInfo::getExtend(retTy)
+                : cir::ABIArgInfo::getDirect());
+  }
+
+  llvm_unreachable("Aggregate types NYI for AMDGPU");
 }
 
 cir::ABIArgInfo AMDGPUABIInfo::classifyArgumentType(QualType ty) const {
-  llvm_unreachable("not yet implemented");
+  ty = useFirstFieldIfTransparentUnion(ty);
+
+  if (!isAggregateTypeForABI(ty)) {
+    return (isPromotableIntegerTypeForABI(ty) ? cir::ABIArgInfo::getExtend(ty)
+                                              : cir::ABIArgInfo::getDirect());
+  }
+
+  llvm_unreachable("Aggregate types NYI for AMDGPU");
+}
+
+// TODO(CIR): This method is not currently called in AST->CIR translation.
+cir::ABIArgInfo AMDGPUABIInfo::classifyKernelArgumentType(QualType ty) const {
+  ty = useFirstFieldIfTransparentUnion(ty);
+
+  if (isAggregateTypeForABI(ty)) {
+    llvm_unreachable("Aggregate types NYI for AMDGPU");
+  }
+
+  // For kernels, all parameters passed directly via special buffer. It doesn't
+  // make sense to pass anything byval, so everything must be direct.
+  // Set CanBeFlattened=false to prevent struct expansion.
+  return cir::ABIArgInfo::getDirect(/*T=*/nullptr, /*Offset=*/0,
+                                    /*Padding=*/nullptr,
+                                    /*CanBeFlattened=*/false);
 }
 
 ABIInfo::~ABIInfo() {}

clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp

@@ -901,6 +901,8 @@ mlir::LLVM::CConv convertCallingConv(cir::CallingConv callingConv) {
     llvm_unreachable("NYI");
   case CIR::PTXKernel:
     return LLVM::PTX_Kernel;
+  case CIR::AMDGPUKernel:
+    return LLVM::AMDGPU_KERNEL;
   }
   llvm_unreachable("Unknown calling convention");
 }

clang/test/CIR/CodeGen/HIP/amdgpu-abi-simple-types.cl

@@ -0,0 +1,42 @@
+// RUN: %clang_cc1 %s -fclangir -triple amdgcn-amd-amdhsa -emit-cir -o %t.cir
+// RUN: FileCheck --input-file=%t.cir %s
+
+// Test int return type
+// CHECK: cir.func{{.*}} @return_int() -> !s32i
+int return_int() { return 42; }
+
+// Test void return type
+// CHECK: cir.func{{.*}} @return_void()
+void return_void() {}
+
+// Test char argument
+// CHECK: cir.func{{.*}} @char_arg(%arg{{[0-9]+}}: !s8i{{.*}})
+int char_arg(char c) { return c; }
+
+// Test short argument
+// CHECK: cir.func{{.*}} @short_arg(%arg{{[0-9]+}}: !s16i{{.*}})
+int short_arg(short s) { return s; }
+
+// Test int argument
+// CHECK: cir.func{{.*}} @int_arg(%arg{{[0-9]+}}: !s32i{{.*}})
+int int_arg(int i) { return i; }
+
+// Test long argument
+// CHECK: cir.func{{.*}} @long_arg(%arg{{[0-9]+}}: !s64i{{.*}})
+long long_arg(long l) { return l; }
+
+// Test float argument
+// CHECK: cir.func{{.*}} @float_arg(%arg{{[0-9]+}}: !cir.float{{.*}})
+float float_arg(float f) { return f; }
+
+// Test double argument
+// CHECK: cir.func{{.*}} @double_arg(%arg{{[0-9]+}}: !cir.double{{.*}})
+double double_arg(double d) { return d; }
+
+// Test pointer argument
+// CHECK: cir.func{{.*}} @ptr_arg(%arg{{[0-9]+}}: !cir.ptr<!s32i{{.*}}>{{.*}})
+int* ptr_arg(int* p) { return p; }
+
+// Test multiple arguments
+// CHECK: cir.func{{.*}} @multi_arg(%arg{{[0-9]+}}: !s32i{{.*}}, %arg{{[0-9]+}}: !cir.float{{.*}}, %arg{{[0-9]+}}: !cir.ptr<!s32i{{.*}}>{{.*}})
+int multi_arg(int a, float b, int* c) { return a; }

clang/test/CIR/CodeGen/HIP/amdgpu-hip-kernel-abi.hip

@@ -0,0 +1,28 @@
+#include "../Inputs/cuda.h"
+
+// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \
+// RUN:            -fcuda-is-device -fhip-new-launch-api \
+// RUN:            -I%S/../../Inputs/ -emit-cir %s -o %t.cir
+// RUN: FileCheck --input-file=%t.cir %s
+
+// Test simple kernel
+// CHECK: cir.func{{.*}} @_Z13simple_kerneli(%arg{{[0-9]+}}: !s32i{{.*}}) cc(amdgpu_kernel)
+__global__ void simple_kernel(int x) {}
+
+// Test kernel with pointer
+// CHECK: cir.func{{.*}} @_Z15kernel_with_ptrPi(%arg{{[0-9]+}}: !cir.ptr<!s32i{{.*}}>{{.*}}) cc(amdgpu_kernel)
+__global__ void kernel_with_ptr(int *ptr) {}
+
+// Test kernel with multiple args
+// CHECK: cir.func{{.*}} @_Z16kernel_multi_argifdPi(%arg{{[0-9]+}}: !s32i{{.*}}, %arg{{[0-9]+}}: !cir.float{{.*}}, %arg{{[0-9]+}}: !cir.double{{.*}}, %arg{{[0-9]+}}: !cir.ptr<!s32i{{.*}}>{{.*}}) cc(amdgpu_kernel)
+__global__ void kernel_multi_arg(int a, float b, double c, int *d) {}
+
+// Test device function (NOT a kernel)
+// CHECK: cir.func{{.*}} @_Z9device_fni(%arg{{[0-9]+}}: !s32i{{.*}})
+// CHECK-NOT: cc(amdgpu_kernel)
+__device__ void device_fn(int x) {}
+
+// Test device function with return value
+// CHECK: cir.func{{.*}} @_Z13device_fn_reti(%arg{{[0-9]+}}: !s32i{{.*}}) -> !s32i
+// CHECK-NOT: cc(amdgpu_kernel)
+__device__ int device_fn_ret(int x) { return x + 1; }

clang/test/CIR/CodeGen/HIP/amdgpu-kernel-calling-conv.cl

@@ -0,0 +1,19 @@
+// RUN: %clang_cc1 %s -fclangir -triple amdgcn-amd-amdhsa -emit-cir -o %t.cir
+// RUN: FileCheck --input-file=%t.cir %s
+
+// Test kernel function with amdgpu_kernel calling convention
+// CHECK: cir.func{{.*}} @simple_kernel() cc(amdgpu_kernel)
+__kernel void simple_kernel() {}
+
+// Test kernel with simple integer argument
+// CHECK: cir.func{{.*}} @kernel_with_int(%arg{{[0-9]+}}: !s32i{{.*}}) cc(amdgpu_kernel)
+__kernel void kernel_with_int(int x) {}
+
+// Test kernel with pointer argument (should be in global address space)
+// CHECK: cir.func{{.*}} @kernel_with_ptr(%arg{{[0-9]+}}: !cir.ptr<!s32i, lang_address_space(offload_global)>{{.*}}) cc(amdgpu_kernel)
+__kernel void kernel_with_ptr(global int *ptr) {}
+
+// Test device function (should NOT have amdgpu_kernel calling convention)
+// CHECK: cir.func{{.*}} @device_fn
+// CHECK-NOT: cc(amdgpu_kernel)
+void device_fn(int x) {}

clang/test/CIR/CodeGen/HIP/simple.cpp

@@ -64,8 +64,8 @@ __device__ void device_fn(int* a, double b, float c) {}
 // CIR-DEVICE: cir.func {{.*}} @_Z9device_fnPidf
 
 __global__ void global_fn(int a) {}
-// CIR-DEVICE: @_Z9global_fni
-// LLVM-DEVICE: define dso_local void @_Z9global_fni
+// CIR-DEVICE: @_Z9global_fni{{.*}} cc(amdgpu_kernel)
+// LLVM-DEVICE: define dso_local amdgpu_kernel void @_Z9global_fni
 // OGCG-DEVICE: define dso_local amdgpu_kernel void @_Z9global_fni 
 
 // CIR-HOST: @_Z24__device_stub__global_fni{{.*}}extra([[Kernel]])