Author: Acim Maravic Date: 2026-10-21T10:55:32+02:00 New Revision: 816cbba1fb9d3937a0ffbe6feffb16a5355c8a04
URL: https://github.com/llvm/llvm-project/commit/816cbba1fb9d3937a0ffbe6feffb16a5355c8a04 DIFF: https://github.com/llvm/llvm-project/commit/816cbba1fb9d3937a0ffbe6feffb16a5355c8a04.diff LOG: [AMDGPU] Fix builtins using long int for 64-bit operands (#231500) 'long int' is 32-bit on Windows (LLP64) and 64-bit on Linux (LP64), and HIP device compilation inherits the host integer widths via AMDGPUTargetInfo::setAuxTarget(). Builtins declared with 'long int' but mapped to intrinsics taking i64 therefore produce an invalid 'bitcast i32 to i64' on Windows hosts. Added: clang/test/CodeGenHIP/builtins-amdgcn-gfx1250-int64.hip Modified: clang/include/clang/Basic/BuiltinsAMDGPU.td Removed: ################################################################################ diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td index 34a1bd2d301a3..5a0b71fa50c19 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.td +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td @@ -1036,8 +1036,8 @@ def __builtin_amdgcn_global_store_async_from_lds_b8 : AMDGPUBuiltin<"void(char a def __builtin_amdgcn_global_store_async_from_lds_b32 : AMDGPUBuiltin<"void(int address_space<1> *, int address_space<3> *, _Constant int, _Constant int)", [Const], "async-store-from-lds-insts">; def __builtin_amdgcn_global_store_async_from_lds_b64 : AMDGPUBuiltin<"void(_ExtVector<2, int> address_space<1> *, _ExtVector<2, int> address_space<3> *, _Constant int, _Constant int)", [Const], "async-store-from-lds-insts">; def __builtin_amdgcn_global_store_async_from_lds_b128 : AMDGPUBuiltin<"void(_ExtVector<4, int> address_space<1> *, _ExtVector<4, int> address_space<3> *, _Constant int, _Constant int)", [Const], "async-store-from-lds-insts">; -def __builtin_amdgcn_ds_atomic_async_barrier_arrive_b64 : AMDGPUBuiltin<"void(long int address_space<3> *)", [Const], "gfx1250-insts">; -def __builtin_amdgcn_ds_atomic_barrier_arrive_rtn_b64 : AMDGPUBuiltin<"long int(long int address_space<3> *, long int)", [Const], "gfx1250-insts">; +def __builtin_amdgcn_ds_atomic_async_barrier_arrive_b64 : AMDGPUBuiltin<"void(int64_t address_space<3> *)", [Const], "gfx1250-insts">; +def __builtin_amdgcn_ds_atomic_barrier_arrive_rtn_b64 : AMDGPUBuiltin<"int64_t(int64_t address_space<3> *, int64_t)", [Const], "gfx1250-insts">; def __builtin_amdgcn_tensor_load_to_lds : AMDGPUBuiltin<"void(_ExtVector<4, unsigned int>, _ExtVector<8, unsigned int>, _ExtVector<4, unsigned int>, _ExtVector<4, unsigned int>, _ExtVector<8, unsigned int>, _Constant int)", [Const], "tdm-insts"> { let Documentation = [DocTensorLoadToLDS_GFX1250]; @@ -1171,8 +1171,8 @@ def __builtin_amdgcn_permlane_xor : AMDGPUBuiltin<"int(int, int, int)", [Const], def __builtin_amdgcn_permlane_idx_gen : AMDGPUBuiltin<"int(int, int)", [Const], "gfx1250-insts,wavefrontsize32">; def __builtin_amdgcn_perm_pk16_b4_u4 : AMDGPUBuiltin<"_ExtVector<2, unsigned int>(unsigned int, unsigned int, _ExtVector<2, unsigned int>)", [Const], "tensor-cvt-lut-insts">; -def __builtin_amdgcn_perm_pk16_b6_u4 : AMDGPUBuiltin<"_ExtVector<3, unsigned int>(unsigned int, unsigned long int, _ExtVector<2, unsigned int>)", [Const], "tensor-cvt-lut-insts">; -def __builtin_amdgcn_perm_pk16_b8_u4 : AMDGPUBuiltin<"_ExtVector<4, unsigned int>(unsigned long int, unsigned long int, _ExtVector<2, unsigned int>)", [Const], "tensor-cvt-lut-insts">; +def __builtin_amdgcn_perm_pk16_b6_u4 : AMDGPUBuiltin<"_ExtVector<3, unsigned int>(unsigned int, uint64_t, _ExtVector<2, unsigned int>)", [Const], "tensor-cvt-lut-insts">; +def __builtin_amdgcn_perm_pk16_b8_u4 : AMDGPUBuiltin<"_ExtVector<4, unsigned int>(uint64_t, uint64_t, _ExtVector<2, unsigned int>)", [Const], "tensor-cvt-lut-insts">; def __builtin_amdgcn_add_max_i32 : AMDGPUBuiltin<"int(int, int, int, _Constant bool)", [Const], "add-min-max-insts">; def __builtin_amdgcn_add_max_u32 : AMDGPUBuiltin<"unsigned int(unsigned int, unsigned int, unsigned int, _Constant bool)", [Const], "add-min-max-insts">; @@ -1276,7 +1276,7 @@ def __builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4 : AMDGPUBuiltin<"_ExtVector let Documentation = [DocWMMA_scale_GFX1250]; let ArgNames = ["matrix_a_fmt", "a", "matrix_b_fmt", "b", "c_mod", "c", "matrix_a_scale", "matrix_a_scale_fmt", "matrix_a_scale_exp", "matrix_b_scale", "matrix_b_scale_fmt", "matrix_b_scale_exp", "matrix_a_reuse", "matrix_b_reuse"]; } -def __builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4 : AMDGPUBuiltin<"_ExtVector<8, float>(_Constant int, _ExtVector<16, int>, _Constant int, _ExtVector<16, int>, _Constant short, _ExtVector<8, float>, _Constant int, _Constant int, long int, _Constant int, _Constant int, long int, _Constant bool, _Constant bool)", [Const], "wmma-n16-insts,wavefrontsize32"> { +def __builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4 : AMDGPUBuiltin<"_ExtVector<8, float>(_Constant int, _ExtVector<16, int>, _Constant int, _ExtVector<16, int>, _Constant short, _ExtVector<8, float>, _Constant int, _Constant int, int64_t, _Constant int, _Constant int, int64_t, _Constant bool, _Constant bool)", [Const], "wmma-n16-insts,wavefrontsize32"> { let Documentation = [DocWMMA_scale16_GFX1250]; let ArgNames = ["matrix_a_fmt", "a", "matrix_b_fmt", "b", "c_mod", "c", "matrix_a_scale", "matrix_a_scale_fmt", "matrix_a_scale_exp", "matrix_b_scale", "matrix_b_scale_fmt", "matrix_b_scale_exp", "matrix_a_reuse", "matrix_b_reuse"]; } @@ -1296,7 +1296,7 @@ def __builtin_amdgcn_wmma_scale_f32_32x16x128_f4 : AMDGPUBuiltin<"_ExtVector<16, let Documentation = [DocWMMA_scale_GFX1250]; let ArgNames = ["a", "b", "c_mod", "c", "matrix_a_scale", "matrix_a_scale_fmt", "matrix_a_scale_exp", "matrix_b_scale", "matrix_b_scale_fmt", "matrix_b_scale_exp", "matrix_a_reuse", "matrix_b_reuse"]; } -def __builtin_amdgcn_wmma_scale16_f32_32x16x128_f4 : AMDGPUBuiltin<"_ExtVector<16, float>(_ExtVector<16, int>, _ExtVector<8, int>, _Constant short, _ExtVector<16, float>, _Constant int, _Constant int, long int, _Constant int, _Constant int, long int, _Constant bool, _Constant bool)", [Const], "wmma-f4-insts,wavefrontsize32"> { +def __builtin_amdgcn_wmma_scale16_f32_32x16x128_f4 : AMDGPUBuiltin<"_ExtVector<16, float>(_ExtVector<16, int>, _ExtVector<8, int>, _Constant short, _ExtVector<16, float>, _Constant int, _Constant int, int64_t, _Constant int, _Constant int, int64_t, _Constant bool, _Constant bool)", [Const], "wmma-f4-insts,wavefrontsize32"> { let Documentation = [DocWMMA_scale16_f4_GFX1250]; let ArgNames = ["a", "b", "c_mod", "c", "matrix_a_scale", "matrix_a_scale_fmt", "matrix_a_scale_exp", "matrix_b_scale", "matrix_b_scale_fmt", "matrix_b_scale_exp", "matrix_a_reuse", "matrix_b_reuse"]; } diff --git a/clang/test/CodeGenHIP/builtins-amdgcn-gfx1250-int64.hip b/clang/test/CodeGenHIP/builtins-amdgcn-gfx1250-int64.hip new file mode 100644 index 0000000000000..286958632ce4a --- /dev/null +++ b/clang/test/CodeGenHIP/builtins-amdgcn-gfx1250-int64.hip @@ -0,0 +1,47 @@ +// RUN: %clang_cc1 -triple amdgpu12.50-amd-amdhsa -aux-triple x86_64-unknown-linux-gnu -x hip -emit-llvm -fcuda-is-device -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgpu12.50-amd-amdhsa -aux-triple x86_64-pc-windows-msvc -x hip -emit-llvm -fcuda-is-device -o - %s | FileCheck %s + +#define __device__ __attribute__((device)) + +typedef unsigned int u32x2 __attribute__((ext_vector_type(2))); +typedef unsigned int u32x3 __attribute__((ext_vector_type(3))); +typedef unsigned int u32x4 __attribute__((ext_vector_type(4))); +typedef int v8i __attribute__((ext_vector_type(8))); +typedef int v16i __attribute__((ext_vector_type(16))); +typedef float v8f __attribute__((ext_vector_type(8))); +typedef float v16f __attribute__((ext_vector_type(16))); + +// CHECK-LABEL: @_Z20test_perm_pk16_b6_u4jyDv2_j +// CHECK: call <3 x i32> @llvm.amdgcn.perm.pk16.b6.u4(i32 %{{.*}}, i64 %{{.*}}, <2 x i32> %{{.*}}) +__device__ u32x3 test_perm_pk16_b6_u4(unsigned int a, unsigned long long b, u32x2 c) { + return __builtin_amdgcn_perm_pk16_b6_u4(a, b, c); +} +// CHECK-LABEL: @_Z20test_perm_pk16_b8_u4yyDv2_j +// CHECK: call <4 x i32> @llvm.amdgcn.perm.pk16.b8.u4(i64 %{{.*}}, i64 %{{.*}}, <2 x i32> %{{.*}}) +__device__ u32x4 test_perm_pk16_b8_u4(unsigned long long a, unsigned long long b, u32x2 c) { + return __builtin_amdgcn_perm_pk16_b8_u4(a, b, c); +} + +// CHECK-LABEL: @_Z{{[0-9]+}}test_ds_atomic_async_barrier_arrive_b64 +// CHECK: call void @llvm.amdgcn.ds.atomic.async.barrier.arrive.b64(ptr addrspace(3) %{{.*}}) +__device__ void test_ds_atomic_async_barrier_arrive_b64(__attribute__((address_space(3))) __INT64_TYPE__ *p) { + __builtin_amdgcn_ds_atomic_async_barrier_arrive_b64(p); +} + +// CHECK-LABEL: @_Z{{[0-9]+}}test_ds_atomic_barrier_arrive_rtn_b64 +// CHECK: call i64 @llvm.amdgcn.ds.atomic.barrier.arrive.rtn.b64(ptr addrspace(3) %{{.*}}, i64 %{{.*}}) +__device__ long long test_ds_atomic_barrier_arrive_rtn_b64(__attribute__((address_space(3))) __INT64_TYPE__ *p, long long v) { + return __builtin_amdgcn_ds_atomic_barrier_arrive_rtn_b64(p, v); +} + +// CHECK-LABEL: @_Z{{[0-9]+}}test_wmma_scale16_f32_16x16x128_f8f6f4 +// CHECK: call contract <8 x float> @llvm.amdgcn.wmma.scale16.f32.16x16x128.f8f6f4.v8f32.v16i32.v16i32(i32 0, <16 x i32> %{{.*}}, i32 0, <16 x i32> %{{.*}}, i16 0, <8 x float> %{{.*}}, i32 0, i32 0, i64 %{{.*}}, i32 0, i32 0, i64 %{{.*}}, i1 false, i1 false) +__device__ v8f test_wmma_scale16_f32_16x16x128_f8f6f4(v16i a, v16i b, v8f c, long long s0, long long s1) { + return __builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4(0, a, 0, b, 0, c, 0, 0, s0, 0, 0, s1, false, false); +} + +// CHECK-LABEL: @_Z{{[0-9]+}}test_wmma_scale16_f32_32x16x128_f4 +// CHECK: call contract <16 x float> @llvm.amdgcn.wmma.scale16.f32.32x16x128.f4.v16f32.v16i32.v8i32(<16 x i32> %{{.*}}, <8 x i32> %{{.*}}, i16 0, <16 x float> %{{.*}}, i32 0, i32 0, i64 %{{.*}}, i32 0, i32 0, i64 %{{.*}}, i1 false, i1 false) +__device__ v16f test_wmma_scale16_f32_32x16x128_f4(v16i a, v8i b, v16f c, long long s0, long long s1) { + return __builtin_amdgcn_wmma_scale16_f32_32x16x128_f4(a, b, 0, c, 0, 0, s0, 0, 0, s1, false, false); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
