[llvm-branch-commits] [clang] [CIR] Enable CallConvLowering for AMDGPU triples (PR #233822)
via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Tue Oct 20 23:12:36 PDT 2026
https://github.com/skc7 updated https://github.com/llvm/llvm-project/pull/233822
>From fcfded0531fde09ce00e336fd11305e5843acb92 Mon Sep 17 00:00:00 2001
From: skc7 <Krishna.Sankisa at amd.com>
Date: Tue, 1 Sep 2026 14:31:15 +0530
Subject: [PATCH] [CIR] Enable CallConvLowering for AMDGPU triples
---
clang/lib/CIR/Lowering/CIRPasses.cpp | 6 +-
.../call-conv-lowering-amdgpu-arg-types.c | 165 ++++++++++++++++++
.../call-conv-lowering-amdgpu-cxx-records.cpp | 74 ++++++++
.../call-conv-lowering-amdgpu-return-types.c | 158 +++++++++++++++++
.../CodeGenHIP/amdgcn-buffer-rsrc-type.hip | 8 +-
.../call-conv-lowering-amdgpu-kernel-args.hip | 137 +++++++++++++++
.../CodeGenHIP/cleanup-alloca-addrspace.hip | 8 +-
.../CIR/CodeGenHIP/lambda-reference-var.hip | 8 +-
.../CodeGenSYCL/kernel-caller-entry-point.cpp | 6 +-
9 files changed, 558 insertions(+), 12 deletions(-)
create mode 100644 clang/test/CIR/CodeGen/call-conv-lowering-amdgpu-arg-types.c
create mode 100644 clang/test/CIR/CodeGen/call-conv-lowering-amdgpu-cxx-records.cpp
create mode 100644 clang/test/CIR/CodeGen/call-conv-lowering-amdgpu-return-types.c
create mode 100644 clang/test/CIR/CodeGenHIP/call-conv-lowering-amdgpu-kernel-args.hip
diff --git a/clang/lib/CIR/Lowering/CIRPasses.cpp b/clang/lib/CIR/Lowering/CIRPasses.cpp
index 6e75ce03d14ffd..5e49a3f5601870 100644
--- a/clang/lib/CIR/Lowering/CIRPasses.cpp
+++ b/clang/lib/CIR/Lowering/CIRPasses.cpp
@@ -26,6 +26,8 @@ static CallConvTarget getCallConvTarget(const llvm::Triple &triple) {
// Windows is not supported. UEFI shares its convention.
if (triple.getArch() == llvm::Triple::x86_64 && !triple.isOSWindowsOrUEFI())
return CallConvTarget::X86_64;
+ if (triple.isAMDGPU())
+ return CallConvTarget::AMDGPU;
return CallConvTarget::None;
}
@@ -126,8 +128,8 @@ runCIRToCIRPasses(mlir::ModuleOp theModule, mlir::MLIRContext &mlirContext,
if (enableCallConvLowering) {
// CallConvLowering rewrites signatures and call sites using the classifier,
// so it must run after CXXABILowering has lowered C++ ABI types to plain
- // records the classifier can handle. Only the x86_64 System V classifier
- // is implemented; other targets are left unchanged.
+ // records the classifier can handle. Only the x86_64 System V and AMDGPU
+ // classifiers are implemented. Other targets are left unchanged.
CallConvTarget target = getCallConvTarget(triple);
if (target != CallConvTarget::None) {
// Source the ABI-compatibility version from the module's serialized
diff --git a/clang/test/CIR/CodeGen/call-conv-lowering-amdgpu-arg-types.c b/clang/test/CIR/CodeGen/call-conv-lowering-amdgpu-arg-types.c
new file mode 100644
index 00000000000000..ef65e90a528166
--- /dev/null
+++ b/clang/test/CIR/CodeGen/call-conv-lowering-amdgpu-arg-types.c
@@ -0,0 +1,165 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fclangir -fclangir-call-conv-lowering -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
+// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fclangir -fclangir-call-conv-lowering -emit-llvm %s -o %t-cir.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s
+// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
+
+// Checks that CallConvLowering classifies arguments of ordinary AMDGPU
+// functions, which use the default convention rather than the kernel one.
+
+// TODO(cir): Add the argument cases still NYI in CallConvLowering:
+// - Aggregates of 33 to 64 bits, which classic coerces to [2 x i32].
+// abiTypeToCIR has no array case yet.
+// - Aggregates over 64 bits that fit the 16-register budget. The classifier
+// returns Direct with no coerce type, which the bridge rejects.
+// - Aggregates past the register budget, passed byref in addrspace(5).
+// - Structs with a flexible array member and transparent unions. The bridge
+// does not set those record flags yet.
+// - Arguments passed through an ellipsis.
+
+typedef struct {} empty_t;
+typedef struct { char c; } i8_wrapper_t;
+typedef struct { int i; } i32_wrapper_t;
+typedef struct { float f; } f32_wrapper_t;
+typedef struct { int *p; } ptr_wrapper_t;
+typedef struct { char a, b; } char2_t;
+typedef struct { char a, b, c; } small_t;
+typedef union { int i; float f; } int_float_t;
+typedef short short4 __attribute__((ext_vector_type(4)));
+typedef float float3 __attribute__((ext_vector_type(3)));
+
+void arg_void(void) {}
+
+// CIR: cir.func {{.*}}@arg_void()
+// LLVM: define {{.*}}void @arg_void()
+
+// Sub-word integers and bool carry signext or zeroext per their signedness.
+void arg_bool(_Bool b) {}
+
+// CIR: cir.func {{.*}}@arg_bool(%arg0: !cir.bool {{.*}}llvm.zeroext
+// LLVM: define {{.*}}void @arg_bool(i1 noundef zeroext %{{.*}})
+
+void arg_char(char c) {}
+
+// CIR: cir.func {{.*}}@arg_char(%arg0: !s8i {{.*}}llvm.signext
+// LLVM: define {{.*}}void @arg_char(i8 noundef signext %{{.*}})
+
+void arg_uchar(unsigned char c) {}
+
+// CIR: cir.func {{.*}}@arg_uchar(%arg0: !u8i {{.*}}llvm.zeroext
+// LLVM: define {{.*}}void @arg_uchar(i8 noundef zeroext %{{.*}})
+
+void arg_short(short s) {}
+
+// CIR: cir.func {{.*}}@arg_short(%arg0: !s16i {{.*}}llvm.signext
+// LLVM: define {{.*}}void @arg_short(i16 noundef signext %{{.*}})
+
+// Word-sized and wider scalars pass through unchanged.
+void arg_int(int i) {}
+
+// CIR: cir.func {{.*}}@arg_int(%arg0: !s32i {{.*}})
+// LLVM: define {{.*}}void @arg_int(i32 noundef %{{.*}})
+
+void arg_long(long l) {}
+
+// CIR: cir.func {{.*}}@arg_long(%arg0: !s64i {{.*}})
+// LLVM: define {{.*}}void @arg_long(i64 noundef %{{.*}})
+
+void arg_float(float f) {}
+
+// CIR: cir.func {{.*}}@arg_float(%arg0: !cir.float {{.*}})
+// LLVM: define {{.*}}void @arg_float(float noundef %{{.*}})
+
+void arg_double(double d) {}
+
+// CIR: cir.func {{.*}}@arg_double(%arg0: !cir.double {{.*}})
+// LLVM: define {{.*}}void @arg_double(double noundef %{{.*}})
+
+void arg_half(_Float16 h) {}
+
+// CIR: cir.func {{.*}}@arg_half(%arg0: !cir.f16 {{.*}})
+// LLVM: define {{.*}}void @arg_half(half noundef %{{.*}})
+
+// Outside a HIP kernel a generic pointer stays generic.
+void arg_ptr(int *p) {}
+
+// CIR: cir.func {{.*}}@arg_ptr(%arg0: !cir.ptr<!s32i> {{.*}})
+// LLVM: define {{.*}}void @arg_ptr(ptr noundef %{{.*}})
+
+// A _BitInt wider than 64 bits and vectors, including 16-bit and 3-element
+// ones, pass whole.
+void arg_bitint65(_BitInt(65) i) {}
+
+// CIR: cir.func {{.*}}@arg_bitint65(%arg0: !cir.int<s, 65, bitint> {{.*}})
+// LLVM: define {{.*}}void @arg_bitint65(i65 noundef %{{.*}})
+
+void arg_short4(short4 v) {}
+
+// CIR: cir.func {{.*}}@arg_short4(%arg0: !cir.vector<4 x !s16i> {{.*}})
+// LLVM: define {{.*}}void @arg_short4(<4 x i16> noundef %{{.*}})
+
+void arg_float3(float3 v) {}
+
+// CIR: cir.func {{.*}}@arg_float3(%arg0: !cir.vector<3 x !cir.float> {{.*}})
+// LLVM: define {{.*}}void @arg_float3(<3 x float> noundef %{{.*}})
+
+// An empty struct is dropped from the signature.
+void arg_empty(empty_t e) {}
+
+// CIR: cir.func {{.*}}@arg_empty()
+// LLVM: define {{.*}}void @arg_empty()
+
+// Single-element structs pass as their element. The callee rebuilds the
+// struct in a private (addrspace 5) slot.
+void arg_i8_wrapper(i8_wrapper_t w) {}
+
+// CIR: cir.func {{.*}}@arg_i8_wrapper(%arg0: !s8i{{.*}})
+// CIR: cir.alloca "coerce" {{.*}} : !cir.ptr<{{.*}}, target_address_space(5)>
+// LLVM: define {{.*}}void @arg_i8_wrapper(i8 %{{.*}})
+
+void arg_i32_wrapper(i32_wrapper_t w) {}
+
+// CIR: cir.func {{.*}}@arg_i32_wrapper(%arg0: !s32i{{.*}})
+// LLVM: define {{.*}}void @arg_i32_wrapper(i32 %{{.*}})
+
+void arg_f32_wrapper(f32_wrapper_t w) {}
+
+// CIR: cir.func {{.*}}@arg_f32_wrapper(%arg0: !cir.float{{.*}})
+// LLVM: define {{.*}}void @arg_f32_wrapper(float %{{.*}})
+
+void arg_ptr_wrapper(ptr_wrapper_t w) {}
+
+// CIR: cir.func {{.*}}@arg_ptr_wrapper(%arg0: !cir.ptr<{{[^,]*}}>{{.*}})
+// LLVM: define {{.*}}void @arg_ptr_wrapper(ptr %{{.*}})
+
+// Aggregates up to 32 bits are packed into one integer register.
+void arg_char2(char2_t s) {}
+
+// CIR: cir.func {{.*}}@arg_char2(%arg0: !u16i{{.*}})
+// LLVM: define {{.*}}void @arg_char2(i16 %{{.*}})
+
+void arg_small(small_t s) {}
+
+// CIR: cir.func {{.*}}@arg_small(%arg0: !u32i{{.*}})
+// LLVM: define {{.*}}void @arg_small(i32 %{{.*}})
+
+void arg_union(int_float_t u) {}
+
+// CIR: cir.func {{.*}}@arg_union(%arg0: !u32i{{.*}})
+// LLVM: define {{.*}}void @arg_union(i32 %{{.*}})
+
+// The call site is coerced the same way as the callee.
+void call_small(small_t s) { arg_small(s); }
+
+// CIR: cir.func {{.*}}@call_small(%arg0: !u32i{{.*}})
+// CIR: cir.call @arg_small(%{{.*}}){{.*}}: (!u32i{{.*}}) -> ()
+// LLVM: define {{.*}}void @call_small(i32 %{{.*}})
+// LLVM: call void @arg_small(i32 %{{.*}})
+
+// The call site carries the same extension as the callee.
+void call_char(char c) { arg_char(c); }
+
+// LLVM: define {{.*}}void @call_char(i8 noundef signext %{{.*}})
+// LLVM: call void @arg_char(i8 noundef signext %{{.*}})
diff --git a/clang/test/CIR/CodeGen/call-conv-lowering-amdgpu-cxx-records.cpp b/clang/test/CIR/CodeGen/call-conv-lowering-amdgpu-cxx-records.cpp
new file mode 100644
index 00000000000000..e434f9d37bcfc1
--- /dev/null
+++ b/clang/test/CIR/CodeGen/call-conv-lowering-amdgpu-cxx-records.cpp
@@ -0,0 +1,74 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fclangir -fclangir-call-conv-lowering -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
+// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fclangir -fclangir-call-conv-lowering -emit-llvm %s -o %t-cir.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s
+// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
+
+// Checks that CallConvLowering classifies C++ records on ordinary AMDGPU
+// functions, where C++ rules on triviality and layout change the result.
+
+// TODO(cir): Add the C++ record cases still NYI in CallConvLowering:
+// - Non-trivial records as arguments. Classic passes them indirectly as
+// ptr addrspace(5), which the bridge cannot express yet.
+// - Member function pointers, lowered to a 128-bit pair that the bridge
+// cannot coerce yet.
+// - Records of 33 to 64 bits, such as a base plus a field, which classic
+// coerces to [2 x i32].
+
+struct Empty {};
+
+struct NonTrivialDtor {
+ int i;
+ ~NonTrivialDtor();
+};
+
+struct NonCopyable {
+ int i;
+ NonCopyable();
+ NonCopyable(const NonCopyable &) = delete;
+};
+
+struct TrivialBase { int i; };
+struct DerivedSingle : TrivialBase {};
+struct MemberPtrHolder { int TrivialBase::*p; };
+
+// A record that cannot be passed in registers returns through a generic sret.
+NonTrivialDtor ret_non_trivial_dtor() { return {1}; }
+
+// CIR: cir.func {{.*}}@_Z20ret_non_trivial_dtorv(%arg0: !cir.ptr<!rec_NonTrivialDtor> {{.*}}llvm.sret = !rec_NonTrivialDtor
+// LLVM: define {{.*}}void @_Z20ret_non_trivial_dtorv(ptr dead_on_unwind noalias writable sret(%struct.NonTrivialDtor) align 4 %{{.*}})
+
+NonCopyable ret_non_copyable() { return NonCopyable(); }
+
+// CIR: cir.func {{.*}}@_Z16ret_non_copyablev(%arg0: !cir.ptr<!rec_NonCopyable> {{.*}}llvm.sret = !rec_NonCopyable
+// LLVM: define {{.*}}void @_Z16ret_non_copyablev(ptr dead_on_unwind noalias writable sret(%struct.NonCopyable) align 4 %{{.*}})
+
+// A base subobject holding the only field still counts as a single element.
+DerivedSingle ret_derived_single() { return {}; }
+
+// CIR: cir.func {{.*}}@_Z18ret_derived_singlev() -> !s32i
+// LLVM: define {{.*}}i32 @_Z18ret_derived_singlev()
+
+void arg_derived_single(DerivedSingle d) {}
+
+// CIR: cir.func {{.*}}@_Z18arg_derived_single13DerivedSingle(%arg0: !s32i{{.*}})
+// LLVM: define {{.*}}void @_Z18arg_derived_single13DerivedSingle(i32 %{{.*}})
+
+// A data member pointer is lowered to an i64 offset before classification.
+MemberPtrHolder ret_member_ptr() { return {nullptr}; }
+
+// CIR: cir.func {{.*}}@_Z14ret_member_ptrv() -> !s64i
+// LLVM: define {{.*}}i64 @_Z14ret_member_ptrv()
+
+void arg_member_ptr(MemberPtrHolder h) {}
+
+// CIR: cir.func {{.*}}@_Z14arg_member_ptr15MemberPtrHolder(%arg0: !s64i{{.*}})
+// LLVM: define {{.*}}void @_Z14arg_member_ptr15MemberPtrHolder(i64 %{{.*}})
+
+// An empty C++ record is one byte in size but is still dropped.
+void arg_empty(Empty e) {}
+
+// CIR: cir.func {{.*}}@_Z9arg_empty5Empty()
+// LLVM: define {{.*}}void @_Z9arg_empty5Empty()
diff --git a/clang/test/CIR/CodeGen/call-conv-lowering-amdgpu-return-types.c b/clang/test/CIR/CodeGen/call-conv-lowering-amdgpu-return-types.c
new file mode 100644
index 00000000000000..e62674bd5a1702
--- /dev/null
+++ b/clang/test/CIR/CodeGen/call-conv-lowering-amdgpu-return-types.c
@@ -0,0 +1,158 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fclangir -fclangir-call-conv-lowering -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
+// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fclangir -fclangir-call-conv-lowering -emit-llvm %s -o %t-cir.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s
+// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
+
+// Checks that CallConvLowering classifies return values of ordinary AMDGPU
+// functions, which use the default convention rather than the kernel one.
+
+// TODO(cir): Add the return cases still NYI in CallConvLowering:
+// - Aggregates of 33 to 64 bits, which classic coerces to [2 x i32].
+// abiTypeToCIR has no array case yet.
+// - Aggregates over 64 bits that fit the 16-register budget. The classifier
+// returns Direct with no coerce type, which the bridge rejects.
+// - Larger aggregates, returned through sret in addrspace(5). The bridge
+// cannot express an indirect address space yet.
+// - Structs with a flexible array member. The bridge does not set that
+// record flag yet.
+// - 3-element vector returns, blocked on CIRGen's vec3 load and store NYI.
+
+typedef struct {} empty_t;
+typedef struct { char c; } i8_wrapper_t;
+typedef struct { int i; } i32_wrapper_t;
+typedef struct { float f; } f32_wrapper_t;
+typedef struct { int *p; } ptr_wrapper_t;
+typedef struct { char a, b; } char2_t;
+typedef struct { char a, b, c; } small_t;
+typedef union { int i; float f; } int_float_t;
+
+void ret_void(void) {}
+
+// CIR: cir.func {{.*}}@ret_void()
+// LLVM: define {{.*}}void @ret_void()
+
+// Sub-word integers and bool carry signext or zeroext per their signedness.
+_Bool ret_bool(void) { return 0; }
+
+// CIR: cir.func {{.*}}@ret_bool() -> (!cir.bool {{.*}}llvm.zeroext
+// LLVM: define {{.*}}zeroext i1 @ret_bool()
+
+char ret_char(void) { return 0; }
+
+// CIR: cir.func {{.*}}@ret_char() -> (!s8i {{.*}}llvm.signext
+// LLVM: define {{.*}}signext i8 @ret_char()
+
+unsigned char ret_uchar(void) { return 0; }
+
+// CIR: cir.func {{.*}}@ret_uchar() -> (!u8i {{.*}}llvm.zeroext
+// LLVM: define {{.*}}zeroext i8 @ret_uchar()
+
+short ret_short(void) { return 0; }
+
+// CIR: cir.func {{.*}}@ret_short() -> (!s16i {{.*}}llvm.signext
+// LLVM: define {{.*}}signext i16 @ret_short()
+
+// Word-sized and wider scalars and pointers return unchanged.
+int ret_int(void) { return 0; }
+
+// CIR: cir.func {{.*}}@ret_int() -> {{\(?}}!s32i
+// LLVM: define {{.*}}i32 @ret_int()
+
+long ret_long(void) { return 0; }
+
+// CIR: cir.func {{.*}}@ret_long() -> {{\(?}}!s64i
+// LLVM: define {{.*}}i64 @ret_long()
+
+float ret_float(void) { return 0; }
+
+// CIR: cir.func {{.*}}@ret_float() -> {{\(?}}!cir.float
+// LLVM: define {{.*}}float @ret_float()
+
+double ret_double(void) { return 0; }
+
+// CIR: cir.func {{.*}}@ret_double() -> {{\(?}}!cir.double
+// LLVM: define {{.*}}double @ret_double()
+
+_Float16 ret_half(void) { return 0; }
+
+// CIR: cir.func {{.*}}@ret_half() -> {{\(?}}!cir.f16
+// LLVM: define {{.*}}half @ret_half()
+
+int *ret_ptr(void) { return 0; }
+
+// CIR: cir.func {{.*}}@ret_ptr() -> {{\(?}}!cir.ptr<!s32i>
+// LLVM: define {{.*}}ptr @ret_ptr()
+_BitInt(65) ret_bitint65(void) { return 0; }
+
+// CIR: cir.func {{.*}}@ret_bitint65() -> {{\(?}}!cir.int<s, 65, bitint>
+// LLVM: define {{.*}}i65 @ret_bitint65()
+
+// An empty struct return becomes void.
+empty_t ret_empty(void) {
+ empty_t e;
+ return e;
+}
+
+// CIR: cir.func {{.*}}@ret_empty()
+// LLVM: define {{.*}}void @ret_empty()
+
+// Single-element structs return as their element.
+i8_wrapper_t ret_i8_wrapper(void) {
+ i8_wrapper_t w = {1};
+ return w;
+}
+
+// CIR: cir.func {{.*}}@ret_i8_wrapper() -> !s8i
+// LLVM: define {{.*}}i8 @ret_i8_wrapper()
+
+i32_wrapper_t ret_i32_wrapper(void) {
+ i32_wrapper_t w = {1};
+ return w;
+}
+
+// CIR: cir.func {{.*}}@ret_i32_wrapper() -> !s32i
+// LLVM: define {{.*}}i32 @ret_i32_wrapper()
+
+f32_wrapper_t ret_f32_wrapper(void) {
+ f32_wrapper_t w = {1.0f};
+ return w;
+}
+
+// CIR: cir.func {{.*}}@ret_f32_wrapper() -> !cir.float
+// LLVM: define {{.*}}float @ret_f32_wrapper()
+
+ptr_wrapper_t ret_ptr_wrapper(void) {
+ ptr_wrapper_t w = {0};
+ return w;
+}
+
+// CIR: cir.func {{.*}}@ret_ptr_wrapper() -> !cir.ptr<{{[^>]*}}>
+// LLVM: define {{.*}}ptr @ret_ptr_wrapper()
+
+// Aggregates up to 32 bits are packed into one integer register.
+char2_t ret_char2(void) {
+ char2_t s = {1, 2};
+ return s;
+}
+
+// CIR: cir.func {{.*}}@ret_char2() -> !u16i
+// LLVM: define {{.*}}i16 @ret_char2()
+
+small_t ret_small(void) {
+ small_t s = {1, 2, 3};
+ return s;
+}
+
+// CIR: cir.func {{.*}}@ret_small() -> !u32i
+// LLVM: define {{.*}}i32 @ret_small()
+
+int_float_t ret_union(void) {
+ int_float_t u = {1};
+ return u;
+}
+
+// CIR: cir.func {{.*}}@ret_union() -> !u32i
+// LLVM: define {{.*}}i32 @ret_union()
diff --git a/clang/test/CIR/CodeGenHIP/amdgcn-buffer-rsrc-type.hip b/clang/test/CIR/CodeGenHIP/amdgcn-buffer-rsrc-type.hip
index e03fa4db5c20df..0045a731c579f6 100644
--- a/clang/test/CIR/CodeGenHIP/amdgcn-buffer-rsrc-type.hip
+++ b/clang/test/CIR/CodeGenHIP/amdgcn-buffer-rsrc-type.hip
@@ -1,12 +1,16 @@
#include "../CodeGenCUDA/Inputs/cuda.h"
// REQUIRES: amdgpu-registered-target
+
+// TODO(cir): drop -fno-clangir-call-conv-lowering once CallConvLowering
+// supports the AMDGPU direct-in-registers aggregate return of a buffer
+// resource struct.
// RUN: %clang_cc1 -triple amdgpu11.00-amd-amdhsa -x hip -std=c++11 -fclangir \
-// RUN: -fcuda-is-device -emit-cir %s -o %t.cir
+// RUN: -fcuda-is-device -fno-clangir-call-conv-lowering -emit-cir %s -o %t.cir
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
// RUN: %clang_cc1 -triple amdgpu11.00-amd-amdhsa -x hip -std=c++11 -fclangir \
-// RUN: -fcuda-is-device -emit-llvm %s -o %t-cir.ll
+// RUN: -fcuda-is-device -fno-clangir-call-conv-lowering -emit-llvm %s -o %t-cir.ll
// RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s
// RUN: %clang_cc1 -triple amdgpu11.00-amd-amdhsa -x hip -std=c++11 \
diff --git a/clang/test/CIR/CodeGenHIP/call-conv-lowering-amdgpu-kernel-args.hip b/clang/test/CIR/CodeGenHIP/call-conv-lowering-amdgpu-kernel-args.hip
new file mode 100644
index 00000000000000..d713a3bd4cf7e2
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/call-conv-lowering-amdgpu-kernel-args.hip
@@ -0,0 +1,137 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -x hip -fcuda-is-device -fclangir -fclangir-call-conv-lowering -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
+// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -x hip -fcuda-is-device -fclangir -fclangir-call-conv-lowering -emit-llvm %s -o %t-cir.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s
+// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -x hip -fcuda-is-device -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
+
+// Checks that CallConvLowering classifies AMDGPU kernel arguments under the
+// kernel convention.
+
+// TODO(cir): Add multi-field aggregates, passed byref in addrspace(4), once
+// CallConvLowering supports IndirectAliased.
+
+#define __global__ __attribute__((global))
+#define __device__ __attribute__((device))
+
+typedef struct { int i; } i32_wrapper_t;
+typedef struct { i32_wrapper_t w; } nested_t;
+typedef struct { float f[1]; } f32_array1_t;
+typedef struct { float *x; } ptr_wrapper_t;
+typedef float float3 __attribute__((ext_vector_type(3)));
+typedef float float4 __attribute__((ext_vector_type(4)));
+typedef short short4 __attribute__((ext_vector_type(4)));
+
+// Word-sized scalars pass through unchanged.
+__global__ void kern_scalar(int i, float f) {}
+
+// CIR: cir.func {{.*}}@_Z11kern_scalarif(%arg0: !s32i {{.*}}, %arg1: !cir.float {{.*}}){{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z11kern_scalarif(i32 noundef %{{.*}}, float noundef %{{.*}})
+
+__global__ void kern_wide(long l, double d) {}
+
+// CIR: cir.func {{.*}}@_Z9kern_wideld(%arg0: !s64i {{.*}}, %arg1: !cir.double {{.*}}){{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z9kern_wideld(i64 noundef %{{.*}}, double noundef %{{.*}})
+
+__global__ void kern_half(_Float16 h) {}
+
+// CIR: cir.func {{.*}}@_Z9kern_halfDF16_(%arg0: !cir.f16 {{.*}}){{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z9kern_halfDF16_(half noundef %{{.*}})
+
+__global__ void kern_bitint(_BitInt(65) b) {}
+
+// CIR: cir.func {{.*}}@_Z11kern_bitintDB65_(%arg0: !cir.int<s, 65, bitint> {{.*}}){{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z11kern_bitintDB65_(i65 noundef %{{.*}})
+
+// Vectors, including 16-bit and 3-element ones, pass whole.
+__global__ void kern_float4(float4 v) {}
+
+// CIR: cir.func {{.*}}@_Z11kern_float4Dv4_f(%arg0: !cir.vector<4 x !cir.float> {{.*}}){{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z11kern_float4Dv4_f(<4 x float> noundef %{{.*}})
+
+__global__ void kern_short4(short4 v) {}
+
+// CIR: cir.func {{.*}}@_Z11kern_short4Dv4_s(%arg0: !cir.vector<4 x !s16i> {{.*}}){{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z11kern_short4Dv4_s(<4 x i16> noundef %{{.*}})
+
+__global__ void kern_float3(float3 v) {}
+
+// CIR: cir.func {{.*}}@_Z11kern_float3Dv3_f(%arg0: !cir.vector<3 x !cir.float> {{.*}}){{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z11kern_float3Dv3_f(<3 x float> noundef %{{.*}})
+
+// A pointer already in a non-generic address space keeps it.
+__global__ void kern_global_ptr(__attribute__((address_space(1))) int *p) {}
+
+// CIR: cir.func {{.*}}@_Z15kern_global_ptrPU3AS1i(%arg0: !cir.ptr<!s32i, target_address_space(1)> {{.*}}){{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z15kern_global_ptrPU3AS1i(ptr addrspace(1) noundef %{{.*}})
+
+// HIP passes a generic pointer or reference as a global pointer and casts it
+// back to generic in the kernel.
+__global__ void kern_ptr(int *p) {}
+
+// CIR: cir.func {{.*}}@_Z8kern_ptrPi(%[[P:[^:]+]]: !cir.ptr<!s32i, target_address_space(1)> {{.*}}){{.*}}cc(amdgpu_kernel)
+// CIR: cir.cast address_space %[[P]] : !cir.ptr<!s32i, target_address_space(1)> -> !cir.ptr<!s32i>
+// LLVM: define {{.*}}amdgpu_kernel void @_Z8kern_ptrPi(ptr addrspace(1) noundef %{{.*}})
+
+__global__ void kern_ref(int &r) {}
+
+// CIR: cir.func {{.*}}@_Z8kern_refRi(%{{.*}}: !cir.ptr<!s32i, target_address_space(1)> {{.*}}){{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z8kern_refRi(ptr addrspace(1) noundef nonnull align 4 dereferenceable(4) %{{.*}})
+
+// A struct holding only a pointer is unwrapped first, so it is coerced too.
+__global__ void kern_ptr_wrap(ptr_wrapper_t w) {}
+
+// CIR: cir.func {{.*}}@_Z13kern_ptr_wrap13ptr_wrapper_t(%{{.*}}: !cir.ptr<{{.*}}, target_address_space(1)>{{.*}}){{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z13kern_ptr_wrap13ptr_wrapper_t(ptr addrspace(1) %{{.*}})
+
+// A device function keeps the generic pointer.
+__device__ void dev_ptr(int *p) {}
+
+// CIR: cir.func {{.*}}@_Z7dev_ptrPi(%{{.*}}: !cir.ptr<!s32i> {{.*}})
+// LLVM: define {{.*}}void @_Z7dev_ptrPi(ptr noundef %{{.*}})
+
+// A single-element struct kernel argument is passed as its element.
+__global__ void kern_wrapper(i32_wrapper_t w) {}
+
+// CIR: cir.func {{.*}}@_Z12kern_wrapper13i32_wrapper_t(%arg0: !s32i{{.*}}){{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z12kern_wrapper13i32_wrapper_t(i32 %{{.*}})
+
+// The element is found through nested structs and one-element arrays.
+__global__ void kern_nested(nested_t n) {}
+
+// CIR: cir.func {{.*}}@_Z11kern_nested8nested_t(%arg0: !s32i{{.*}}){{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z11kern_nested8nested_t(i32 %{{.*}})
+
+__global__ void kern_array1(f32_array1_t a) {}
+
+// CIR: cir.func {{.*}}@_Z11kern_array112f32_array1_t(%arg0: !cir.float{{.*}}){{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z11kern_array112f32_array1_t(float %{{.*}})
+
+// Kernel arguments are read from the kernarg segment, so sub-word integers
+// and bool are not extended, unlike in a device function.
+__global__ void kern_char(char c) {}
+
+// CIR: cir.func {{.*}}@_Z9kern_charc(%arg0: !s8i {llvm.noundef}{{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z9kern_charc(i8 noundef %{{.*}})
+
+__global__ void kern_uchar(unsigned char c) {}
+
+// CIR: cir.func {{.*}}@_Z10kern_ucharh(%arg0: !u8i {llvm.noundef}{{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z10kern_ucharh(i8 noundef %{{.*}})
+
+__global__ void kern_short(short s) {}
+
+// CIR: cir.func {{.*}}@_Z10kern_shorts(%arg0: !s16i {llvm.noundef}{{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z10kern_shorts(i16 noundef %{{.*}})
+
+__global__ void kern_bool(bool b) {}
+
+// CIR: cir.func {{.*}}@_Z9kern_boolb(%arg0: !cir.bool {llvm.noundef}{{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z9kern_boolb(i1 noundef %{{.*}})
+
+// The same argument to a device function is extended.
+__device__ void dev_char(char c) {}
+
+// CIR: cir.func {{.*}}@_Z8dev_charc(%arg0: !s8i {{.*}}llvm.signext
+// LLVM: define {{.*}}void @_Z8dev_charc(i8 noundef signext %{{.*}})
diff --git a/clang/test/CIR/CodeGenHIP/cleanup-alloca-addrspace.hip b/clang/test/CIR/CodeGenHIP/cleanup-alloca-addrspace.hip
index c55cd55a98d675..7cbb8ff8301ec0 100644
--- a/clang/test/CIR/CodeGenHIP/cleanup-alloca-addrspace.hip
+++ b/clang/test/CIR/CodeGenHIP/cleanup-alloca-addrspace.hip
@@ -7,9 +7,10 @@
// correct address space after flatten cfg. This is important for address-space
// aware targets like amdgpu.
-// CIR-FLAT-LABEL: cir.func {{.*}} @_Z1fv
+// S is not trivially destructible, so f returns it through a generic sret
+// pointer rather than a __retval alloca.
+// CIR-FLAT-LABEL: cir.func {{.*}} @_Z1fv(%arg0: !cir.ptr<!rec_S> {{.*}}llvm.sret = !rec_S
// CIR-FLAT: cir.alloca "__cleanup_dest_slot" {{.*}} : !cir.ptr<!s32i, target_address_space(5)>
-// CIR-FLAT: cir.alloca "__retval" {{.*}} : !cir.ptr<!rec_S, target_address_space(5)>
// CIR-FLAT: cir.alloca "nrvo" {{.*}} : !cir.ptr<!cir.bool, target_address_space(5)>
// CIR-FLAT-LABEL: cir.func {{.*}} @_ZN1SD1Ev
@@ -18,9 +19,8 @@
// CIR-FLAT-LABEL: cir.func {{.*}} @_ZN1SD2Ev
// CIR-FLAT: cir.alloca "this" {{.*}} : !cir.ptr<!cir.ptr<!rec_S>, target_address_space(5)>
-// LLVM-LABEL: define {{.*}} @_Z1fv
+// LLVM-LABEL: define {{.*}} void @_Z1fv(ptr {{.*}}sret(%struct.S)
// LLVM: alloca i32, align 4, addrspace(5)
-// LLVM: alloca %struct.S, align 1, addrspace(5)
// LLVM: alloca i8, align 1, addrspace(5)
// LLVM-LABEL: define {{.*}} @_ZN1SD1Ev
diff --git a/clang/test/CIR/CodeGenHIP/lambda-reference-var.hip b/clang/test/CIR/CodeGenHIP/lambda-reference-var.hip
index 5b3ad34df9a2a8..60c94010940c8a 100644
--- a/clang/test/CIR/CodeGenHIP/lambda-reference-var.hip
+++ b/clang/test/CIR/CodeGenHIP/lambda-reference-var.hip
@@ -15,11 +15,15 @@
// RUN: %s -o %t-host.ll
// RUN: FileCheck --check-prefix=HOST --input-file=%t-host.ll %s
+// TODO(cir): Drop -fno-clangir-call-conv-lowering once kernel byref arguments
+// are supported.
// RUN: %clang_cc1 -x hip -std=c++11 -triple amdgcn-amd-amdhsa -fcuda-is-device \
-// RUN: -fclangir -emit-cir -I%S/../CodeGenCUDA/Inputs %s -o %t-dev.cir
+// RUN: -fclangir -fno-clangir-call-conv-lowering -emit-cir \
+// RUN: -I%S/../CodeGenCUDA/Inputs %s -o %t-dev.cir
// RUN: FileCheck --check-prefix=CIR-DEV --input-file=%t-dev.cir %s
// RUN: %clang_cc1 -x hip -std=c++11 -triple amdgcn-amd-amdhsa -fcuda-is-device \
-// RUN: -fclangir -emit-llvm -I%S/../CodeGenCUDA/Inputs %s -o %t-dev-cir.ll
+// RUN: -fclangir -fno-clangir-call-conv-lowering -emit-llvm \
+// RUN: -I%S/../CodeGenCUDA/Inputs %s -o %t-dev-cir.ll
// RUN: FileCheck --check-prefix=DEV --input-file=%t-dev-cir.ll %s
// RUN: %clang_cc1 -x hip -std=c++11 -triple amdgcn-amd-amdhsa -fcuda-is-device \
// RUN: -emit-llvm -I%S/../CodeGenCUDA/Inputs %s -o %t-dev.ll
diff --git a/clang/test/CIR/CodeGenSYCL/kernel-caller-entry-point.cpp b/clang/test/CIR/CodeGenSYCL/kernel-caller-entry-point.cpp
index cf9a923d466144..4b0d591bde61b1 100644
--- a/clang/test/CIR/CodeGenSYCL/kernel-caller-entry-point.cpp
+++ b/clang/test/CIR/CodeGenSYCL/kernel-caller-entry-point.cpp
@@ -17,9 +17,11 @@
// On AMDGPU and NVPTX, the kernel caller entry point uses the target's device
// kernel calling convention.
-// RUN: %clang_cc1 -std=c++20 -fsycl-is-device -triple amdgpu-amd-amdhsa -fclangir -emit-cir %s -o %t-amdgcn.cir
+// TODO(cir): drop -fno-clangir-call-conv-lowering once CallConvLowering
+// supports byref (IndirectAliased) kernel arguments on AMDGPU.
+// RUN: %clang_cc1 -std=c++20 -fsycl-is-device -triple amdgpu-amd-amdhsa -fclangir -fno-clangir-call-conv-lowering -emit-cir %s -o %t-amdgcn.cir
// RUN: FileCheck --input-file=%t-amdgcn.cir %s -check-prefix=CIR-AMDGCN
-// RUN: %clang_cc1 -std=c++20 -fsycl-is-device -triple amdgpu-amd-amdhsa -fclangir -emit-llvm %s -o %t-amdgcn-cir.ll
+// RUN: %clang_cc1 -std=c++20 -fsycl-is-device -triple amdgpu-amd-amdhsa -fclangir -fno-clangir-call-conv-lowering -emit-llvm %s -o %t-amdgcn-cir.ll
// RUN: FileCheck --input-file=%t-amdgcn-cir.ll %s -check-prefix=LLVM-OGCG-AMDGCN
// RUN: %clang_cc1 -std=c++20 -fsycl-is-device -triple amdgpu-amd-amdhsa -emit-llvm %s -o %t-amdgcn.ll
// RUN: FileCheck --input-file=%t-amdgcn.ll %s -check-prefix=LLVM-OGCG-AMDGCN
More information about the llvm-branch-commits
mailing list