Skip to content

Commit b025864

Browse files
OutOfCacheOutOfCache
authored
[AMDGPU] - Add clang builtins for tied WMMA intrinsics (#70669)
Add clang builtins for the new tied wmma intrinsics. These variations tie the destination accumulator matrix to the input accumulator matrix. See #69903 for context.
1 parent 93ae263 commit b025864

File tree

8 files changed

+214
-0
lines changed

8 files changed

+214
-0
lines changed

‎clang/include/clang/Basic/BuiltinsAMDGPU.def‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -292,13 +292,17 @@ TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_f16_w32, "V8fV16hV16hV8f", "nc
292292
TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32, "V8fV16sV16sV8f", "nc", "gfx11-insts")
293293
TARGET_BUILTIN(__builtin_amdgcn_wmma_f16_16x16x16_f16_w32, "V16hV16hV16hV16hIb", "nc", "gfx11-insts")
294294
TARGET_BUILTIN(__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32, "V16sV16sV16sV16sIb", "nc", "gfx11-insts")
295+
TARGET_BUILTIN(__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32, "V16hV16hV16hV16hIb", "nc", "gfx11-insts")
296+
TARGET_BUILTIN(__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32, "V16sV16sV16sV16sIb", "nc", "gfx11-insts")
295297
TARGET_BUILTIN(__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32, "V8iIbV4iIbV4iV8iIb", "nc", "gfx11-insts")
296298
TARGET_BUILTIN(__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32, "V8iIbV2iIbV2iV8iIb", "nc", "gfx11-insts")
297299

298300
TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_f16_w64, "V4fV16hV16hV4f", "nc", "gfx11-insts")
299301
TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64, "V4fV16sV16sV4f", "nc", "gfx11-insts")
300302
TARGET_BUILTIN(__builtin_amdgcn_wmma_f16_16x16x16_f16_w64, "V8hV16hV16hV8hIb", "nc", "gfx11-insts")
301303
TARGET_BUILTIN(__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64, "V8sV16sV16sV8sIb", "nc", "gfx11-insts")
304+
TARGET_BUILTIN(__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w64, "V8hV16hV16hV8hIb", "nc", "gfx11-insts")
305+
TARGET_BUILTIN(__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w64, "V8sV16sV16sV8sIb", "nc", "gfx11-insts")
302306
TARGET_BUILTIN(__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64, "V4iIbV4iIbV4iV4iIb", "nc", "gfx11-insts")
303307
TARGET_BUILTIN(__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64, "V4iIbV2iIbV2iV4iIb", "nc", "gfx11-insts")
304308

‎clang/lib/CodeGen/CGBuiltin.cpp‎

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -18024,9 +18024,13 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
1802418024
}
1802518025

1802618026
case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32:
18027+
case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32:
1802718028
case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64:
18029+
case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w64:
1802818030
case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w32:
18031+
case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32:
1802918032
case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w64:
18033+
case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w64:
1803018034
case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32:
1803118035
case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64:
1803218036
case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w32:
@@ -18064,6 +18068,16 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
1806418068
ArgForMatchingRetType = 2;
1806518069
BuiltinWMMAOp = Intrinsic::amdgcn_wmma_bf16_16x16x16_bf16;
1806618070
break;
18071+
case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32:
18072+
case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w64:
18073+
ArgForMatchingRetType = 2;
18074+
BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f16_16x16x16_f16_tied;
18075+
break;
18076+
case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32:
18077+
case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w64:
18078+
ArgForMatchingRetType = 2;
18079+
BuiltinWMMAOp = Intrinsic::amdgcn_wmma_bf16_16x16x16_bf16_tied;
18080+
break;
1806718081
case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32:
1806818082
case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64:
1806918083
ArgForMatchingRetType = 4;

‎clang/test/CodeGenOpenCL/builtins-amdgcn-wmma-w32-gfx10-err.cl‎

