diff options
Diffstat (limited to 'llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.bf16.ll')
| -rw-r--r-- | llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.bf16.ll | 587 |
1 files changed, 497 insertions, 90 deletions
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.bf16.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.bf16.ll index 1585a2c..303ea50 100644 --- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.bf16.ll +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.bf16.ll @@ -1,6 +1,7 @@ -; RUN: llc -mtriple=amdgcn -mcpu=gfx908 < %s | FileCheck -enable-var-scope --check-prefixes=GCN,GFX908 %s -; RUN: llc -mtriple=amdgcn -mcpu=gfx908 -mattr=-mfma-inline-literal-bug < %s | FileCheck -enable-var-scope --check-prefixes=GCN,GFX908 %s -; RUN: llc -mtriple=amdgcn -mcpu=gfx90a < %s | FileCheck -enable-var-scope --check-prefixes=GCN,GFX90A %s +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 +; RUN: llc -mtriple=amdgcn -mcpu=gfx908 < %s | FileCheck --check-prefixes=GCN,GFX908 %s +; RUN: llc -mtriple=amdgcn -mcpu=gfx908 -mattr=-mfma-inline-literal-bug < %s | FileCheck --check-prefixes=GCN,GFX908 %s +; RUN: llc -mtriple=amdgcn -mcpu=gfx90a < %s | FileCheck --check-prefixes=GCN,GFX90A %s declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x2bf16(<2 x i16>, <2 x i16>, <32 x float>, i32, i32, i32) declare <16 x float> @llvm.amdgcn.mfma.f32.16x16x2bf16(<2 x i16>, <2 x i16>, <16 x float>, i32, i32, i32) @@ -9,50 +10,199 @@ declare <16 x float> @llvm.amdgcn.mfma.f32.32x32x4bf16(<2 x i16>, <2 x i16>, <16 declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x8bf16(<2 x i16>, <2 x i16>, <4 x float>, i32, i32, i32) declare i32 @llvm.amdgcn.workitem.id.x() -; GCN-LABEL: {{^}}test_mfma_f32_32x32x2bf16: -; GCN-DAG: v_mov_b32_e32 [[TWO:v[0-9]+]], 2 -; GCN-DAG: v_mov_b32_e32 [[ONE:v[0-9]+]], 1 -; GCN-DAG: s_load_dwordx16 -; GCN-DAG: s_load_dwordx16 -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX90A-COUNT-32: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}} -; GCN: v_mfma_f32_32x32x2bf16 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3 -; GFX908-COUNT-32: v_accvgpr_read_b32 -; GFX908: global_store_dwordx4 -; GFX90A-NOT: v_accvgpr_read_b32 -; GFX90A-COUNT-8: global_store_dwordx4 v{{[0-9]+}}, a[{{[0-9:]+}}], define amdgpu_kernel void @test_mfma_f32_32x32x2bf16(ptr addrspace(1) %arg) #0 { +; GFX908-LABEL: test_mfma_f32_32x32x2bf16: +; GFX908: ; %bb.0: ; %bb +; GFX908-NEXT: s_load_dwordx2 s[34:35], s[4:5], 0x24 +; GFX908-NEXT: v_mov_b32_e32 v4, 0 +; GFX908-NEXT: s_waitcnt lgkmcnt(0) +; GFX908-NEXT: s_load_dwordx16 s[16:31], s[34:35], 0x0 +; GFX908-NEXT: s_load_dwordx16 s[0:15], s[34:35], 0x40 +; GFX908-NEXT: s_waitcnt lgkmcnt(0) +; GFX908-NEXT: v_mov_b32_e32 v0, s16 +; GFX908-NEXT: v_mov_b32_e32 v1, s17 +; GFX908-NEXT: v_mov_b32_e32 v2, s18 +; GFX908-NEXT: v_accvgpr_write_b32 a0, v0 +; GFX908-NEXT: v_accvgpr_write_b32 a1, v1 +; GFX908-NEXT: v_accvgpr_write_b32 a2, v2 +; GFX908-NEXT: v_mov_b32_e32 v0, s21 +; GFX908-NEXT: v_mov_b32_e32 v1, s22 +; GFX908-NEXT: v_mov_b32_e32 v2, s23 +; GFX908-NEXT: v_accvgpr_write_b32 a5, v0 +; GFX908-NEXT: v_accvgpr_write_b32 a6, v1 +; GFX908-NEXT: v_accvgpr_write_b32 a7, v2 +; GFX908-NEXT: v_mov_b32_e32 v0, s24 +; GFX908-NEXT: v_mov_b32_e32 v1, s25 +; GFX908-NEXT: v_mov_b32_e32 v2, s26 +; GFX908-NEXT: v_accvgpr_write_b32 a8, v0 +; GFX908-NEXT: v_accvgpr_write_b32 a9, v1 +; GFX908-NEXT: v_accvgpr_write_b32 a10, v2 +; GFX908-NEXT: v_mov_b32_e32 v0, s27 +; GFX908-NEXT: v_mov_b32_e32 v1, s28 +; GFX908-NEXT: v_mov_b32_e32 v2, s29 +; GFX908-NEXT: v_accvgpr_write_b32 a11, v0 +; GFX908-NEXT: v_accvgpr_write_b32 a12, v1 +; GFX908-NEXT: v_accvgpr_write_b32 a13, v2 +; GFX908-NEXT: v_mov_b32_e32 v0, s30 +; GFX908-NEXT: v_mov_b32_e32 v1, s31 +; GFX908-NEXT: v_mov_b32_e32 v2, s0 +; GFX908-NEXT: v_accvgpr_write_b32 a14, v0 +; GFX908-NEXT: v_accvgpr_write_b32 a15, v1 +; GFX908-NEXT: v_accvgpr_write_b32 a16, v2 +; GFX908-NEXT: v_mov_b32_e32 v0, s1 +; GFX908-NEXT: v_mov_b32_e32 v1, s2 +; GFX908-NEXT: v_mov_b32_e32 v2, s3 +; GFX908-NEXT: v_accvgpr_write_b32 a17, v0 +; GFX908-NEXT: v_accvgpr_write_b32 a18, v1 +; GFX908-NEXT: v_accvgpr_write_b32 a19, v2 +; GFX908-NEXT: v_mov_b32_e32 v0, s4 +; GFX908-NEXT: v_mov_b32_e32 v1, s5 +; GFX908-NEXT: v_mov_b32_e32 v2, s6 +; GFX908-NEXT: v_accvgpr_write_b32 a20, v0 +; GFX908-NEXT: v_accvgpr_write_b32 a21, v1 +; GFX908-NEXT: v_accvgpr_write_b32 a22, v2 +; GFX908-NEXT: v_mov_b32_e32 v0, s7 +; GFX908-NEXT: v_mov_b32_e32 v1, s8 +; GFX908-NEXT: v_mov_b32_e32 v2, s9 +; GFX908-NEXT: v_mov_b32_e32 v3, s19 +; GFX908-NEXT: v_accvgpr_write_b32 a23, v0 +; GFX908-NEXT: v_accvgpr_write_b32 a24, v1 +; GFX908-NEXT: v_accvgpr_write_b32 a25, v2 +; GFX908-NEXT: v_mov_b32_e32 v0, s10 +; GFX908-NEXT: v_mov_b32_e32 v1, s11 +; GFX908-NEXT: v_mov_b32_e32 v2, s12 +; GFX908-NEXT: v_mov_b32_e32 v5, s20 +; GFX908-NEXT: v_accvgpr_write_b32 a3, v3 +; GFX908-NEXT: v_accvgpr_write_b32 a26, v0 +; GFX908-NEXT: v_accvgpr_write_b32 a27, v1 +; GFX908-NEXT: v_accvgpr_write_b32 a28, v2 +; GFX908-NEXT: v_mov_b32_e32 v0, s13 +; GFX908-NEXT: v_mov_b32_e32 v1, s14 +; GFX908-NEXT: v_mov_b32_e32 v2, s15 +; GFX908-NEXT: v_mov_b32_e32 v3, 1 +; GFX908-NEXT: v_accvgpr_write_b32 a4, v5 +; GFX908-NEXT: v_accvgpr_write_b32 a29, v0 +; GFX908-NEXT: v_accvgpr_write_b32 a30, v1 +; GFX908-NEXT: v_accvgpr_write_b32 a31, v2 +; GFX908-NEXT: v_mov_b32_e32 v0, 2 +; GFX908-NEXT: s_nop 1 +; GFX908-NEXT: v_mfma_f32_32x32x2bf16 a[0:31], v3, v0, a[0:31] cbsz:1 abid:2 blgp:3 +; GFX908-NEXT: s_nop 7 +; GFX908-NEXT: s_nop 7 +; GFX908-NEXT: s_nop 1 +; GFX908-NEXT: v_accvgpr_read_b32 v3, a27 +; GFX908-NEXT: v_accvgpr_read_b32 v2, a26 +; GFX908-NEXT: v_accvgpr_read_b32 v1, a25 +; GFX908-NEXT: v_accvgpr_read_b32 v0, a24 +; GFX908-NEXT: s_nop 1 +; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:96 +; GFX908-NEXT: s_nop 0 +; GFX908-NEXT: v_accvgpr_read_b32 v3, a31 +; GFX908-NEXT: v_accvgpr_read_b32 v2, a30 +; GFX908-NEXT: v_accvgpr_read_b32 v1, a29 +; GFX908-NEXT: v_accvgpr_read_b32 v0, a28 +; GFX908-NEXT: s_nop 1 +; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:112 +; GFX908-NEXT: s_nop 0 +; GFX908-NEXT: v_accvgpr_read_b32 v3, a19 +; GFX908-NEXT: v_accvgpr_read_b32 v2, a18 +; GFX908-NEXT: v_accvgpr_read_b32 v1, a17 +; GFX908-NEXT: v_accvgpr_read_b32 v0, a16 +; GFX908-NEXT: s_nop 1 +; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:64 +; GFX908-NEXT: s_nop 0 +; GFX908-NEXT: v_accvgpr_read_b32 v3, a23 +; GFX908-NEXT: v_accvgpr_read_b32 v2, a22 +; GFX908-NEXT: v_accvgpr_read_b32 v1, a21 +; GFX908-NEXT: v_accvgpr_read_b32 v0, a20 +; GFX908-NEXT: s_nop 1 +; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:80 +; GFX908-NEXT: s_nop 0 +; GFX908-NEXT: v_accvgpr_read_b32 v3, a11 +; GFX908-NEXT: v_accvgpr_read_b32 v2, a10 +; GFX908-NEXT: v_accvgpr_read_b32 v1, a9 +; GFX908-NEXT: v_accvgpr_read_b32 v0, a8 +; GFX908-NEXT: s_nop 1 +; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:32 +; GFX908-NEXT: s_nop 0 +; GFX908-NEXT: v_accvgpr_read_b32 v3, a15 +; GFX908-NEXT: v_accvgpr_read_b32 v2, a14 +; GFX908-NEXT: v_accvgpr_read_b32 v1, a13 +; GFX908-NEXT: v_accvgpr_read_b32 v0, a12 +; GFX908-NEXT: s_nop 1 +; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:48 +; GFX908-NEXT: s_nop 0 +; GFX908-NEXT: v_accvgpr_read_b32 v3, a3 +; GFX908-NEXT: v_accvgpr_read_b32 v2, a2 +; GFX908-NEXT: v_accvgpr_read_b32 v1, a1 +; GFX908-NEXT: v_accvgpr_read_b32 v0, a0 +; GFX908-NEXT: s_nop 1 +; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] +; GFX908-NEXT: s_nop 0 +; GFX908-NEXT: v_accvgpr_read_b32 v3, a7 +; GFX908-NEXT: v_accvgpr_read_b32 v2, a6 +; GFX908-NEXT: v_accvgpr_read_b32 v1, a5 +; GFX908-NEXT: v_accvgpr_read_b32 v0, a4 +; GFX908-NEXT: s_nop 1 +; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:16 +; GFX908-NEXT: s_endpgm +; +; GFX90A-LABEL: test_mfma_f32_32x32x2bf16: +; GFX90A: ; %bb.0: ; %bb +; GFX90A-NEXT: s_load_dwordx2 s[34:35], s[4:5], 0x24 +; GFX90A-NEXT: v_mov_b32_e32 v1, 1 +; GFX90A-NEXT: v_mov_b32_e32 v2, 2 +; GFX90A-NEXT: v_mov_b32_e32 v0, 0 +; GFX90A-NEXT: s_waitcnt lgkmcnt(0) +; GFX90A-NEXT: s_load_dwordx16 s[16:31], s[34:35], 0x0 +; GFX90A-NEXT: s_load_dwordx16 s[0:15], s[34:35], 0x40 +; GFX90A-NEXT: s_waitcnt lgkmcnt(0) +; GFX90A-NEXT: v_accvgpr_write_b32 a0, s16 +; GFX90A-NEXT: v_accvgpr_write_b32 a1, s17 +; GFX90A-NEXT: v_accvgpr_write_b32 a2, s18 +; GFX90A-NEXT: v_accvgpr_write_b32 a3, s19 +; GFX90A-NEXT: v_accvgpr_write_b32 a4, s20 +; GFX90A-NEXT: v_accvgpr_write_b32 a5, s21 +; GFX90A-NEXT: v_accvgpr_write_b32 a6, s22 +; GFX90A-NEXT: v_accvgpr_write_b32 a7, s23 +; GFX90A-NEXT: v_accvgpr_write_b32 a8, s24 +; GFX90A-NEXT: v_accvgpr_write_b32 a9, s25 +; GFX90A-NEXT: v_accvgpr_write_b32 a10, s26 +; GFX90A-NEXT: v_accvgpr_write_b32 a11, s27 +; GFX90A-NEXT: v_accvgpr_write_b32 a12, s28 +; GFX90A-NEXT: v_accvgpr_write_b32 a13, s29 +; GFX90A-NEXT: v_accvgpr_write_b32 a14, s30 +; GFX90A-NEXT: v_accvgpr_write_b32 a15, s31 +; GFX90A-NEXT: v_accvgpr_write_b32 a16, s0 +; GFX90A-NEXT: v_accvgpr_write_b32 a17, s1 +; GFX90A-NEXT: v_accvgpr_write_b32 a18, s2 +; GFX90A-NEXT: v_accvgpr_write_b32 a19, s3 +; GFX90A-NEXT: v_accvgpr_write_b32 a20, s4 +; GFX90A-NEXT: v_accvgpr_write_b32 a21, s5 +; GFX90A-NEXT: v_accvgpr_write_b32 a22, s6 +; GFX90A-NEXT: v_accvgpr_write_b32 a23, s7 +; GFX90A-NEXT: v_accvgpr_write_b32 a24, s8 +; GFX90A-NEXT: v_accvgpr_write_b32 a25, s9 +; GFX90A-NEXT: v_accvgpr_write_b32 a26, s10 +; GFX90A-NEXT: v_accvgpr_write_b32 a27, s11 +; GFX90A-NEXT: v_accvgpr_write_b32 a28, s12 +; GFX90A-NEXT: v_accvgpr_write_b32 a29, s13 +; GFX90A-NEXT: v_accvgpr_write_b32 a30, s14 +; GFX90A-NEXT: v_accvgpr_write_b32 a31, s15 +; GFX90A-NEXT: s_nop 1 +; GFX90A-NEXT: v_mfma_f32_32x32x2bf16 a[0:31], v1, v2, a[0:31] cbsz:1 abid:2 blgp:3 +; GFX90A-NEXT: s_nop 7 +; GFX90A-NEXT: s_nop 7 +; GFX90A-NEXT: s_nop 2 +; GFX90A-NEXT: global_store_dwordx4 v0, a[24:27], s[34:35] offset:96 +; GFX90A-NEXT: global_store_dwordx4 v0, a[28:31], s[34:35] offset:112 +; GFX90A-NEXT: global_store_dwordx4 v0, a[16:19], s[34:35] offset:64 +; GFX90A-NEXT: global_store_dwordx4 v0, a[20:23], s[34:35] offset:80 +; GFX90A-NEXT: global_store_dwordx4 v0, a[8:11], s[34:35] offset:32 +; GFX90A-NEXT: global_store_dwordx4 v0, a[12:15], s[34:35] offset:48 +; GFX90A-NEXT: global_store_dwordx4 v0, a[0:3], s[34:35] +; GFX90A-NEXT: global_store_dwordx4 v0, a[4:7], s[34:35] offset:16 +; GFX90A-NEXT: s_endpgm bb: %in.1 = load <32 x float>, ptr addrspace(1) %arg %a = bitcast i32 1 to <2 x i16> @@ -62,18 +212,109 @@ bb: ret void } -; GCN-LABEL: {{^}}test_mfma_f32_16x16x2bf16: -; GCN-DAG: v_mov_b32_e32 [[TWO:v[0-9]+]], 2 -; GCN-DAG: v_mov_b32_e32 [[ONE:v[0-9]+]], 1 -; GCN-DAG: s_load_dwordx16 -; GFX908-DAG-COUNT-16: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX90A-COUNT-16: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}} -; GCN: v_mfma_f32_16x16x2bf16 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3 -; GFX908-COUNT-16: v_accvgpr_read_b32 -; GFX908: global_store_dwordx4 -; GFX90A-NOT: v_accvgpr_read_b32 -; GFX90A-COUNT-4: global_store_dwordx4 v{{[0-9]+}}, a[{{[0-9:]+}}], define amdgpu_kernel void @test_mfma_f32_16x16x2bf16(ptr addrspace(1) %arg) #0 { +; GFX908-LABEL: test_mfma_f32_16x16x2bf16: +; GFX908: ; %bb.0: ; %bb +; GFX908-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24 +; GFX908-NEXT: v_mov_b32_e32 v0, 1 +; GFX908-NEXT: v_mov_b32_e32 v12, 0 +; GFX908-NEXT: s_waitcnt lgkmcnt(0) +; GFX908-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0 +; GFX908-NEXT: s_waitcnt lgkmcnt(0) +; GFX908-NEXT: v_mov_b32_e32 v13, s0 +; GFX908-NEXT: v_mov_b32_e32 v1, s1 +; GFX908-NEXT: v_mov_b32_e32 v2, s2 +; GFX908-NEXT: v_accvgpr_write_b32 a0, v13 +; GFX908-NEXT: v_mov_b32_e32 v13, s3 +; GFX908-NEXT: v_accvgpr_write_b32 a1, v1 +; GFX908-NEXT: v_accvgpr_write_b32 a2, v2 +; GFX908-NEXT: v_accvgpr_write_b32 a3, v13 +; GFX908-NEXT: v_mov_b32_e32 v1, s4 +; GFX908-NEXT: v_mov_b32_e32 v2, s5 +; GFX908-NEXT: v_mov_b32_e32 v13, s6 +; GFX908-NEXT: v_accvgpr_write_b32 a4, v1 +; GFX908-NEXT: v_accvgpr_write_b32 a5, v2 +; GFX908-NEXT: v_accvgpr_write_b32 a6, v13 +; GFX908-NEXT: v_mov_b32_e32 v1, s7 +; GFX908-NEXT: v_mov_b32_e32 v2, s8 +; GFX908-NEXT: v_mov_b32_e32 v13, s9 +; GFX908-NEXT: v_accvgpr_write_b32 a7, v1 +; GFX908-NEXT: v_accvgpr_write_b32 a8, v2 +; GFX908-NEXT: v_accvgpr_write_b32 a9, v13 +; GFX908-NEXT: v_mov_b32_e32 v1, s10 +; GFX908-NEXT: v_mov_b32_e32 v2, s11 +; GFX908-NEXT: v_mov_b32_e32 v13, s12 +; GFX908-NEXT: v_accvgpr_write_b32 a10, v1 +; GFX908-NEXT: v_accvgpr_write_b32 a11, v2 +; GFX908-NEXT: v_accvgpr_write_b32 a12, v13 +; GFX908-NEXT: v_mov_b32_e32 v1, s13 +; GFX908-NEXT: v_mov_b32_e32 v2, s14 +; GFX908-NEXT: v_mov_b32_e32 v13, s15 +; GFX908-NEXT: v_accvgpr_write_b32 a13, v1 +; GFX908-NEXT: v_accvgpr_write_b32 a14, v2 +; GFX908-NEXT: v_accvgpr_write_b32 a15, v13 +; GFX908-NEXT: v_mov_b32_e32 v1, 2 +; GFX908-NEXT: s_nop 1 +; GFX908-NEXT: v_mfma_f32_16x16x2bf16 a[0:15], v0, v1, a[0:15] cbsz:1 abid:2 blgp:3 +; GFX908-NEXT: s_nop 7 +; GFX908-NEXT: s_nop 1 +; GFX908-NEXT: v_accvgpr_read_b32 v3, a15 +; GFX908-NEXT: v_accvgpr_read_b32 v2, a14 +; GFX908-NEXT: v_accvgpr_read_b32 v1, a13 +; GFX908-NEXT: v_accvgpr_read_b32 v0, a12 +; GFX908-NEXT: v_accvgpr_read_b32 v7, a11 +; GFX908-NEXT: v_accvgpr_read_b32 v6, a10 +; GFX908-NEXT: v_accvgpr_read_b32 v5, a9 +; GFX908-NEXT: v_accvgpr_read_b32 v4, a8 +; GFX908-NEXT: v_accvgpr_read_b32 v11, a7 +; GFX908-NEXT: v_accvgpr_read_b32 v10, a6 +; GFX908-NEXT: v_accvgpr_read_b32 v9, a5 +; GFX908-NEXT: v_accvgpr_read_b32 v8, a4 +; GFX908-NEXT: global_store_dwordx4 v12, v[0:3], s[16:17] offset:48 +; GFX908-NEXT: global_store_dwordx4 v12, v[4:7], s[16:17] offset:32 +; GFX908-NEXT: global_store_dwordx4 v12, v[8:11], s[16:17] offset:16 +; GFX908-NEXT: v_accvgpr_read_b32 v3, a3 +; GFX908-NEXT: v_accvgpr_read_b32 v2, a2 +; GFX908-NEXT: v_accvgpr_read_b32 v1, a1 +; GFX908-NEXT: v_accvgpr_read_b32 v0, a0 +; GFX908-NEXT: s_nop 1 +; GFX908-NEXT: global_store_dwordx4 v12, v[0:3], s[16:17] +; GFX908-NEXT: s_endpgm +; +; GFX90A-LABEL: test_mfma_f32_16x16x2bf16: +; GFX90A: ; %bb.0: ; %bb +; GFX90A-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24 +; GFX90A-NEXT: v_mov_b32_e32 v0, 1 +; GFX90A-NEXT: v_mov_b32_e32 v1, 2 +; GFX90A-NEXT: s_waitcnt lgkmcnt(0) +; GFX90A-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0 +; GFX90A-NEXT: s_waitcnt lgkmcnt(0) +; GFX90A-NEXT: v_accvgpr_write_b32 a0, s0 +; GFX90A-NEXT: v_accvgpr_write_b32 a1, s1 +; GFX90A-NEXT: v_accvgpr_write_b32 a2, s2 +; GFX90A-NEXT: v_accvgpr_write_b32 a3, s3 +; GFX90A-NEXT: v_accvgpr_write_b32 a4, s4 +; GFX90A-NEXT: v_accvgpr_write_b32 a5, s5 +; GFX90A-NEXT: v_accvgpr_write_b32 a6, s6 +; GFX90A-NEXT: v_accvgpr_write_b32 a7, s7 +; GFX90A-NEXT: v_accvgpr_write_b32 a8, s8 +; GFX90A-NEXT: v_accvgpr_write_b32 a9, s9 +; GFX90A-NEXT: v_accvgpr_write_b32 a10, s10 +; GFX90A-NEXT: v_accvgpr_write_b32 a11, s11 +; GFX90A-NEXT: v_accvgpr_write_b32 a12, s12 +; GFX90A-NEXT: v_accvgpr_write_b32 a13, s13 +; GFX90A-NEXT: v_accvgpr_write_b32 a14, s14 +; GFX90A-NEXT: v_accvgpr_write_b32 a15, s15 +; GFX90A-NEXT: s_nop 1 +; GFX90A-NEXT: v_mfma_f32_16x16x2bf16 a[0:15], v0, v1, a[0:15] cbsz:1 abid:2 blgp:3 +; GFX90A-NEXT: v_mov_b32_e32 v0, 0 +; GFX90A-NEXT: s_nop 7 +; GFX90A-NEXT: s_nop 1 +; GFX90A-NEXT: global_store_dwordx4 v0, a[12:15], s[16:17] offset:48 +; GFX90A-NEXT: global_store_dwordx4 v0, a[8:11], s[16:17] offset:32 +; GFX90A-NEXT: global_store_dwordx4 v0, a[4:7], s[16:17] offset:16 +; GFX90A-NEXT: global_store_dwordx4 v0, a[0:3], s[16:17] +; GFX90A-NEXT: s_endpgm bb: %in.1 = load <16 x float>, ptr addrspace(1) %arg %a = bitcast i32 1 to <2 x i16> @@ -83,18 +324,53 @@ bb: ret void } -; GCN-LABEL: {{^}}test_mfma_f32_4x4x2bf16: -; GCN-DAG: v_mov_b32_e32 [[TWO:v[0-9]+]], 2 -; GCN-DAG: v_mov_b32_e32 [[ONE:v[0-9]+]], 1 -; GCN: s_load_dwordx4 -; GFX908-COUNT-4: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX90A-COUNT-4: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}} -; GCN: v_mfma_f32_4x4x2bf16 [[RES:a\[[0-9]+:[0-9]+\]]], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3 -; GFX908-COUNT-4: v_accvgpr_read_b32 -; GFX908: global_store_dwordx4 -; GFX90A-NOT: v_accvgpr_read_b32 -; GFX90A: global_store_dwordx4 v{{[0-9]+}}, [[RES]], define amdgpu_kernel void @test_mfma_f32_4x4x2bf16(ptr addrspace(1) %arg) #0 { +; GFX908-LABEL: test_mfma_f32_4x4x2bf16: +; GFX908: ; %bb.0: ; %bb +; GFX908-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24 +; GFX908-NEXT: v_mov_b32_e32 v0, 1 +; GFX908-NEXT: v_mov_b32_e32 v1, 2 +; GFX908-NEXT: v_mov_b32_e32 v4, 0 +; GFX908-NEXT: s_waitcnt lgkmcnt(0) +; GFX908-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0 +; GFX908-NEXT: s_waitcnt lgkmcnt(0) +; GFX908-NEXT: v_mov_b32_e32 v5, s0 +; GFX908-NEXT: v_mov_b32_e32 v2, s1 +; GFX908-NEXT: v_mov_b32_e32 v3, s2 +; GFX908-NEXT: v_accvgpr_write_b32 a0, v5 +; GFX908-NEXT: v_mov_b32_e32 v5, s3 +; GFX908-NEXT: v_accvgpr_write_b32 a1, v2 +; GFX908-NEXT: v_accvgpr_write_b32 a2, v3 +; GFX908-NEXT: v_accvgpr_write_b32 a3, v5 +; GFX908-NEXT: s_nop 0 +; GFX908-NEXT: v_mfma_f32_4x4x2bf16 a[0:3], v0, v1, a[0:3] cbsz:1 abid:2 blgp:3 +; GFX908-NEXT: s_nop 3 +; GFX908-NEXT: v_accvgpr_read_b32 v0, a0 +; GFX908-NEXT: v_accvgpr_read_b32 v1, a1 +; GFX908-NEXT: v_accvgpr_read_b32 v2, a2 +; GFX908-NEXT: v_accvgpr_read_b32 v3, a3 +; GFX908-NEXT: s_nop 1 +; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[6:7] +; GFX908-NEXT: s_endpgm +; +; GFX90A-LABEL: test_mfma_f32_4x4x2bf16: +; GFX90A: ; %bb.0: ; %bb +; GFX90A-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24 +; GFX90A-NEXT: v_mov_b32_e32 v0, 1 +; GFX90A-NEXT: v_mov_b32_e32 v2, 2 +; GFX90A-NEXT: v_mov_b32_e32 v1, 0 +; GFX90A-NEXT: s_waitcnt lgkmcnt(0) +; GFX90A-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0 +; GFX90A-NEXT: s_waitcnt lgkmcnt(0) +; GFX90A-NEXT: v_accvgpr_write_b32 a0, s0 +; GFX90A-NEXT: v_accvgpr_write_b32 a1, s1 +; GFX90A-NEXT: v_accvgpr_write_b32 a2, s2 +; GFX90A-NEXT: v_accvgpr_write_b32 a3, s3 +; GFX90A-NEXT: s_nop 1 +; GFX90A-NEXT: v_mfma_f32_4x4x2bf16 a[0:3], v0, v2, a[0:3] cbsz:1 abid:2 blgp:3 +; GFX90A-NEXT: s_nop 4 +; GFX90A-NEXT: global_store_dwordx4 v1, a[0:3], s[6:7] +; GFX90A-NEXT: s_endpgm bb: %in.1 = load <4 x float>, ptr addrspace(1) %arg %a = bitcast i32 1 to <2 x i16> @@ -104,18 +380,110 @@ bb: ret void } -; GCN-LABEL: {{^}}test_mfma_f32_32x32x4bf16: -; GCN-DAG: v_mov_b32_e32 [[TWO:v[0-9]+]], 2 -; GCN-DAG: v_mov_b32_e32 [[ONE:v[0-9]+]], 1 -; GCN-DAG: s_load_dwordx16 -; GFX908-DAG-COUNT-16: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX90A-COUNT-4: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}} -; GCN: v_mfma_f32_32x32x4bf16 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3 -; GFX908-COUNT-16: v_accvgpr_read_b32 -; GFX908: global_store_dwordx4 -; GFX90A-NOT: v_accvgpr_read_b32 -; GFX90A-COUNT-4: global_store_dwordx4 v{{[0-9]+}}, a[{{[0-9:]+}}], define amdgpu_kernel void @test_mfma_f32_32x32x4bf16(ptr addrspace(1) %arg) #0 { +; GFX908-LABEL: test_mfma_f32_32x32x4bf16: +; GFX908: ; %bb.0: ; %bb +; GFX908-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24 +; GFX908-NEXT: v_mov_b32_e32 v0, 1 +; GFX908-NEXT: v_mov_b32_e32 v16, 0 +; GFX908-NEXT: s_waitcnt lgkmcnt(0) +; GFX908-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0 +; GFX908-NEXT: s_waitcnt lgkmcnt(0) +; GFX908-NEXT: v_mov_b32_e32 v17, s0 +; GFX908-NEXT: v_mov_b32_e32 v1, s1 +; GFX908-NEXT: v_mov_b32_e32 v2, s2 +; GFX908-NEXT: v_accvgpr_write_b32 a0, v17 +; GFX908-NEXT: v_mov_b32_e32 v17, s3 +; GFX908-NEXT: v_accvgpr_write_b32 a1, v1 +; GFX908-NEXT: v_accvgpr_write_b32 a2, v2 +; GFX908-NEXT: v_accvgpr_write_b32 a3, v17 +; GFX908-NEXT: v_mov_b32_e32 v1, s4 +; GFX908-NEXT: v_mov_b32_e32 v2, s5 +; GFX908-NEXT: v_mov_b32_e32 v17, s6 +; GFX908-NEXT: v_accvgpr_write_b32 a4, v1 +; GFX908-NEXT: v_accvgpr_write_b32 a5, v2 +; GFX908-NEXT: v_accvgpr_write_b32 a6, v17 +; GFX908-NEXT: v_mov_b32_e32 v1, s7 +; GFX908-NEXT: v_mov_b32_e32 v2, s8 +; GFX908-NEXT: v_mov_b32_e32 v17, s9 +; GFX908-NEXT: v_accvgpr_write_b32 a7, v1 +; GFX908-NEXT: v_accvgpr_write_b32 a8, v2 +; GFX908-NEXT: v_accvgpr_write_b32 a9, v17 +; GFX908-NEXT: v_mov_b32_e32 v1, s10 +; GFX908-NEXT: v_mov_b32_e32 v2, s11 +; GFX908-NEXT: v_mov_b32_e32 v17, s12 +; GFX908-NEXT: v_accvgpr_write_b32 a10, v1 +; GFX908-NEXT: v_accvgpr_write_b32 a11, v2 +; GFX908-NEXT: v_accvgpr_write_b32 a12, v17 +; GFX908-NEXT: v_mov_b32_e32 v1, s13 +; GFX908-NEXT: v_mov_b32_e32 v2, s14 +; GFX908-NEXT: v_mov_b32_e32 v17, s15 +; GFX908-NEXT: v_accvgpr_write_b32 a13, v1 +; GFX908-NEXT: v_accvgpr_write_b32 a14, v2 +; GFX908-NEXT: v_accvgpr_write_b32 a15, v17 +; GFX908-NEXT: v_mov_b32_e32 v1, 2 +; GFX908-NEXT: s_nop 1 +; GFX908-NEXT: v_mfma_f32_32x32x4bf16 a[0:15], v0, v1, a[0:15] cbsz:1 abid:2 blgp:3 +; GFX908-NEXT: s_nop 7 +; GFX908-NEXT: s_nop 7 +; GFX908-NEXT: s_nop 1 +; GFX908-NEXT: v_accvgpr_read_b32 v3, a15 +; GFX908-NEXT: v_accvgpr_read_b32 v2, a14 +; GFX908-NEXT: v_accvgpr_read_b32 v1, a13 +; GFX908-NEXT: v_accvgpr_read_b32 v0, a12 +; GFX908-NEXT: v_accvgpr_read_b32 v7, a11 +; GFX908-NEXT: v_accvgpr_read_b32 v6, a10 +; GFX908-NEXT: v_accvgpr_read_b32 v5, a9 +; GFX908-NEXT: v_accvgpr_read_b32 v4, a8 +; GFX908-NEXT: v_accvgpr_read_b32 v11, a7 +; GFX908-NEXT: v_accvgpr_read_b32 v10, a6 +; GFX908-NEXT: v_accvgpr_read_b32 v9, a5 +; GFX908-NEXT: v_accvgpr_read_b32 v8, a4 +; GFX908-NEXT: v_accvgpr_read_b32 v15, a3 +; GFX908-NEXT: v_accvgpr_read_b32 v14, a2 +; GFX908-NEXT: v_accvgpr_read_b32 v13, a1 +; GFX908-NEXT: v_accvgpr_read_b32 v12, a0 +; GFX908-NEXT: global_store_dwordx4 v16, v[0:3], s[16:17] offset:48 +; GFX908-NEXT: global_store_dwordx4 v16, v[4:7], s[16:17] offset:32 +; GFX908-NEXT: global_store_dwordx4 v16, v[8:11], s[16:17] offset:16 +; GFX908-NEXT: global_store_dwordx4 v16, v[12:15], s[16:17] +; GFX908-NEXT: s_endpgm +; +; GFX90A-LABEL: test_mfma_f32_32x32x4bf16: +; GFX90A: ; %bb.0: ; %bb +; GFX90A-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24 +; GFX90A-NEXT: v_mov_b32_e32 v0, 1 +; GFX90A-NEXT: v_mov_b32_e32 v1, 2 +; GFX90A-NEXT: s_waitcnt lgkmcnt(0) +; GFX90A-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0 +; GFX90A-NEXT: s_waitcnt lgkmcnt(0) +; GFX90A-NEXT: v_accvgpr_write_b32 a0, s0 +; GFX90A-NEXT: v_accvgpr_write_b32 a1, s1 +; GFX90A-NEXT: v_accvgpr_write_b32 a2, s2 +; GFX90A-NEXT: v_accvgpr_write_b32 a3, s3 +; GFX90A-NEXT: v_accvgpr_write_b32 a4, s4 +; GFX90A-NEXT: v_accvgpr_write_b32 a5, s5 +; GFX90A-NEXT: v_accvgpr_write_b32 a6, s6 +; GFX90A-NEXT: v_accvgpr_write_b32 a7, s7 +; GFX90A-NEXT: v_accvgpr_write_b32 a8, s8 +; GFX90A-NEXT: v_accvgpr_write_b32 a9, s9 +; GFX90A-NEXT: v_accvgpr_write_b32 a10, s10 +; GFX90A-NEXT: v_accvgpr_write_b32 a11, s11 +; GFX90A-NEXT: v_accvgpr_write_b32 a12, s12 +; GFX90A-NEXT: v_accvgpr_write_b32 a13, s13 +; GFX90A-NEXT: v_accvgpr_write_b32 a14, s14 +; GFX90A-NEXT: v_accvgpr_write_b32 a15, s15 +; GFX90A-NEXT: s_nop 1 +; GFX90A-NEXT: v_mfma_f32_32x32x4bf16 a[0:15], v0, v1, a[0:15] cbsz:1 abid:2 blgp:3 +; GFX90A-NEXT: v_mov_b32_e32 v0, 0 +; GFX90A-NEXT: s_nop 7 +; GFX90A-NEXT: s_nop 7 +; GFX90A-NEXT: s_nop 1 +; GFX90A-NEXT: global_store_dwordx4 v0, a[12:15], s[16:17] offset:48 +; GFX90A-NEXT: global_store_dwordx4 v0, a[8:11], s[16:17] offset:32 +; GFX90A-NEXT: global_store_dwordx4 v0, a[4:7], s[16:17] offset:16 +; GFX90A-NEXT: global_store_dwordx4 v0, a[0:3], s[16:17] +; GFX90A-NEXT: s_endpgm bb: %in.1 = load <16 x float>, ptr addrspace(1) %arg %a = bitcast i32 1 to <2 x i16> @@ -125,18 +493,55 @@ bb: ret void } -; GCN-LABEL: {{^}}test_mfma_f32_16x16x8bf16: -; GCN-DAG: v_mov_b32_e32 [[TWO:v[0-9]+]], 2 -; GCN-DAG: v_mov_b32_e32 [[ONE:v[0-9]+]], 1 -; GCN: s_load_dwordx4 -; GFX908-COUNT-4: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} -; GFX90A-COUNT-4: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}} -; GCN: v_mfma_f32_16x16x8bf16 [[RES:a\[[0-9]+:[0-9]+\]]], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3 -; GFX908-COUNT-4: v_accvgpr_read_b32 -; GFX908: global_store_dwordx4 -; GFX90A-NOT: v_accvgpr_read_b32 -; GFX90A: global_store_dwordx4 v{{[0-9]+}}, [[RES]], define amdgpu_kernel void @test_mfma_f32_16x16x8bf16(ptr addrspace(1) %arg) #0 { +; GFX908-LABEL: test_mfma_f32_16x16x8bf16: +; GFX908: ; %bb.0: ; %bb +; GFX908-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24 +; GFX908-NEXT: v_mov_b32_e32 v0, 1 +; GFX908-NEXT: v_mov_b32_e32 v1, 2 +; GFX908-NEXT: v_mov_b32_e32 v4, 0 +; GFX908-NEXT: s_waitcnt lgkmcnt(0) +; GFX908-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0 +; GFX908-NEXT: s_waitcnt lgkmcnt(0) +; GFX908-NEXT: v_mov_b32_e32 v5, s0 +; GFX908-NEXT: v_mov_b32_e32 v2, s1 +; GFX908-NEXT: v_mov_b32_e32 v3, s2 +; GFX908-NEXT: v_accvgpr_write_b32 a0, v5 +; GFX908-NEXT: v_mov_b32_e32 v5, s3 +; GFX908-NEXT: v_accvgpr_write_b32 a1, v2 +; GFX908-NEXT: v_accvgpr_write_b32 a2, v3 +; GFX908-NEXT: v_accvgpr_write_b32 a3, v5 +; GFX908-NEXT: s_nop 0 +; GFX908-NEXT: v_mfma_f32_16x16x8bf16 a[0:3], v0, v1, a[0:3] cbsz:1 abid:2 blgp:3 +; GFX908-NEXT: s_nop 7 +; GFX908-NEXT: s_nop 1 +; GFX908-NEXT: v_accvgpr_read_b32 v0, a0 +; GFX908-NEXT: v_accvgpr_read_b32 v1, a1 +; GFX908-NEXT: v_accvgpr_read_b32 v2, a2 +; GFX908-NEXT: v_accvgpr_read_b32 v3, a3 +; GFX908-NEXT: s_nop 1 +; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[6:7] +; GFX908-NEXT: s_endpgm +; +; GFX90A-LABEL: test_mfma_f32_16x16x8bf16: +; GFX90A: ; %bb.0: ; %bb +; GFX90A-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24 +; GFX90A-NEXT: v_mov_b32_e32 v0, 1 +; GFX90A-NEXT: v_mov_b32_e32 v2, 2 +; GFX90A-NEXT: v_mov_b32_e32 v1, 0 +; GFX90A-NEXT: s_waitcnt lgkmcnt(0) +; GFX90A-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0 +; GFX90A-NEXT: s_waitcnt lgkmcnt(0) +; GFX90A-NEXT: v_accvgpr_write_b32 a0, s0 +; GFX90A-NEXT: v_accvgpr_write_b32 a1, s1 +; GFX90A-NEXT: v_accvgpr_write_b32 a2, s2 +; GFX90A-NEXT: v_accvgpr_write_b32 a3, s3 +; GFX90A-NEXT: s_nop 1 +; GFX90A-NEXT: v_mfma_f32_16x16x8bf16 a[0:3], v0, v2, a[0:3] cbsz:1 abid:2 blgp:3 +; GFX90A-NEXT: s_nop 7 +; GFX90A-NEXT: s_nop 2 +; GFX90A-NEXT: global_store_dwordx4 v1, a[0:3], s[6:7] +; GFX90A-NEXT: s_endpgm bb: %in.1 = load <4 x float>, ptr addrspace(1) %arg %a = bitcast i32 1 to <2 x i16> @@ -147,3 +552,5 @@ bb: } attributes #0 = { "amdgpu-flat-work-group-size"="1,256" } +;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line: +; GCN: {{.*}} |