shylie/llvm-project
1
0
Fork 0
Code Issues Pull Requests Actions 6 Packages Projects Releases Wiki Activity
llvm-project/clang/test/CodeGenHIP/amdgcnspirv-uses-amdgpu-abi.cpp
Alex Voicu 68fea00dfb
[SPIRV] Use AMDGPU ABI for AMDGCN flavoured SPIRV (#169865) 
At the moment AMDGCN flavoured SPIRV uses the SPIRV ABI with some tweaks
revolving around passing aggregates as direct. This is problematic in
multiple ways:

- it leads to divergence from code compiled for a concrete target, which
makes it difficult to debug;
- it incurs a run time cost, when dealing with larger aggregates;
- it incurs a compile time cost, when dealing with larger aggregates.

This patch switches over AMDGCN flavoured SPIRV to implement the AMDGPU
ABI (except for dealing with variadic functions, which will be added in
the future). One additional complication (and the primary motivation
behind the current less than ideal state of affairs) stems from `byref`,
which AMDGPU uses, not being expressible in SPIR-V. We deal with this by
CodeGen-ing for `byref`, lowering it to the `FuncParamAttr ByVal` in
SPIR-V, and restoring it when doing reverse translation from AMDGCN
flavoured SPIR-V.
2025-12-07 17:15:48 +00:00

322 lines
13 KiB
C++
Raw Blame History

// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -emit-llvm -fcuda-is-device -O3 \
// RUN: -o - %s | FileCheck --check-prefix=AMDGCNSPIRV %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -target-cpu gfx906 -emit-llvm -fcuda-is-device -O3 \
// RUN: -o - %s | FileCheck --check-prefix=AMDGPU %s
#define __global__ __attribute__((global))
#define __device__ __attribute__((device))
union Transparent { unsigned x; };
using V1 = unsigned __attribute__((ext_vector_type(1)));
using V2 = unsigned __attribute__((ext_vector_type(2)));
using V3 = unsigned __attribute__((ext_vector_type(3)));
using V4 = unsigned __attribute__((ext_vector_type(4)));
struct SingleElement { unsigned x; };
struct ByRef { unsigned x[17]; };
// AMDGCNSPIRV-LABEL: define spir_kernel void @_Z2k0s(
// AMDGCNSPIRV-SAME: i16 noundef [[TMP0:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR0:[0-9]+]] !max_work_group_size [[META9:![0-9]+]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: ret void
//
// AMDGPU-LABEL: define dso_local amdgpu_kernel void @_Z2k0s(
// AMDGPU-SAME: i16 noundef [[TMP0:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: ret void
//
__global__ void k0(short) { }
// AMDGCNSPIRV-LABEL: define spir_kernel void @_Z2k1j(
// AMDGCNSPIRV-SAME: i32 noundef [[TMP0:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR0]] !max_work_group_size [[META9]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: ret void
//
// AMDGPU-LABEL: define dso_local amdgpu_kernel void @_Z2k1j(
// AMDGPU-SAME: i32 noundef [[TMP0:%.*]]) local_unnamed_addr #[[ATTR0]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: ret void
//
__global__ void k1(unsigned) { }
// AMDGCNSPIRV-LABEL: define spir_kernel void @_Z2k2d(
// AMDGCNSPIRV-SAME: double noundef [[TMP0:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR0]] !max_work_group_size [[META9]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: ret void
//
// AMDGPU-LABEL: define dso_local amdgpu_kernel void @_Z2k2d(
// AMDGPU-SAME: double noundef [[TMP0:%.*]]) local_unnamed_addr #[[ATTR0]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: ret void
//
__global__ void k2(double) { }
// AMDGCNSPIRV-LABEL: define spir_kernel void @_Z2k311Transparent(
// AMDGCNSPIRV-SAME: i32 [[DOTCOERCE:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR0]] !max_work_group_size [[META9]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: ret void
//
// AMDGPU-LABEL: define dso_local amdgpu_kernel void @_Z2k311Transparent(
// AMDGPU-SAME: i32 [[DOTCOERCE:%.*]]) local_unnamed_addr #[[ATTR0]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: ret void
//
__global__ void k3(Transparent) { }
// AMDGCNSPIRV-LABEL: define spir_kernel void @_Z2k413SingleElement(
// AMDGCNSPIRV-SAME: i32 [[DOTCOERCE:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR0]] !max_work_group_size [[META9]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: ret void
//
// AMDGPU-LABEL: define dso_local amdgpu_kernel void @_Z2k413SingleElement(
// AMDGPU-SAME: i32 [[DOTCOERCE:%.*]]) local_unnamed_addr #[[ATTR0]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: ret void
//
__global__ void k4(SingleElement) { }
// AMDGCNSPIRV-LABEL: define spir_kernel void @_Z2k55ByRef(
// AMDGCNSPIRV-SAME: ptr addrspace(2) noundef readnone byref([[STRUCT_BYREF:%.*]]) align 4 captures(none) [[TMP0:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR0]] !max_work_group_size [[META9]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: ret void
//
// AMDGPU-LABEL: define dso_local amdgpu_kernel void @_Z2k55ByRef(
// AMDGPU-SAME: ptr addrspace(4) noundef readnone byref([[STRUCT_BYREF:%.*]]) align 4 captures(none) [[TMP0:%.*]]) local_unnamed_addr #[[ATTR0]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: ret void
//
__global__ void k5(ByRef) { }
// AMDGCNSPIRV-LABEL: define spir_kernel void @_Z2k6Dv1_jDv2_jDv3_jDv4_j(
// AMDGCNSPIRV-SAME: <1 x i32> noundef [[TMP0:%.*]], <2 x i32> noundef [[TMP1:%.*]], <3 x i32> noundef [[TMP2:%.*]], <4 x i32> noundef [[TMP3:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR0]] !max_work_group_size [[META9]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: ret void
//
// AMDGPU-LABEL: define dso_local amdgpu_kernel void @_Z2k6Dv1_jDv2_jDv3_jDv4_j(
// AMDGPU-SAME: <1 x i32> noundef [[TMP0:%.*]], <2 x i32> noundef [[TMP1:%.*]], <3 x i32> noundef [[TMP2:%.*]], <4 x i32> noundef [[TMP3:%.*]]) local_unnamed_addr #[[ATTR0]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: ret void
//
__global__ void k6(V1, V2, V3, V4) { }
// AMDGCNSPIRV-LABEL: define spir_kernel void @_Z2k7Pj(
// AMDGCNSPIRV-SAME: ptr addrspace(1) noundef readnone captures(none) [[DOTCOERCE:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR0]] !max_work_group_size [[META9]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: ret void
//
// AMDGPU-LABEL: define dso_local amdgpu_kernel void @_Z2k7Pj(
// AMDGPU-SAME: ptr addrspace(1) noundef readnone captures(none) [[DOTCOERCE:%.*]]) local_unnamed_addr #[[ATTR0]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: ret void
//
__global__ void k7(unsigned*) { }
// AMDGCNSPIRV-LABEL: define spir_func void @_Z2f0s(
// AMDGCNSPIRV-SAME: i16 noundef signext [[TMP0:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR1:[0-9]+]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: ret void
//
// AMDGPU-LABEL: define dso_local void @_Z2f0s(
// AMDGPU-SAME: i16 noundef signext [[TMP0:%.*]]) local_unnamed_addr #[[ATTR1:[0-9]+]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: ret void
//
__device__ void f0(short) { }
// AMDGCNSPIRV-LABEL: define spir_func void @_Z2f1j(
// AMDGCNSPIRV-SAME: i32 noundef [[TMP0:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR1]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: ret void
//
// AMDGPU-LABEL: define dso_local void @_Z2f1j(
// AMDGPU-SAME: i32 noundef [[TMP0:%.*]]) local_unnamed_addr #[[ATTR1]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: ret void
//
__device__ void f1(unsigned) { }
// AMDGCNSPIRV-LABEL: define spir_func void @_Z2f2d(
// AMDGCNSPIRV-SAME: double noundef [[TMP0:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR1]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: ret void
//
// AMDGPU-LABEL: define dso_local void @_Z2f2d(
// AMDGPU-SAME: double noundef [[TMP0:%.*]]) local_unnamed_addr #[[ATTR1]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: ret void
//
__device__ void f2(double) { }
// AMDGCNSPIRV-LABEL: define spir_func void @_Z2f311Transparent(
// AMDGCNSPIRV-SAME: i32 [[DOTCOERCE:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR1]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: ret void
//
// AMDGPU-LABEL: define dso_local void @_Z2f311Transparent(
// AMDGPU-SAME: i32 [[DOTCOERCE:%.*]]) local_unnamed_addr #[[ATTR1]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: ret void
//
__device__ void f3(Transparent) { }
// AMDGCNSPIRV-LABEL: define spir_func void @_Z2f413SingleElement(
// AMDGCNSPIRV-SAME: i32 [[DOTCOERCE:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR1]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: ret void
//
// AMDGPU-LABEL: define dso_local void @_Z2f413SingleElement(
// AMDGPU-SAME: i32 [[DOTCOERCE:%.*]]) local_unnamed_addr #[[ATTR1]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: ret void
//
__device__ void f4(SingleElement) { }
// AMDGCNSPIRV-LABEL: define spir_func void @_Z2f55ByRef(
// AMDGCNSPIRV-SAME: ptr noundef readnone byref([[STRUCT_BYREF:%.*]]) align 4 captures(none) [[TMP0:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR1]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: ret void
//
// AMDGPU-LABEL: define dso_local void @_Z2f55ByRef(
// AMDGPU-SAME: ptr addrspace(5) noundef readnone byref([[STRUCT_BYREF:%.*]]) align 4 captures(none) [[TMP0:%.*]]) local_unnamed_addr #[[ATTR1]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: ret void
//
__device__ void f5(ByRef) { }
// AMDGCNSPIRV-LABEL: define spir_func void @_Z2f6Dv1_jDv2_jDv3_jDv4_j(
// AMDGCNSPIRV-SAME: <1 x i32> noundef [[TMP0:%.*]], <2 x i32> noundef [[TMP1:%.*]], <3 x i32> noundef [[TMP2:%.*]], <4 x i32> noundef [[TMP3:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR1]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: ret void
//
// AMDGPU-LABEL: define dso_local void @_Z2f6Dv1_jDv2_jDv3_jDv4_j(
// AMDGPU-SAME: <1 x i32> noundef [[TMP0:%.*]], <2 x i32> noundef [[TMP1:%.*]], <3 x i32> noundef [[TMP2:%.*]], <4 x i32> noundef [[TMP3:%.*]]) local_unnamed_addr #[[ATTR1]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: ret void
//
__device__ void f6(V1, V2, V3, V4) { }
// AMDGCNSPIRV-LABEL: define spir_func noundef signext i16 @_Z2f7v(
// AMDGCNSPIRV-SAME: ) local_unnamed_addr addrspace(4) #[[ATTR1]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: ret i16 0
//
// AMDGPU-LABEL: define dso_local noundef signext i16 @_Z2f7v(
// AMDGPU-SAME: ) local_unnamed_addr #[[ATTR1]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: ret i16 0
//
__device__ short f7() { return 0; }
// AMDGCNSPIRV-LABEL: define spir_func noundef i32 @_Z2f8v(
// AMDGCNSPIRV-SAME: ) local_unnamed_addr addrspace(4) #[[ATTR1]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: ret i32 0
//
// AMDGPU-LABEL: define dso_local noundef i32 @_Z2f8v(
// AMDGPU-SAME: ) local_unnamed_addr #[[ATTR1]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: ret i32 0
//
__device__ unsigned f8() { return 0; }
// AMDGCNSPIRV-LABEL: define spir_func noundef double @_Z2f9v(
// AMDGCNSPIRV-SAME: ) local_unnamed_addr addrspace(4) #[[ATTR1]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: ret double 0.000000e+00
//
// AMDGPU-LABEL: define dso_local noundef double @_Z2f9v(
// AMDGPU-SAME: ) local_unnamed_addr #[[ATTR1]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: ret double 0.000000e+00
//
__device__ double f9() { return 0.; }
// AMDGCNSPIRV-LABEL: define spir_func noundef i32 @_Z3f10v(
// AMDGCNSPIRV-SAME: ) local_unnamed_addr addrspace(4) #[[ATTR1]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: ret i32 0
//
// AMDGPU-LABEL: define dso_local noundef i32 @_Z3f10v(
// AMDGPU-SAME: ) local_unnamed_addr #[[ATTR1]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: ret i32 0
//
__device__ Transparent f10() { return {}; }
// AMDGCNSPIRV-LABEL: define spir_func noundef i32 @_Z3f11v(
// AMDGCNSPIRV-SAME: ) local_unnamed_addr addrspace(4) #[[ATTR1]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: ret i32 0
//
// AMDGPU-LABEL: define dso_local noundef i32 @_Z3f11v(
// AMDGPU-SAME: ) local_unnamed_addr #[[ATTR1]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: ret i32 0
//
__device__ SingleElement f11() { return {}; }
// AMDGCNSPIRV-LABEL: define spir_func void @_Z3f12v(
// AMDGCNSPIRV-SAME: ptr dead_on_unwind noalias writable writeonly sret([[STRUCT_BYREF:%.*]]) align 4 captures(none) initializes((0, 68)) [[AGG_RESULT:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR2:[0-9]+]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: tail call addrspace(4) void @llvm.memset.p0.i64(ptr noundef nonnull align 4 dereferenceable(68) [[AGG_RESULT]], i8 0, i64 68, i1 false)
// AMDGCNSPIRV-NEXT: ret void
//
// AMDGPU-LABEL: define dso_local void @_Z3f12v(
// AMDGPU-SAME: ptr addrspace(5) dead_on_unwind noalias writable writeonly sret([[STRUCT_BYREF:%.*]]) align 4 captures(none) initializes((0, 68)) [[AGG_RESULT:%.*]]) local_unnamed_addr #[[ATTR2:[0-9]+]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: tail call void @llvm.memset.p5.i64(ptr addrspace(5) noundef align 4 dereferenceable(68) [[AGG_RESULT]], i8 0, i64 68, i1 false)
// AMDGPU-NEXT: ret void
//
__device__ ByRef f12() { return {}; }
// AMDGCNSPIRV-LABEL: define spir_func noundef <1 x i32> @_Z3f13v(
// AMDGCNSPIRV-SAME: ) local_unnamed_addr addrspace(4) #[[ATTR1]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: ret <1 x i32> zeroinitializer
//
// AMDGPU-LABEL: define dso_local noundef <1 x i32> @_Z3f13v(
// AMDGPU-SAME: ) local_unnamed_addr #[[ATTR1]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: ret <1 x i32> zeroinitializer
//
__device__ V1 f13() { return {}; }
// AMDGCNSPIRV-LABEL: define spir_func noundef <2 x i32> @_Z3f14v(
// AMDGCNSPIRV-SAME: ) local_unnamed_addr addrspace(4) #[[ATTR1]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: ret <2 x i32> zeroinitializer
//
// AMDGPU-LABEL: define dso_local noundef <2 x i32> @_Z3f14v(
// AMDGPU-SAME: ) local_unnamed_addr #[[ATTR1]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: ret <2 x i32> zeroinitializer
//
__device__ V2 f14() { return {}; }
// AMDGCNSPIRV-LABEL: define spir_func noundef <3 x i32> @_Z3f15v(
// AMDGCNSPIRV-SAME: ) local_unnamed_addr addrspace(4) #[[ATTR1]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: ret <3 x i32> zeroinitializer
//
// AMDGPU-LABEL: define dso_local noundef <3 x i32> @_Z3f15v(
// AMDGPU-SAME: ) local_unnamed_addr #[[ATTR1]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: ret <3 x i32> zeroinitializer
//
__device__ V3 f15() { return {}; }
// AMDGCNSPIRV-LABEL: define spir_func noundef <4 x i32> @_Z3f16v(
// AMDGCNSPIRV-SAME: ) local_unnamed_addr addrspace(4) #[[ATTR1]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: ret <4 x i32> zeroinitializer
//
// AMDGPU-LABEL: define dso_local noundef <4 x i32> @_Z3f16v(
// AMDGPU-SAME: ) local_unnamed_addr #[[ATTR1]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: ret <4 x i32> zeroinitializer
//
__device__ V4 f16() { return {}; }
//.
// AMDGCNSPIRV: [[META9]] = !{i32 1024, i32 1, i32 1}
//.
Reference in New Issue View Git Blame Copy Permalink