Lines changed: 34 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,34 @@
1+
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
2+
// REQUIRES: amdgpu-registered-target
3+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1030 \
4+
// RUN: -verify -DWMMA_GFX1100_TESTS -S -o - %s
5+
6+
7+
typedef float v4f __attribute__((ext_vector_type(4)));
8+
typedef float v8f __attribute__((ext_vector_type(8)));
9+
typedef half v16h __attribute__((ext_vector_type(16)));
10+
typedef int v2i __attribute__((ext_vector_type(2)));
11+
typedef int v4i __attribute__((ext_vector_type(4)));
12+
typedef int v8i __attribute__((ext_vector_type(8)));
13+
typedef short v16s __attribute__((ext_vector_type(16)));
14+
15+
#ifdef WMMA_GFX1100_TESTS
16+
17+
// Wave32
18+
19+
void test_amdgcn_wmma_f32_16x16x16_bf16_w32(global v8f* out8f, v16s a16s, v16s b16s, v8f c8f,
20+
global v16h* out16h, v16h a16h, v16h b16h, v16h c16h,
21+
global v16s* out16s, v2i a2i, v2i b2i, v16s c16s,
22+
global v8i* out8i, v4i a4i, v4i b4i, v8i c8i)
23+
{
24+
*out8f = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32(a16h, b16h, c8f); // expected-error{{'__builtin_amdgcn_wmma_f32_16x16x16_f16_w32' needs target feature gfx11-insts}}
25+
*out8f = __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32(a16s, b16s, c8f); // expected-error{{'__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32' needs target feature gfx11-insts}}
26+
*out16h = __builtin_amdgcn_wmma_f16_16x16x16_f16_w32(a16h, b16h, c16h, true); // expected-error{{'__builtin_amdgcn_wmma_f16_16x16x16_f16_w32' needs target feature gfx11-insts}}
27+
*out16s = __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32(a16s, b16s, c16s, true); // expected-error{{'__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32' needs target feature gfx11-insts}}
28+
*out16h = __builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32(a16h, b16h, c16h, true); // expected-error{{'__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32' needs target feature gfx11-insts}}
29+
*out16s = __builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32(a16s, b16s, c16s, true); // expected-error{{'__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32' needs target feature gfx11-insts}}
30+
*out8i = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32(true, a4i, true, b4i, c8i, false); // expected-error{{'__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32' needs target feature gfx11-insts}}
31+
*out8i = __builtin_amdgcn_wmma_i32_16x16x16_iu4_w32(true, a2i, true, b2i, c8i, false); // expected-error{{'__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32' needs target feature gfx11-insts}}
32+
}
33+
34+
#endif

‎clang/test/CodeGenOpenCL/builtins-amdgcn-wmma-w32-param-err.cl‎

Lines changed: 34 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,34 @@
1+
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
2+
// REQUIRES: amdgpu-registered-target
3+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 \
4+
// RUN: -verify -DWMMA_GFX1100_TESTS -S -o - %s
5+
6+
7+
typedef float v4f __attribute__((ext_vector_type(4)));
8+
typedef float v8f __attribute__((ext_vector_type(8)));
9+
typedef half v16h __attribute__((ext_vector_type(16)));
10+
typedef int v2i __attribute__((ext_vector_type(2)));
11+
typedef int v4i __attribute__((ext_vector_type(4)));
12+
typedef int v8i __attribute__((ext_vector_type(8)));
13+
typedef short v16s __attribute__((ext_vector_type(16)));
14+
15+
#ifdef WMMA_GFX1100_TESTS
16+
17+
// Wave32
18+
19+
void test_amdgcn_wmma_f32_16x16x16_bf16_w32(v16s a16s, v16s b16s, unsigned int i,
20+
global v16h* out16h, v16h a16h, v16h b16h, v16h c16h,
21+
global v16s* out16s, v2i a2i, v2i b2i, v16s c16s,
22+
global v8i* out8i, v4i a4i, v4i b4i, v8i c8i)
23+
{
24+
*out16h = __builtin_amdgcn_wmma_f16_16x16x16_f16_w32(a16h, b16h, c16h, i); // expected-error{{argument to '__builtin_amdgcn_wmma_f16_16x16x16_f16_w32' must be a constant integer}}
25+
*out16s = __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32(a16s, b16s, c16s, i); // expected-error{{argument to '__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32' must be a constant integer}}
26+
*out16h = __builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32(a16h, b16h, c16h, i); // expected-error{{argument to '__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32' must be a constant integer}}
27+
*out16s = __builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32(a16s, b16s, c16s, i); // expected-error{{argument to '__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32' must be a constant integer}}
28+
*out8i = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32(i, a4i, true, b4i, c8i, false); // expected-error{{argument to '__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32' must be a constant integer}}
29+
*out8i = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32(true, a4i, i, b4i, c8i, false); // expected-error{{argument to '__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32' must be a constant integer}}
30+
*out8i = __builtin_amdgcn_wmma_i32_16x16x16_iu4_w32(i, a2i, true, b2i, c8i, false); // expected-error{{argument to '__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32' must be a constant integer}}
31+
*out8i = __builtin_amdgcn_wmma_i32_16x16x16_iu4_w32(true, a2i, i, b2i, c8i, false); // expected-error{{argument to '__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32' must be a constant integer}}
32+
}
33+
34+
#endif

