[clang] [AMDGPU][NFC] Change AMDGPU builtins to use ExtVector (PR #176033)
via cfe-commits
cfe-commits at lists.llvm.org
Wed Jan 14 12:53:48 PST 2026
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang
Author: Joseph Huber (jhuber6)
<details>
<summary>Changes</summary>
Summary:
These currently use the GNU vectors, not the OpenCL vectors, which is
strange.
---
Patch is 188.81 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/176033.diff
1 Files Affected:
- (modified) clang/include/clang/Basic/BuiltinsAMDGPU.td (+470-470)
``````````diff
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td
index b7d1236549eee..12ffad305e7c0 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td
@@ -163,26 +163,26 @@ def __builtin_amdgcn_sad_hi_u8 : AMDGPUBuiltin<"unsigned int(unsigned int, unsig
def __builtin_amdgcn_sad_u16 : AMDGPUBuiltin<"unsigned int(unsigned int, unsigned int, unsigned int)", [Const], "sad-insts">;
def __builtin_amdgcn_qsad_pk_u16_u8 : AMDGPUBuiltin<"uint64_t(uint64_t, unsigned int, uint64_t)", [Const], "qsad-insts">;
def __builtin_amdgcn_mqsad_pk_u16_u8 : AMDGPUBuiltin<"uint64_t(uint64_t, unsigned int, uint64_t)", [Const]>;
-def __builtin_amdgcn_mqsad_u32_u8 : AMDGPUBuiltin<"_Vector<4, unsigned int>(uint64_t, unsigned int, _Vector<4, unsigned int>)", [Const]>;
+def __builtin_amdgcn_mqsad_u32_u8 : AMDGPUBuiltin<"_ExtVector<4, unsigned int>(uint64_t, unsigned int, _ExtVector<4, unsigned int>)", [Const]>;
def __builtin_amdgcn_make_buffer_rsrc : AMDGPUBuiltin<"__amdgpu_buffer_rsrc_t(void *, short, int64_t, int)", [Const]>;
def __builtin_amdgcn_raw_buffer_store_b8 : AMDGPUBuiltin<"void(unsigned char, __amdgpu_buffer_rsrc_t, int, int, _Constant int)">;
def __builtin_amdgcn_raw_buffer_store_b16 : AMDGPUBuiltin<"void(unsigned short, __amdgpu_buffer_rsrc_t, int, int, _Constant int)">;
def __builtin_amdgcn_raw_buffer_store_b32 : AMDGPUBuiltin<"void(unsigned int, __amdgpu_buffer_rsrc_t, int, int, _Constant int)">;
-def __builtin_amdgcn_raw_buffer_store_b64 : AMDGPUBuiltin<"void(_Vector<2, unsigned int>, __amdgpu_buffer_rsrc_t, int, int, _Constant int)">;
-def __builtin_amdgcn_raw_buffer_store_b96 : AMDGPUBuiltin<"void(_Vector<3, unsigned int>, __amdgpu_buffer_rsrc_t, int, int, _Constant int)">;
-def __builtin_amdgcn_raw_buffer_store_b128 : AMDGPUBuiltin<"void(_Vector<4, unsigned int>, __amdgpu_buffer_rsrc_t, int, int, _Constant int)">;
+def __builtin_amdgcn_raw_buffer_store_b64 : AMDGPUBuiltin<"void(_ExtVector<2, unsigned int>, __amdgpu_buffer_rsrc_t, int, int, _Constant int)">;
+def __builtin_amdgcn_raw_buffer_store_b96 : AMDGPUBuiltin<"void(_ExtVector<3, unsigned int>, __amdgpu_buffer_rsrc_t, int, int, _Constant int)">;
+def __builtin_amdgcn_raw_buffer_store_b128 : AMDGPUBuiltin<"void(_ExtVector<4, unsigned int>, __amdgpu_buffer_rsrc_t, int, int, _Constant int)">;
def __builtin_amdgcn_raw_buffer_load_b8 : AMDGPUBuiltin<"unsigned char(__amdgpu_buffer_rsrc_t, int, int, _Constant int)">;
def __builtin_amdgcn_raw_buffer_load_b16 : AMDGPUBuiltin<"unsigned short(__amdgpu_buffer_rsrc_t, int, int, _Constant int)">;
def __builtin_amdgcn_raw_buffer_load_b32 : AMDGPUBuiltin<"unsigned int(__amdgpu_buffer_rsrc_t, int, int, _Constant int)">;
-def __builtin_amdgcn_raw_buffer_load_b64 : AMDGPUBuiltin<"_Vector<2, unsigned int>(__amdgpu_buffer_rsrc_t, int, int, _Constant int)">;
-def __builtin_amdgcn_raw_buffer_load_b96 : AMDGPUBuiltin<"_Vector<3, unsigned int>(__amdgpu_buffer_rsrc_t, int, int, _Constant int)">;
-def __builtin_amdgcn_raw_buffer_load_b128 : AMDGPUBuiltin<"_Vector<4, unsigned int>(__amdgpu_buffer_rsrc_t, int, int, _Constant int)">;
+def __builtin_amdgcn_raw_buffer_load_b64 : AMDGPUBuiltin<"_ExtVector<2, unsigned int>(__amdgpu_buffer_rsrc_t, int, int, _Constant int)">;
+def __builtin_amdgcn_raw_buffer_load_b96 : AMDGPUBuiltin<"_ExtVector<3, unsigned int>(__amdgpu_buffer_rsrc_t, int, int, _Constant int)">;
+def __builtin_amdgcn_raw_buffer_load_b128 : AMDGPUBuiltin<"_ExtVector<4, unsigned int>(__amdgpu_buffer_rsrc_t, int, int, _Constant int)">;
def __builtin_amdgcn_raw_ptr_buffer_atomic_add_i32 : AMDGPUBuiltin<"int(int, __amdgpu_buffer_rsrc_t, int, int, _Constant int)">;
def __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32 : AMDGPUBuiltin<"float(float, __amdgpu_buffer_rsrc_t, int, int, _Constant int)", [], "atomic-fadd-rtn-insts">;
-def __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16 : AMDGPUBuiltin<"_Vector<2, _Float16>(_Vector<2, _Float16>, __amdgpu_buffer_rsrc_t, int, int, _Constant int)", [], "atomic-buffer-global-pk-add-f16-insts">;
+def __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16 : AMDGPUBuiltin<"_ExtVector<2, _Float16>(_ExtVector<2, _Float16>, __amdgpu_buffer_rsrc_t, int, int, _Constant int)", [], "atomic-buffer-global-pk-add-f16-insts">;
def __builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f32 : AMDGPUBuiltin<"float(float, __amdgpu_buffer_rsrc_t, int, int, _Constant int)", [], "atomic-fmin-fmax-global-f32">;
def __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32 : AMDGPUBuiltin<"float(float, __amdgpu_buffer_rsrc_t, int, int, _Constant int)", [], "atomic-fmin-fmax-global-f32">;
@@ -270,7 +270,7 @@ def __builtin_amdgcn_fmed3h : AMDGPUBuiltin<"__fp16(__fp16, __fp16, __fp16)", [C
def __builtin_amdgcn_global_atomic_fadd_f64 : AMDGPUBuiltin<"double(double address_space<1> *, double)", [], "gfx90a-insts">;
def __builtin_amdgcn_global_atomic_fadd_f32 : AMDGPUBuiltin<"float(float address_space<1> *, float)", [], "atomic-fadd-rtn-insts">;
-def __builtin_amdgcn_global_atomic_fadd_v2f16 : AMDGPUBuiltin<"_Vector<2, _Float16>(_Vector<2, _Float16 address_space<1> *>, _Vector<2, _Float16>)", [CustomTypeChecking], "atomic-buffer-global-pk-add-f16-insts">;
+def __builtin_amdgcn_global_atomic_fadd_v2f16 : AMDGPUBuiltin<"_ExtVector<2, _Float16>(_ExtVector<2, _Float16 address_space<1> *>, _ExtVector<2, _Float16>)", [CustomTypeChecking], "atomic-buffer-global-pk-add-f16-insts">;
def __builtin_amdgcn_global_atomic_fmin_f64 : AMDGPUBuiltin<"double(double address_space<1> *, double)", [], "gfx90a-insts">;
def __builtin_amdgcn_global_atomic_fmax_f64 : AMDGPUBuiltin<"double(double address_space<1> *, double)", [], "gfx90a-insts">;
@@ -282,11 +282,11 @@ def __builtin_amdgcn_ds_atomic_fadd_f64 : AMDGPUBuiltin<"double(double address_s
def __builtin_amdgcn_ds_atomic_fadd_f32 : AMDGPUBuiltin<"float(float address_space<3> *, float)", [], "gfx8-insts">;
def __builtin_amdgcn_flat_atomic_fadd_f32 : AMDGPUBuiltin<"float(float address_space<0> *, float)", [], "gfx940-insts">;
-def __builtin_amdgcn_flat_atomic_fadd_v2f16 : AMDGPUBuiltin<"_Vector<2, _Float16>(_Vector<2, _Float16 address_space<0> *>, _Vector<2, _Float16>)", [CustomTypeChecking], "atomic-flat-pk-add-16-insts">;
-def __builtin_amdgcn_flat_atomic_fadd_v2bf16 : AMDGPUBuiltin<"_Vector<2, short>(_Vector<2, short address_space<0> *>, _Vector<2, short>)", [CustomTypeChecking], "atomic-flat-pk-add-16-insts">;
-def __builtin_amdgcn_global_atomic_fadd_v2bf16 : AMDGPUBuiltin<"_Vector<2, short>(_Vector<2, short address_space<1> *>, _Vector<2, short>)", [CustomTypeChecking], "atomic-global-pk-add-bf16-inst">;
-def __builtin_amdgcn_ds_atomic_fadd_v2bf16 : AMDGPUBuiltin<"_Vector<2, short>(_Vector<2, short address_space<3> *>, _Vector<2, short>)", [CustomTypeChecking], "atomic-ds-pk-add-16-insts">;
-def __builtin_amdgcn_ds_atomic_fadd_v2f16 : AMDGPUBuiltin<"_Vector<2, _Float16>(_Vector<2, _Float16 address_space<3> *>, _Vector<2, _Float16>)", [CustomTypeChecking], "atomic-ds-pk-add-16-insts">;
+def __builtin_amdgcn_flat_atomic_fadd_v2f16 : AMDGPUBuiltin<"_ExtVector<2, _Float16>(_ExtVector<2, _Float16 address_space<0> *>, _ExtVector<2, _Float16>)", [CustomTypeChecking], "atomic-flat-pk-add-16-insts">;
+def __builtin_amdgcn_flat_atomic_fadd_v2bf16 : AMDGPUBuiltin<"_ExtVector<2, short>(_ExtVector<2, short address_space<0> *>, _ExtVector<2, short>)", [CustomTypeChecking], "atomic-flat-pk-add-16-insts">;
+def __builtin_amdgcn_global_atomic_fadd_v2bf16 : AMDGPUBuiltin<"_ExtVector<2, short>(_ExtVector<2, short address_space<1> *>, _ExtVector<2, short>)", [CustomTypeChecking], "atomic-global-pk-add-bf16-inst">;
+def __builtin_amdgcn_ds_atomic_fadd_v2bf16 : AMDGPUBuiltin<"_ExtVector<2, short>(_ExtVector<2, short address_space<3> *>, _ExtVector<2, short>)", [CustomTypeChecking], "atomic-ds-pk-add-16-insts">;
+def __builtin_amdgcn_ds_atomic_fadd_v2f16 : AMDGPUBuiltin<"_ExtVector<2, _Float16>(_ExtVector<2, _Float16 address_space<3> *>, _ExtVector<2, _Float16>)", [CustomTypeChecking], "atomic-ds-pk-add-16-insts">;
def __builtin_amdgcn_load_to_lds : AMDGPUBuiltin<"void(void *, void address_space<3> *, _Constant unsigned int, _Constant int, _Constant unsigned int)", [], "vmem-to-lds-load-insts">;
def __builtin_amdgcn_global_load_lds : AMDGPUBuiltin<"void(void address_space<1> *, void address_space<3> *, _Constant unsigned int, _Constant int, _Constant unsigned int)", [], "vmem-to-lds-load-insts">;
@@ -294,12 +294,12 @@ def __builtin_amdgcn_global_load_lds : AMDGPUBuiltin<"void(void address_space<1>
// Deep learning builtins.
//===----------------------------------------------------------------------===//
-def __builtin_amdgcn_fdot2 : AMDGPUBuiltin<"float(_Vector<2, _Float16>, _Vector<2, _Float16>, float, _Constant bool)", [Const], "dot10-insts">;
-def __builtin_amdgcn_fdot2_f16_f16 : AMDGPUBuiltin<"_Float16(_Vector<2, _Float16>, _Vector<2, _Float16>, _Float16)", [Const], "dot9-insts">;
-def __builtin_amdgcn_fdot2_bf16_bf16 : AMDGPUBuiltin<"short(_Vector<2, short>, _Vector<2, short>, short)", [Const], "dot9-insts">;
-def __builtin_amdgcn_fdot2_f32_bf16 : AMDGPUBuiltin<"float(_Vector<2, short>, _Vector<2, short>, float, _Constant bool)", [Const], "dot12-insts">;
-def __builtin_amdgcn_sdot2 : AMDGPUBuiltin<"int(_Vector<2, short>, _Vector<2, short>, int, _Constant bool)", [Const], "dot2-insts">;
-def __builtin_amdgcn_udot2 : AMDGPUBuiltin<"unsigned int(_Vector<2, unsigned short>, _Vector<2, unsigned short>, unsigned int, _Constant bool)", [Const], "dot2-insts">;
+def __builtin_amdgcn_fdot2 : AMDGPUBuiltin<"float(_ExtVector<2, _Float16>, _ExtVector<2, _Float16>, float, _Constant bool)", [Const], "dot10-insts">;
+def __builtin_amdgcn_fdot2_f16_f16 : AMDGPUBuiltin<"_Float16(_ExtVector<2, _Float16>, _ExtVector<2, _Float16>, _Float16)", [Const], "dot9-insts">;
+def __builtin_amdgcn_fdot2_bf16_bf16 : AMDGPUBuiltin<"short(_ExtVector<2, short>, _ExtVector<2, short>, short)", [Const], "dot9-insts">;
+def __builtin_amdgcn_fdot2_f32_bf16 : AMDGPUBuiltin<"float(_ExtVector<2, short>, _ExtVector<2, short>, float, _Constant bool)", [Const], "dot12-insts">;
+def __builtin_amdgcn_sdot2 : AMDGPUBuiltin<"int(_ExtVector<2, short>, _ExtVector<2, short>, int, _Constant bool)", [Const], "dot2-insts">;
+def __builtin_amdgcn_udot2 : AMDGPUBuiltin<"unsigned int(_ExtVector<2, unsigned short>, _ExtVector<2, unsigned short>, unsigned int, _Constant bool)", [Const], "dot2-insts">;
def __builtin_amdgcn_sdot4 : AMDGPUBuiltin<"int(int, int, int, _Constant bool)", [Const], "dot1-insts">;
def __builtin_amdgcn_udot4 : AMDGPUBuiltin<"unsigned int(unsigned int, unsigned int, unsigned int, _Constant bool)", [Const], "dot7-insts">;
def __builtin_amdgcn_sudot4 : AMDGPUBuiltin<"int(_Constant bool, int, _Constant bool, int, int, _Constant bool)", [Const], "dot8-insts">;
@@ -310,7 +310,7 @@ def __builtin_amdgcn_dot4_f32_fp8_bf8 : AMDGPUBuiltin<"float(unsigned int, unsig
def __builtin_amdgcn_dot4_f32_bf8_fp8 : AMDGPUBuiltin<"float(unsigned int, unsigned int, float)", [Const], "dot11-insts">;
def __builtin_amdgcn_dot4_f32_fp8_fp8 : AMDGPUBuiltin<"float(unsigned int, unsigned int, float)", [Const], "dot11-insts">;
def __builtin_amdgcn_dot4_f32_bf8_bf8 : AMDGPUBuiltin<"float(unsigned int, unsigned int, float)", [Const], "dot11-insts">;
-def __builtin_amdgcn_fdot2c_f32_bf16 : AMDGPUBuiltin<"float(_Vector<2, __bf16>, _Vector<2, __bf16>, float, _Constant bool)", [Const], "dot13-insts">;
+def __builtin_amdgcn_fdot2c_f32_bf16 : AMDGPUBuiltin<"float(_ExtVector<2, __bf16>, _ExtVector<2, __bf16>, float, _Constant bool)", [Const], "dot13-insts">;
//===----------------------------------------------------------------------===//
// GFX10+ only builtins.
@@ -326,10 +326,10 @@ def __builtin_amdgcn_s_ttracedata_imm : AMDGPUBuiltin<"void(_Constant short)", [
// Postfix l indicates the 1st argument is i64.
// Postfix h indicates the 4/5-th arguments are half4.
//===----------------------------------------------------------------------===//
-def __builtin_amdgcn_image_bvh_intersect_ray : AMDGPUBuiltin<"_Vector<4, unsigned int>(unsigned int, float, _Vector<4, float>, _Vector<4, float>, _Vector<4, float>, _Vector<4, unsigned int>)", [Const], "gfx10-insts">;
-def __builtin_amdgcn_image_bvh_intersect_ray_h : AMDGPUBuiltin<"_Vector<4, unsigned int>(unsigned int, float, _Vector<4, float>, _Vector<4, _Float16>, _Vector<4, _Float16>, _Vector<4, unsigned int>)", [Const], "gfx10-insts">;
-def __builtin_amdgcn_image_bvh_intersect_ray_l : AMDGPUBuiltin<"_Vector<4, unsigned int>(uint64_t, float, _Vector<4, float>, _Vector<4, float>, _Vector<4, float>, _Vector<4, unsigned int>)", [Const], "gfx10-insts">;
-def __builtin_amdgcn_image_bvh_intersect_ray_lh : AMDGPUBuiltin<"_Vector<4, unsigned int>(uint64_t, float, _Vector<4, float>, _Vector<4, _Float16>, _Vector<4, _Float16>, _Vector<4, unsigned int>)", [Const], "gfx10-insts">;
+def __builtin_amdgcn_image_bvh_intersect_ray : AMDGPUBuiltin<"_ExtVector<4, unsigned int>(unsigned int, float, _ExtVector<4, float>, _ExtVector<4, float>, _ExtVector<4, float>, _ExtVector<4, unsigned int>)", [Const], "gfx10-insts">;
+def __builtin_amdgcn_image_bvh_intersect_ray_h : AMDGPUBuiltin<"_ExtVector<4, unsigned int>(unsigned int, float, _ExtVector<4, float>, _ExtVector<4, _Float16>, _ExtVector<4, _Float16>, _ExtVector<4, unsigned int>)", [Const], "gfx10-insts">;
+def __builtin_amdgcn_image_bvh_intersect_ray_l : AMDGPUBuiltin<"_ExtVector<4, unsigned int>(uint64_t, float, _ExtVector<4, float>, _ExtVector<4, float>, _ExtVector<4, float>, _ExtVector<4, unsigned int>)", [Const], "gfx10-insts">;
+def __builtin_amdgcn_image_bvh_intersect_ray_lh : AMDGPUBuiltin<"_ExtVector<4, unsigned int>(uint64_t, float, _ExtVector<4, float>, _ExtVector<4, _Float16>, _ExtVector<4, _Float16>, _ExtVector<4, unsigned int>)", [Const], "gfx10-insts">;
//===----------------------------------------------------------------------===//
@@ -345,28 +345,28 @@ def __builtin_amdgcn_s_wait_event_export_ready : AMDGPUBuiltin<"void()", [], "gf
// Postfix w32 indicates the builtin requires wavefront size of 32.
// Postfix w64 indicates the builtin requires wavefront size of 64.
//===----------------------------------------------------------------------===//
-def __builtin_amdgcn_wmma_f32_16x16x16_f16_w32 : AMDGPUBuiltin<"_Vector<8, float>(_Vector<16, _Float16>, _Vector<16, _Float16>, _Vector<8, float>)", [Const], "gfx11-insts,wavefrontsize32">;
-def __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32 : AMDGPUBuiltin<"_Vector<8, float>(_Vector<16, short>, _Vector<16, short>, _Vector<8, float>)", [Const], "gfx11-insts,wavefrontsize32">;
-def __builtin_amdgcn_wmma_f16_16x16x16_f16_w32 : AMDGPUBuiltin<"_Vector<16, _Float16>(_Vector<16, _Float16>, _Vector<16, _Float16>, _Vector<16, _Float16>, _Constant bool)", [Const], "gfx11-insts,wavefrontsize32">;
-def __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32 : AMDGPUBuiltin<"_Vector<16, short>(_Vector<16, short>, _Vector<16, short>, _Vector<16, short>, _Constant bool)", [Const], "gfx11-insts,wavefrontsize32">;
-def __builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32 : AMDGPUBuiltin<"_Vector<16, _Float16>(_Vector<16, _Float16>, _Vector<16, _Float16>, _Vector<16, _Float16>, _Constant bool)", [Const], "gfx11-insts,wavefrontsize32">;
-def __builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32 : AMDGPUBuiltin<"_Vector<16, short>(_Vector<16, short>, _Vector<16, short>, _Vector<16, short>, _Constant bool)", [Const], "gfx11-insts,wavefrontsize32">;
-def __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32 : AMDGPUBuiltin<"_Vector<8, int>(_Constant bool, _Vector<4, int>, _Constant bool, _Vector<4, int>, _Vector<8, int>, _Constant bool)", [Const], "gfx11-insts,wavefrontsize32">;
-def __builtin_amdgcn_wmma_i32_16x16x16_iu4_w32 : AMDGPUBuiltin<"_Vector<8, int>(_Constant bool, _Vector<2, int>, _Constant bool, _Vector<2, int>, _Vector<8, int>, _Constant bool)", [Const], "gfx11-insts,wavefrontsize32">;
-
-def __builtin_amdgcn_wmma_f32_16x16x16_f16_w64 : AMDGPUBuiltin<"_Vector<4, float>(_Vector<16, _Float16>, _Vector<16, _Float16>, _Vector<4, float>)", [Const], "gfx11-insts,wavefrontsize64">;
-def __builtin_amdgcn_wmma_f32_16x16x16_bf16_w64 : AMDGPUBuiltin<"_Vector<4, float>(_Vector<16, short>, _Vector<16, short>, _Vector<4, float>)", [Const], "gfx11-insts,wavefrontsize64">;
-def __builtin_amdgcn_wmma_f16_16x16x16_f16_w64 : AMDGPUBuiltin<"_Vector<8, _Float16>(_Vector<16, _Float16>, _Vector<16, _Float16>, _Vector<8, _Float16>, _Constant bool)", [Const], "gfx11-insts,wavefrontsize64">;
-def __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64 : AMDGPUBuiltin<"_Vector<8, short>(_Vector<16, short>, _Vector<16, short>, _Vector<8, short>, _Constant bool)", [Const], "gfx11-insts,wavefrontsize64">;
-def __builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w64 : AMDGPUBuiltin<"_Vector<8, _Float16>(_Vector<16, _Float16>, _Vector<16, _Float16>, _Vector<8, _Float16>, _Constant bool)", [Const], "gfx11-insts,wavefrontsize64">;
-def __builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w64 : AMDGPUBuiltin<"_Vector<8, short>(_Vector<16, short>, _Vector<16, short>, _Vector<8, short>, _Constant bool)", [Const], "gfx11-insts,wavefrontsize64">;
-def __builtin_amdgcn_wmma_i32_16x16x16_iu8_w64 : AMDGPUBuiltin<"_Vector<4, int>(_Constant bool, _Vector<4, int>, _Constant bool, _Vector<4, int>, _Vector<4, int>, _Constant bool)", [Const], "gfx11-insts,wavefrontsize64">;
-def __builtin_amdgcn_wmma_i32_16x16x16_iu4_w64 : AMDGPUBuiltin<"_Vector<4, int>(_Constant bool, _Vector<2, int>, _Constant bool, _Vector<2, int>, _Vector<4, int>, _Constant bool)", [Const], "gfx11-insts,wavefrontsize64">;
+def __builtin_amdgcn_wmma_f32_16x16x16_f16_w32 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<16, _Float16>, _ExtVector<16, _Float16>, _ExtVector<8, float>)", [Const], "gfx11-insts,wavefrontsize32">;
+def __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<16, short>, _ExtVector<16, short>, _ExtVector<8, float>)", [Const], "gfx11-insts,wavefrontsize32">;
+def __builtin_amdgcn_wmma_f16_16x16x16_f16_w32 : AMDGPUBuiltin<"_ExtVector<16, _Float16>(_ExtVector<16, _Float16>, _ExtVector<16, _Float16>, _ExtVector<16, _Float16>, _Constant bool)", [Const], "gfx11-insts,wavefrontsize32">;
+def __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32 : AMDGPUBuiltin<"_ExtVector<16, short>(_ExtVector<16, short>, _ExtVector<16, short>, _ExtVector<16, short>, _Constant bool)", [Const], "gfx11-insts,wavefrontsize32">;
+def __builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32 : AMDGPUBuiltin<"_ExtVector<16, _Float16>(_ExtVector<16, _Float16>, _ExtVector<16, _Float16>, _ExtVector<16, _Float16>, _Constant bool)", [Const], "gfx11-insts,wavefrontsize32">;
+def __builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32 : AMDGPUBuiltin<"_ExtVector<16, short>(_ExtVector<16, short>, _ExtVector<16, short>, _ExtVector<16, short>, _Constant bool)", [Const], "gfx11-insts,wavefrontsize32">;
+def __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32 : AMDGPUBuiltin<"_ExtVector<8, int>(_Constant bool, _ExtVector<4, int>, _Constant bool, _ExtVector<4, int>, _ExtVector<8, int>, _Constant bool)", [Const], "gfx11-insts,wavefrontsize32">;
+def __builtin_amdgcn_wmma_i32_16x16x16_iu4_w32 : AMDGPUBuiltin<"_ExtVector<8, int>(_Constant bool, _ExtVector<2, int>, _Constant bool, _ExtVector<2, int>, _ExtVector<8, int>, _Constant bool)", [Const], "gfx11-insts,wavefrontsize32">;
+
+def __builtin_amdgcn_wmma_f32_16x16x16_f16_w64 : AMDGPUBuiltin<"_ExtVector<4, float>(_ExtVector<16, _Float16>, _ExtVector<16, _Float16>, _ExtVector<4, float>)", [Const], "gfx11-insts,wavefrontsize64">;
+def __builtin_amdgcn_wmma_f32_16x16x16_bf16_w64 : AMDGPUBuiltin<"_ExtVector<4, float>(_ExtVector<16, short>, _ExtVector<16, short>, _ExtVector<4, float>)", [Const], "gfx11-insts,wavefrontsize64">;
+def __builtin_amdgcn_wmma_f16_16x16x16_f16_w64 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<16, _Float16>, _ExtVector<16, _Float16>, _ExtVector<8, _Float16>, _Constant bool)", [Const], "gfx11-insts,wavefrontsize64">;
+def __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64 : AMDGPUBuiltin<"_ExtVector<8, short>(_ExtVector<16, short>, _ExtVector<16, short>, _ExtVector<8, short>, _Constant bool)", [Const], "gfx11-insts,wavefrontsize64">;
+def __builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w64 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<16, _Float16>, _ExtVector<16, _Float16>, _ExtVector<8, _Float16>, _Constant bool)", [Const], "gfx11-insts,wavefrontsize64">;
+def __builtin_amdgcn_wmma_bf16_16x16x16...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/176033
More information about the cfe-commits
mailing list