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

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