‎clang/test/CodeGenOpenCL/builtins-amdgcn-wmma-w32.cl‎

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -74,6 +74,36 @@ void test_amdgcn_wmma_bf16_16x16x16_bf16_w32(global v16s* out, v16s a, v16s b, v
7474
*out = __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32(a, b, c, true);
7575
}
7676

77+
//
78+
// amdgcn_wmma_f16_16x16x16_f16_tied
79+
//
80+
81+
// CHECK-GFX1100-LABEL: @test_amdgcn_wmma_f16_16x16x16_f16_tied_w32(
82+
// CHECK-GFX1100-NEXT: entry:
83+
// CHECK-GFX1100-NEXT: [[TMP0:%.*]] = tail call <16 x half> @llvm.amdgcn.wmma.f16.16x16x16.f16.tied.v16f16(<16 x half> [[A:%.*]], <16 x half> [[B:%.*]], <16 x half> [[C:%.*]], i1 true)
84+
// CHECK-GFX1100-NEXT: store <16 x half> [[TMP0]], ptr addrspace(1) [[OUT:%.*]], align 32, !tbaa [[TBAA4]]
85+
// CHECK-GFX1100-NEXT: ret void
86+
//
87+
void test_amdgcn_wmma_f16_16x16x16_f16_tied_w32(global v16h* out, v16h a, v16h b, v16h c)
88+
{
89+
*out = __builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32(a, b, c, true);
90+
}
91+
92+
//
93+
// amdgcn_wmma_bf16_16x16x16_bf16_tied
94+
//
95+
96+
// CHECK-GFX1100-LABEL: @test_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32(
97+
// CHECK-GFX1100-NEXT: entry:
98+
// CHECK-GFX1100-NEXT: [[TMP0:%.*]] = tail call <16 x i16> @llvm.amdgcn.wmma.bf16.16x16x16.bf16.tied.v16i16(<16 x i16> [[A:%.*]], <16 x i16> [[B:%.*]], <16 x i16> [[C:%.*]], i1 true)
99+
// CHECK-GFX1100-NEXT: store <16 x i16> [[TMP0]], ptr addrspace(1) [[OUT:%.*]], align 32, !tbaa [[TBAA4]]
100+
// CHECK-GFX1100-NEXT: ret void
101+
//
102+
void test_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32(global v16s* out, v16s a, v16s b, v16s c)
103+
{
104+
*out = __builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32(a, b, c, true);
105+
}
106+
77107
//
78108
// amdgcn_wmma_i32_16x16x16_iu8
79109
//

‎clang/test/CodeGenOpenCL/builtins-amdgcn-wmma-w64-gfx10-err.cl‎

