llvm/test/CodeGen/AMDGPU/mfma-vgpr-cd-select.ll

   1 ; RUN: llc -mtriple=amdgcn -mcpu=gfx90a -verify-machineinstrs < %s | FileCheck --enable-var-scope --check-prefixes=GCN %s
   2 ; RUN: llc -global-isel -mtriple=amdgcn -mcpu=gfx90a -verify-machineinstrs < %s | FileCheck --enable-var-scope --check-prefixes=GCN %s
   3 ; RUN: llc -mtriple=amdgcn -mcpu=gfx940 -verify-machineinstrs < %s | FileCheck --enable-var-scope --check-prefixes=GCN %s
   4 ; RUN: llc -global-isel -mtriple=amdgcn -mcpu=gfx940 -verify-machineinstrs < %s | FileCheck --enable-var-scope --check-prefixes=GCN %s
   5
   6 declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float, float, <32 x float>, i32, i32, i32)
   7 declare <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float, float, <16 x float>, i32, i32, i32)
   8 declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float, float, <4 x float>, i32, i32, i32)
   9 declare <16 x float> @llvm.amdgcn.mfma.f32.32x32x2f32(float, float, <16 x float>, i32, i32, i32)
  10 declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x4f32(float, float, <4 x float>, i32, i32, i32)
  11 declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x4f16(<4 x half>, <4 x half>, <32 x float>, i32, i32, i32)
  12 declare <16 x float> @llvm.amdgcn.mfma.f32.16x16x4f16(<4 x half>, <4 x half>, <16 x float>, i32, i32, i32)
  13 declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x4f16(<4 x half>, <4 x half>, <4 x float>, i32, i32, i32)
  14 declare <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half>, <4 x half>, <16 x float>, i32, i32, i32)
  15 declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half>, <4 x half>, <4 x float>, i32, i32, i32)
  16 declare <32 x i32> @llvm.amdgcn.mfma.i32.32x32x4i8(i32, i32, <32 x i32>, i32, i32, i32)
  17 declare <16 x i32> @llvm.amdgcn.mfma.i32.16x16x4i8(i32, i32, <16 x i32>, i32, i32, i32)
  18 declare <4 x i32> @llvm.amdgcn.mfma.i32.4x4x4i8(i32, i32, <4 x i32>, i32, i32, i32)
  19
  20 ; GCN-LABEL: {{^}}test_mfma_f32_32x32x1f32:
  21 ; GCN: v_mfma_f32_32x32x1{{.*}} v[{{[0-9]+:[0-9]+}}], v{{[0-9]+}}, v{{[0-9:]+}}, v[{{[0-9]+:[0-9]+}}]
  22 define amdgpu_kernel void @test_mfma_f32_32x32x1f32(ptr addrspace(1) %arg) #0 {
  23 bb:
  24   %in.1 = load <32 x float>, ptr addrspace(1) %arg
  25   %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 1.0, <32 x float> %in.1, i32 0, i32 0, i32 0)
  26   store <32 x float> %mai.1, ptr addrspace(1) %arg
  27   ret void
  28 }
  29
  30 ; GCN-LABEL: {{^}}test_mfma_f32_16x16x1f32:
  31 ; GCN: v_mfma_f32_16x16x1{{.*}} v[{{[0-9]+:[0-9]+}}], v{{[0-9]+}}, v{{[0-9:]+}}, v[{{[0-9]+:[0-9]+}}]
  32 define amdgpu_kernel void @test_mfma_f32_16x16x1f32(ptr addrspace(1) %arg) #0 {
  33 bb:
  34   %in.1 = load <16 x float>, ptr addrspace(1) %arg
  35   %mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float 1.0, float 1.0, <16 x float> %in.1, i32 0, i32 0, i32 0)
  36   store <16 x float> %mai.1, ptr addrspace(1) %arg
  37   ret void
  38 }
  39
  40 ; GCN-LABEL: {{^}}test_mfma_f32_4x4x1f32:
  41 ; GCN: v_mfma_f32_4x4x1{{.*}} v[{{[0-9]+:[0-9]+}}], v{{[0-9]+}}, v{{[0-9]+}}, v[{{[0-9]+:[0-9]+}}]
  42 define amdgpu_kernel void @test_mfma_f32_4x4x1f32(ptr addrspace(1) %arg) #0 {
  43 bb:
  44   %in.1 = load <4 x float>, ptr addrspace(1) %arg
  45   %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 1.0, <4 x float> %in.1, i32 0, i32 0, i32 0)
  46   store <4 x float> %mai.1, ptr addrspace(1) %arg
  47   ret void
  48 }
  49
  50 ; GCN-LABEL: {{^}}test_mfma_f32_32x32x2f32:
  51 ; GCN: v_mfma_f32_32x32x2{{.*}} v[{{[0-9]+:[0-9]+}}], v{{[0-9]+}}, v{{[0-9]+}}, v[{{[0-9]+:[0-9]+}}]
  52 define amdgpu_kernel void @test_mfma_f32_32x32x2f32(ptr addrspace(1) %arg) #0 {
  53 bb:
  54   %in.1 = load <16 x float>, ptr addrspace(1) %arg
  55   %mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x2f32(float 1.0, float 1.0, <16 x float> %in.1, i32 0, i32 0, i32 0)
  56   store <16 x float> %mai.1, ptr addrspace(1) %arg
  57   ret void
  58 }
  59
  60 ; GCN-LABEL: {{^}}test_mfma_f32_16x16x4f32:
  61 ; GCN: v_mfma_f32_16x16x4{{.*}} v[{{[0-9]+:[0-9]+}}], v{{[0-9]+}}, v{{[0-9]+}}, v[{{[0-9]+:[0-9]+}}]
  62 define amdgpu_kernel void @test_mfma_f32_16x16x4f32(ptr addrspace(1) %arg) #0 {
  63 bb:
  64   %in.1 = load <4 x float>, ptr addrspace(1) %arg
  65   %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x4f32(float 1.0, float 1.0, <4 x float> %in.1, i32 0, i32 0, i32 0)
  66   store <4 x float> %mai.1, ptr addrspace(1) %arg
  67   ret void
  68 }
  69
  70 ; GCN-LABEL: {{^}}test_mfma_f32_32x32x4f16:
  71 ; GCN: v_mfma_f32_32x32x4{{.*}} v[{{[0-9]+:[0-9]+}}], v[{{[0-9]+:[0-9]+}}], v[{{[0-9]+:[0-9]+}}], v[{{[0-9]+:[0-9]+}}]
  72 define amdgpu_kernel void @test_mfma_f32_32x32x4f16(ptr addrspace(1) %arg) #0 {
  73 bb:
  74   %in.1 = load <32 x float>, ptr addrspace(1) %arg
  75   %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x4f16(<4 x half> undef, <4 x half> undef, <32 x float> %in.1, i32 0, i32 0, i32 0)
  76   store <32 x float> %mai.1, ptr addrspace(1) %arg
  77   ret void
  78 }
  79
  80 ; GCN-LABEL: {{^}}test_mfma_f32_16x16x4f16:
  81 ; GCN: v_mfma_f32_16x16x4{{.*}} v[{{[0-9]+:[0-9]+}}], v[{{[0-9]+:[0-9]+}}], v[{{[0-9]+:[0-9]+}}], v[{{[0-9]+:[0-9]+}}]
  82 define amdgpu_kernel void @test_mfma_f32_16x16x4f16(ptr addrspace(1) %arg) #0 {
  83 bb:
  84   %in.1 = load <16 x float>, ptr addrspace(1) %arg
  85   %mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.16x16x4f16(<4 x half> undef, <4 x half> undef, <16 x float> %in.1, i32 0, i32 0, i32 0)
  86   store <16 x float> %mai.1, ptr addrspace(1) %arg
  87   ret void
  88 }
  89
  90 ; GCN-LABEL: {{^}}test_mfma_f32_4x4x4f16:
  91 ; GCN: v_mfma_f32_4x4x4{{.*}} v[{{[0-9]+:[0-9]+}}], v[{{[0-9]+:[0-9]+}}], v[{{[0-9]+:[0-9]+}}], v[{{[0-9]+:[0-9]+}}]
  92 define amdgpu_kernel void @test_mfma_f32_4x4x4f16(ptr addrspace(1) %arg) #0 {
  93 bb:
  94   %in.1 = load <4 x float>, ptr addrspace(1) %arg
  95   %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x4f16(<4 x half> undef, <4 x half> undef, <4 x float> %in.1, i32 0, i32 0, i32 0)
  96   store <4 x float> %mai.1, ptr addrspace(1) %arg
  97   ret void
  98 }
  99
 100 ; GCN-LABEL: {{^}}test_mfma_f32_32x32x8f16:
 101 ; GCN: v_mfma_f32_32x32x8{{.*}} v[{{[0-9]+:[0-9]+}}], v[{{[0-9]+:[0-9]+}}], v[{{[0-9]+:[0-9]+}}], v[{{[0-9]+:[0-9]+}}]
 102 define amdgpu_kernel void @test_mfma_f32_32x32x8f16(ptr addrspace(1) %arg) #0 {
 103 bb:
 104   %in.1 = load <16 x float>, ptr addrspace(1) %arg
 105   %mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> undef, <4 x half> undef, <16 x float> %in.1, i32 0, i32 0, i32 0)
 106   store <16 x float> %mai.1, ptr addrspace(1) %arg
 107   ret void
 108 }
 109
 110 ; GCN-LABEL: {{^}}test_mfma_f32_16x16x16f16:
 111 ; GCN: v_mfma_f32_16x16x16{{.*}} v[{{[0-9]+:[0-9]+}}], v[{{[0-9]+:[0-9]+}}], v[{{[0-9]+:[0-9]+}}], v[{{[0-9]+:[0-9]+}}]
 112 define amdgpu_kernel void @test_mfma_f32_16x16x16f16(ptr addrspace(1) %arg) #0 {
 113 bb:
 114   %in.1 = load <4 x float>, ptr addrspace(1) %arg
 115   %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> undef, <4 x half> undef, <4 x float> %in.1, i32 0, i32 0, i32 0)
 116   store <4 x float> %mai.1, ptr addrspace(1) %arg
 117   ret void
 118 }
 119
 120 ; GCN-LABEL: {{^}}test_mfma_i32_32x32x4i8:
 121 ; GCN: v_mfma_i32_32x32x4{{.*}} v[{{[0-9]+:[0-9]+}}], v{{[0-9]+}}, v{{[0-9]+}}, v[{{[0-9]+:[0-9]+}}]
 122 define amdgpu_kernel void @test_mfma_i32_32x32x4i8(ptr addrspace(1) %arg) #0 {
 123 bb:
 124   %in.1 = load <32 x i32>, ptr addrspace(1) %arg
 125   %mai.1 = tail call <32 x i32> @llvm.amdgcn.mfma.i32.32x32x4i8(i32 1, i32 1, <32 x i32> %in.1, i32 0, i32 0, i32 0)
 126   store <32 x i32> %mai.1, ptr addrspace(1) %arg
 127   ret void
 128 }
 129
 130 ; GCN-LABEL: {{^}}test_mfma_i32_16x16x4i8:
 131 ; GCN: v_mfma_i32_16x16x4{{.*}} v[{{[0-9]+:[0-9]+}}], v{{[0-9]+}}, v{{[0-9]+}}, v[{{[0-9]+:[0-9]+}}]
 132 define amdgpu_kernel void @test_mfma_i32_16x16x4i8(ptr addrspace(1) %arg) #0 {
 133 bb:
 134   %in.1 = load <16 x i32>, ptr addrspace(1) %arg
 135   %mai.1 = tail call <16 x i32> @llvm.amdgcn.mfma.i32.16x16x4i8(i32 1, i32 1, <16 x i32> %in.1, i32 0, i32 0, i32 0)
 136   store <16 x i32> %mai.1, ptr addrspace(1) %arg
 137   ret void
 138 }
 139
 140 ; GCN-LABEL: {{^}}test_mfma_i32_4x4x4i8:
 141 ; GCN: v_mfma_i32_4x4x4{{.*}} v[{{[0-9]+:[0-9]+}}], v{{[0-9]+}}, v{{[0-9]+}}, v[{{[0-9]+:[0-9]+}}]
 142 define amdgpu_kernel void @test_mfma_i32_4x4x4i8(ptr addrspace(1) %arg) #0 {
 143 bb:
 144   %in.1 = load <4 x i32>, ptr addrspace(1) %arg
 145   %mai.1 = tail call <4 x i32> @llvm.amdgcn.mfma.i32.4x4x4i8(i32 1, i32 1, <4 x i32> %in.1, i32 0, i32 0, i32 0)
 146   store <4 x i32> %mai.1, ptr addrspace(1) %arg
 147   ret void
 148 }
 149
 150 attributes #0 = { "amdgpu-no-agpr" }