llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.bf16.ll

   1 ; RUN: llc -march=amdgcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck -enable-var-scope --check-prefixes=GCN %s
   2 ; RUN: llc -march=amdgcn -mcpu=gfx908 -mattr=-mfma-inline-literal-bug -verify-machineinstrs < %s | FileCheck -enable-var-scope --check-prefixes=GCN %s
   3 ; RUN: llc -march=amdgcn -mcpu=gfx90a -verify-machineinstrs < %s | FileCheck -enable-var-scope --check-prefixes=GCN,GFX90A %s
   4
   5 declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x2bf16(<2 x i16>, <2 x i16>, <32 x float>, i32, i32, i32)
   6 declare <16 x float> @llvm.amdgcn.mfma.f32.16x16x2bf16(<2 x i16>, <2 x i16>, <16 x float>, i32, i32, i32)
   7 declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16>, <2 x i16>, <4 x float>, i32, i32, i32)
   8 declare <16 x float> @llvm.amdgcn.mfma.f32.32x32x4bf16(<2 x i16>, <2 x i16>, <16 x float>, i32, i32, i32)
   9 declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x8bf16(<2 x i16>, <2 x i16>, <4 x float>, i32, i32, i32)
  10 declare i32 @llvm.amdgcn.workitem.id.x()
  11
  12 ; GCN-LABEL: {{^}}test_mfma_f32_32x32x2bf16:
  13 ; GCN-DAG:         v_mov_b32_e32 [[TWO:v[0-9]+]], 2
  14 ; GCN-DAG:         v_mov_b32_e32 [[ONE:v[0-9]+]], 1
  15 ; GCN-DAG:         s_load_dwordx16
  16 ; GCN-DAG:         s_load_dwordx16
  17 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  18 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  19 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  20 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  21 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  22 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  23 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  24 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  25 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  26 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  27 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  28 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  29 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  30 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  31 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  32 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  33 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  34 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  35 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  36 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  37 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  38 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  39 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  40 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  41 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  42 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  43 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  44 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  45 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  46 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  47 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  48 ; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  49 ; GCN:             v_mfma_f32_32x32x2bf16 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
  50 ; GFX908-COUNT-32: v_accvgpr_read_b32
  51 ; GFX908:          global_store_dwordx4
  52 ; GFX90A-NOT:      v_accvgpr_read_b32
  53 ; GFX90A-COUNT-8:  global_store_dwordx4 v{{[0-9]+}}, a[{{[0-9:]+}}],
  54 define amdgpu_kernel void @test_mfma_f32_32x32x2bf16(<32 x float> addrspace(1)* %arg) {
  55 bb:
  56   %in.1 = load <32 x float>, <32 x float> addrspace(1)* %arg
  57   %a = bitcast i32 1 to <2 x i16>
  58   %b = bitcast i32 2 to <2 x i16>
  59   %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x2bf16(<2 x i16> %a, <2 x i16> %b, <32 x float> %in.1, i32 1, i32 2, i32 3)
  60   store <32 x float> %mai.1, <32 x float> addrspace(1)* %arg
  61   ret void
  62 }
  63
  64 ; GCN-LABEL: {{^}}test_mfma_f32_16x16x2bf16:
  65 ; GCN-DAG:         v_mov_b32_e32 [[TWO:v[0-9]+]], 2
  66 ; GCN-DAG:         v_mov_b32_e32 [[ONE:v[0-9]+]], 1
  67 ; GCN:             s_load_dwordx16
  68 ; GCN-COUNT-16:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  69 ; GCN:             v_mfma_f32_16x16x2bf16 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
  70 ; GFX908-COUNT-16: v_accvgpr_read_b32
  71 ; GFX908:          global_store_dwordx4
  72 ; GFX90A-NOT:      v_accvgpr_read_b32
  73 ; GFX90A-COUNT-4:  global_store_dwordx4 v{{[0-9]+}}, a[{{[0-9:]+}}],
  74 define amdgpu_kernel void @test_mfma_f32_16x16x2bf16(<16 x float> addrspace(1)* %arg) {
  75 bb:
  76   %in.1 = load <16 x float>, <16 x float> addrspace(1)* %arg
  77   %a = bitcast i32 1 to <2 x i16>
  78   %b = bitcast i32 2 to <2 x i16>
  79   %mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.16x16x2bf16(<2 x i16> %a, <2 x i16> %b, <16 x float> %in.1, i32 1, i32 2, i32 3)
  80   store <16 x float> %mai.1, <16 x float> addrspace(1)* %arg
  81   ret void
  82 }
  83
  84 ; GCN-LABEL: {{^}}test_mfma_f32_4x4x2bf16:
  85 ; GCN-DAG:        v_mov_b32_e32 [[TWO:v[0-9]+]], 2
  86 ; GCN-DAG:        v_mov_b32_e32 [[ONE:v[0-9]+]], 1
  87 ; GCN:            s_load_dwordx4
  88 ; GCN-COUNT-4:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
  89 ; GCN:            v_mfma_f32_4x4x2bf16 [[RES:a\[[0-9]+:[0-9]+\]]], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
  90 ; GFX908-COUNT-4: v_accvgpr_read_b32
  91 ; GFX908:         global_store_dwordx4
  92 ; GFX90A-NOT:     v_accvgpr_read_b32
  93 ; GFX90A:         global_store_dwordx4 v{{[0-9]+}}, [[RES]],
  94 define amdgpu_kernel void @test_mfma_f32_4x4x2bf16(<4 x float> addrspace(1)* %arg) {
  95 bb:
  96   %in.1 = load <4 x float>, <4 x float> addrspace(1)* %arg
  97   %a = bitcast i32 1 to <2 x i16>
  98   %b = bitcast i32 2 to <2 x i16>
  99   %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %in.1, i32 1, i32 2, i32 3)
 100   store <4 x float> %mai.1, <4 x float> addrspace(1)* %arg
 101   ret void
 102 }
 103
 104 ; GCN-LABEL: {{^}}test_mfma_f32_32x32x4bf16:
 105 ; GCN-DAG:         v_mov_b32_e32 [[TWO:v[0-9]+]], 2
 106 ; GCN-DAG:         v_mov_b32_e32 [[ONE:v[0-9]+]], 1
 107 ; GCN:             s_load_dwordx16
 108 ; GCN-COUNT-16:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
 109 ; GCN:             v_mfma_f32_32x32x4bf16 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
 110 ; GFX908-COUNT-16: v_accvgpr_read_b32
 111 ; GFX908:          global_store_dwordx4
 112 ; GFX90A-NOT:      v_accvgpr_read_b32
 113 ; GFX90A-COUNT-4:  global_store_dwordx4 v{{[0-9]+}}, a[{{[0-9:]+}}],
 114 define amdgpu_kernel void @test_mfma_f32_32x32x4bf16(<16 x float> addrspace(1)* %arg) {
 115 bb:
 116   %in.1 = load <16 x float>, <16 x float> addrspace(1)* %arg
 117   %a = bitcast i32 1 to <2 x i16>
 118   %b = bitcast i32 2 to <2 x i16>
 119   %mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x4bf16(<2 x i16> %a, <2 x i16> %b, <16 x float> %in.1, i32 1, i32 2, i32 3)
 120   store <16 x float> %mai.1, <16 x float> addrspace(1)* %arg
 121   ret void
 122 }
 123
 124 ; GCN-LABEL: {{^}}test_mfma_f32_16x16x8bf16:
 125 ; GCN-DAG:        v_mov_b32_e32 [[TWO:v[0-9]+]], 2
 126 ; GCN-DAG:        v_mov_b32_e32 [[ONE:v[0-9]+]], 1
 127 ; GCN:            s_load_dwordx4
 128 ; GCN-COUNT-4:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
 129 ; GCN:            v_mfma_f32_16x16x8bf16 [[RES:a\[[0-9]+:[0-9]+\]]], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
 130 ; GFX908-COUNT-4: v_accvgpr_read_b32
 131 ; GFX908:         global_store_dwordx4
 132 ; GFX90A-NOT:     v_accvgpr_read_b32
 133 ; GFX90A:         global_store_dwordx4 v{{[0-9]+}}, [[RES]],
 134 define amdgpu_kernel void @test_mfma_f32_16x16x8bf16(<4 x float> addrspace(1)* %arg) {
 135 bb:
 136   %in.1 = load <4 x float>, <4 x float> addrspace(1)* %arg
 137   %a = bitcast i32 1 to <2 x i16>
 138   %b = bitcast i32 2 to <2 x i16>
 139   %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x8bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %in.1, i32 1, i32 2, i32 3)
 140   store <4 x float> %mai.1, <4 x float> addrspace(1)* %arg
 141   ret void
 142 }