Lines changed: 34 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,34 @@
1+
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
2+
// REQUIRES: amdgpu-registered-target
3+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1030 \
4+
// RUN: -verify -DWMMA_GFX1100_TESTS -S -o - %s
5+
6+
7+
typedef float v4f __attribute__((ext_vector_type(4)));
8+
typedef half v8h __attribute__((ext_vector_type(8)));
9+
typedef half v16h __attribute__((ext_vector_type(16)));
10+
typedef int v2i __attribute__((ext_vector_type(2)));
11+
typedef int v4i __attribute__((ext_vector_type(4)));
12+
typedef short v8s __attribute__((ext_vector_type(8)));
13+
typedef short v16s __attribute__((ext_vector_type(16)));
14+
15+
#ifdef WMMA_GFX1100_TESTS
16+
17+
// Wave64
18+
19+
void test_amdgcn_wmma_f32_16x16x16_bf16_w64(global v4f* out4f, v16h a16h, v16h b16h, v4f c4f,
20+
global v8h* out8h, v16s a16s, v16s b16s, v8h c8h,
21+
global v8s* out8s, v4i a4i, v4i b4i, v8s c8s,
22+
global v4i* out4i, v2i a2i, v2i b2i, v4i c4i)
23+
{
24+
*out4f = __builtin_amdgcn_wmma_f32_16x16x16_f16_w64(a16h, b16h, c4f); // expected-error{{'__builtin_amdgcn_wmma_f32_16x16x16_f16_w64' needs target feature gfx11-insts}}
25+
*out4f = __builtin_amdgcn_wmma_f32_16x16x16_bf16_w64(a16s, b16s, c4f); // expected-error{{'__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64' needs target feature gfx11-insts}}
26+
*out8h = __builtin_amdgcn_wmma_f16_16x16x16_f16_w64(a16h, b16h, c8h, true); // expected-error{{'__builtin_amdgcn_wmma_f16_16x16x16_f16_w64' needs target feature gfx11-insts}}
27+
*out8s = __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64(a16s, b16s, c8s, true); // expected-error{{'__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64' needs target feature gfx11-insts}}
28+
*out8h = __builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w64(a16h, b16h, c8h, true); // expected-error{{'__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w64' needs target feature gfx11-insts}}
29+
*out8s = __builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w64(a16s, b16s, c8s, true); // expected-error{{'__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w64' needs target feature gfx11-insts}}
30+
*out4i = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w64(true, a4i, true, b4i, c4i, false); // expected-error{{'__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64' needs target feature gfx11-insts}}
31+
*out4i = __builtin_amdgcn_wmma_i32_16x16x16_iu4_w64(true, a2i, true, b2i, c4i, false); // expected-error{{'__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64' needs target feature gfx11-insts}}
32+
}
33+
34+
#endif

‎clang/test/CodeGenOpenCL/builtins-amdgcn-wmma-w64-param-err.cl‎

Lines changed: 34 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,34 @@
1+
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
2+
// REQUIRES: amdgpu-registered-target
3+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 \
4+
// RUN: -verify -DWMMA_GFX1100_TESTS -S -o - %s
5+
6+
7+
typedef float v4f __attribute__((ext_vector_type(4)));
8+
typedef half v8h __attribute__((ext_vector_type(8)));
9+
typedef half v16h __attribute__((ext_vector_type(16)));
10+
typedef int v2i __attribute__((ext_vector_type(2)));
11+
typedef int v4i __attribute__((ext_vector_type(4)));
12+
typedef short v8s __attribute__((ext_vector_type(8)));
13+
typedef short v16s __attribute__((ext_vector_type(16)));
14+
15+
#ifdef WMMA_GFX1100_TESTS
16+
17+
// Wave64
18+
19+
void test_amdgcn_wmma_f32_16x16x16_bf16_w64(v16h a16h, v16h b16h, unsigned int i,
20+
global v8h* out8h, v16s a16s, v16s b16s, v8h c8h,
21+
global v8s* out8s, v4i a4i, v4i b4i, v8s c8s,
22+
global v4i* out4i, v2i a2i, v2i b2i, v4i c4i)
23+
{
24+
*out8h = __builtin_amdgcn_wmma_f16_16x16x16_f16_w64(a16h, b16h, c8h, i); // expected-error{{argument to '__builtin_amdgcn_wmma_f16_16x16x16_f16_w64' must be a constant integer}}
25+
*out8s = __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64(a16s, b16s, c8s, i); // expected-error{{argument to '__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64' must be a constant integer}}
26+
*out8h = __builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w64(a16h, b16h, c8h, i); // expected-error{{argument to '__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w64' must be a constant integer}}
27+
*out8s = __builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w64(a16s, b16s, c8s, i); // expected-error{{argument to '__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w64' must be a constant integer}}
28+
*out4i = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w64(i, a4i, true, b4i, c4i, false); // expected-error{{argument to '__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64' must be a constant integer}}
29+
*out4i = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w64(true, a4i, i, b4i, c4i, false); // expected-error{{argument to '__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64' must be a constant integer}}
30+
*out4i = __builtin_amdgcn_wmma_i32_16x16x16_iu4_w64(i, a2i, true, b2i, c4i, false); // expected-error{{argument to '__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64' must be a constant integer}}
31+
*out4i = __builtin_amdgcn_wmma_i32_16x16x16_iu4_w64(true, a2i, i, b2i, c4i, false); // expected-error{{argument to '__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64' must be a constant integer}}
32+
}
33+
34+
#endif

‎clang/test/CodeGenOpenCL/builtins-amdgcn-wmma-w64.cl‎

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -76,6 +76,36 @@ void test_amdgcn_wmma_bf16_16x16x16_bf16_w64(global v8s* out, v16s a, v16s b, v8
7676
*out = __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64(a, b, c, true);
7777
}
7878

79+
//
80+
// amdgcn_wmma_f16_16x16x16_f16_tied
81+
//
82+
83+
// CHECK-GFX1100-LABEL: @test_amdgcn_wmma_f16_16x16x16_f16_tied_w64(
84+
// CHECK-GFX1100-NEXT: entry:
85+
// CHECK-GFX1100-NEXT: [[TMP0:%.*]] = tail call <8 x half> @llvm.amdgcn.wmma.f16.16x16x16.f16.tied.v8f16(<16 x half> [[A:%.*]], <16 x half> [[B:%.*]], <8 x half> [[C:%.*]], i1 true)
86+
// CHECK-GFX1100-NEXT: store <8 x half> [[TMP0]], ptr addrspace(1) [[OUT:%.*]], align 16, !tbaa [[TBAA4]]
87+
// CHECK-GFX1100-NEXT: ret void
88+
//
89+
void test_amdgcn_wmma_f16_16x16x16_f16_tied_w64(global v8h* out, v16h a, v16h b, v8h c)
90+
{
91+
*out = __builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w64(a, b, c, true);
92+
}
93+
94+
//
95+
// amdgcn_wmma_bf16_16x16x16_bf16_tied
96+
//
97+
98+
// CHECK-GFX1100-LABEL: @test_amdgcn_wmma_bf16_16x16x16_bf16_tied_w64(
99+
// CHECK-GFX1100-NEXT: entry:
100+
// CHECK-GFX1100-NEXT: [[TMP0:%.*]] = tail call <8 x i16> @llvm.amdgcn.wmma.bf16.16x16x16.bf16.tied.v8i16(<16 x i16> [[A:%.*]], <16 x i16> [[B:%.*]], <8 x i16> [[C:%.*]], i1 true)
101+
// CHECK-GFX1100-NEXT: store <8 x i16> [[TMP0]], ptr addrspace(1) [[OUT:%.*]], align 16, !tbaa [[TBAA4]]
102+
// CHECK-GFX1100-NEXT: ret void
103+
//
104+
void test_amdgcn_wmma_bf16_16x16x16_bf16_tied_w64(global v8s* out, v16s a, v16s b, v8s c)
105+
{
106+
*out = __builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w64(a, b, c, true);
107+
}
108+
79109
//
80110
// amdgcn_wmma_i32_16x16x16_iu8
81111
//

0 commit comments

Comments
 (0)