[llvm] [AMDGPU] Fix handling of setting register classes in MFMA scheduler rewrite stage (PR #181047)
Shilei Tian via llvm-commits
llvm-commits at lists.llvm.org
Thu Feb 12 10:37:13 PST 2026
================
@@ -0,0 +1,309 @@
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -O1 -amdgpu-disable-rewrite-mfma-form-sched-stage=false -verify-machineinstrs < %s | FileCheck %s
+
+define amdgpu_kernel void @"fused_op_layout_cus_534d8e85a33330539c81dbeb009d324be60530f2_1x30x46x2016xbfloat16_1x30x46x2016xbfloat16_2016x1x1x2016xbfloat16$async_dispatch_0_matmul_like_2016x2016x1380_bf16xbf16xf32"(i32 %0, ptr addrspace(3) %1, ptr addrspace(3) %2, ptr addrspace(3) %3, ptr addrspace(3) %4, ptr addrspace(3) %5, ptr addrspace(3) %6, ptr addrspace(3) %7, ptr addrspace(3) %8, ptr addrspace(3) %9, ptr addrspace(3) %10, ptr addrspace(3) %11, ptr addrspace(3) %12, ptr addrspace(7) %invariant.gep, ptr addrspace(7) %gep, ptr addrspace(7) %gep106, ptr addrspace(7) %gep108, ptr addrspace(7) %gep112, ptr addrspace(7) %gep114, ptr addrspace(7) %gep116, ptr addrspace(7) %gep118, ptr addrspace(7) %gep120, ptr addrspace(7) %gep122, ptr addrspace(7) %gep124, ptr addrspace(3) %13, ptr addrspace(3) %14, ptr addrspace(3) %15, ptr addrspace(3) %16, ptr addrspace(3) %17, ptr addrspace(3) %18, ptr addrspace(3) %19, ptr addrspace(3) %20, ptr addrspace(3) %21, ptr addrspace(3) %22, ptr addrspace(3) %23, ptr addrspace(3) %24, ptr addrspace(3) %25, ptr addrspace(3) %26, ptr addrspace(3) %27, ptr addrspace(3) %28, ptr addrspace(3) %29, ptr addrspace(3) %30, ptr addrspace(3) %31, ptr addrspace(3) %32, ptr addrspace(3) %33, ptr addrspace(3) %34, ptr addrspace(3) %35, ptr addrspace(3) %36, ptr addrspace(3) %37, ptr addrspace(3) %38, ptr addrspace(3) %39, ptr addrspace(3) %40, ptr addrspace(3) %41, ptr addrspace(3) %42, ptr addrspace(3) %43, ptr addrspace(3) %44, ptr addrspace(3) %45, ptr addrspace(3) %46, ptr addrspace(3) %47, ptr addrspace(3) %48, ptr addrspace(3) %49, ptr addrspace(3) %50, ptr addrspace(3) %51, ptr addrspace(3) %52, ptr addrspace(3) %53, ptr addrspace(3) %54, ptr addrspace(3) %55, ptr addrspace(3) %56, ptr addrspace(3) %57, ptr addrspace(3) %58, ptr addrspace(3) %59, <4 x bfloat> %60, <4 x i16> %61, <4 x i16> %62, <4 x i16> %63, <4 x i16> %64, <4 x i16> %65, <8 x i1> %66, <4 x i16> %67, <4 x i16> %68, <4 x i16> %69, <4 x i16> %70, <8 x i1> %71) #0 {
+; This should be rejected for MFMA rewrite by the scheduler so we should see
+; no AGPR operands.
+; CHECK-NOT: v_mfma_f32_16x16x16_bf16 a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}]
+ br label %73
+73: ; preds = %73, %72
+ %74 = phi <1 x bfloat> [ zeroinitializer, %72 ], [ %339, %73 ]
+ %75 = phi [4 x [2 x [1 x <4 x float>]]] [ zeroinitializer, %72 ], [ %304, %73 ]
+ %76 = phi i32 [ 0, %72 ], [ %77, %73 ]
+ %77 = add i32 %76, 1
+ %78 = tail call i32 @llvm.amdgcn.workitem.id.x()
+ %79 = tail call i32 @llvm.umin.i32(i32 %77, i32 %78)
+ %80 = sub i32 0, %79
+ %81 = tail call i32 @llvm.smin.i32(i32 %80, i32 0)
+ %82 = load <8 x bfloat>, ptr addrspace(7) %gep, align 2
+ %83 = sub i32 0, %76
+ %84 = tail call i32 @llvm.umin.i32(i32 %83, i32 1)
+ %85 = tail call i32 @llvm.umin.i32(i32 %84, i32 %0)
+ %86 = load <8 x bfloat>, ptr addrspace(7) %gep106, align 2
+ %87 = or i32 %78, 512
+ %88 = lshr i32 %87, 1
+ %89 = tail call i32 @llvm.umin.i32(i32 %83, i32 %88)
+ %90 = sub i32 0, %89
+ %91 = tail call i32 @llvm.smin.i32(i32 %90, i32 0)
+ %92 = load <8 x bfloat>, ptr addrspace(7) %gep108, align 2
+ %93 = or i32 %78, 768
+ %94 = lshr i32 %93, 1
+ %95 = tail call i32 @llvm.umin.i32(i32 %83, i32 %94)
+ %96 = sub i32 0, %95
+ %97 = tail call i32 @llvm.smin.i32(i32 %96, i32 0)
+ %98 = load <8 x bfloat>, ptr addrspace(7) %invariant.gep, align 2
+ %99 = lshr i32 %78, 1
+ %100 = tail call i32 @llvm.umin.i32(i32 %77, i32 %99)
+ %101 = load <8 x bfloat>, ptr addrspace(7) %gep112, align 2
+ %102 = or i32 %78, 256
+ %103 = tail call i32 @llvm.umin.i32(i32 %83, i32 %102)
+ %104 = sub i32 0, %103
+ %105 = tail call i32 @llvm.smin.i32(i32 %104, i32 0)
+ %106 = load <8 x bfloat>, ptr addrspace(7) %gep114, align 2
+ %107 = tail call i32 @llvm.umin.i32(i32 %83, i32 %87)
+ %108 = load <8 x bfloat>, ptr addrspace(7) %gep116, align 2
+ %109 = tail call i32 @llvm.umin.i32(i32 %83, i32 %93)
+ %110 = load <8 x bfloat>, ptr addrspace(7) %gep118, align 2
+ %111 = or i32 %78, 64
+ %112 = tail call i32 @llvm.umin.i32(i32 %83, i32 %111)
+ %113 = load <8 x bfloat>, ptr addrspace(7) %gep120, align 2
+ %114 = or i32 %78, 80
+ %115 = tail call i32 @llvm.umin.i32(i32 %83, i32 %114)
+ %116 = load <8 x bfloat>, ptr addrspace(7) %gep122, align 2
+ %117 = or i32 %78, 96
+ %118 = tail call i32 @llvm.umin.i32(i32 %83, i32 %117)
+ %119 = load <8 x bfloat>, ptr addrspace(7) %gep124, align 2
+ %120 = or i32 %78, 1
+ %121 = tail call i32 @llvm.umin.i32(i32 %77, i32 %120)
+ %122 = load <8 x bfloat>, ptr addrspace(7) null, align 2
+ fence syncscope("workgroup") release
+ %123 = load <1 x bfloat>, ptr addrspace(3) %13, align 2
+ %124 = load <1 x bfloat>, ptr addrspace(3) %14, align 2
+ %125 = load <1 x bfloat>, ptr addrspace(3) %15, align 2
+ %126 = load <1 x bfloat>, ptr addrspace(3) %16, align 2
+ %127 = load <1 x bfloat>, ptr addrspace(3) %17, align 2
+ %128 = shufflevector <1 x bfloat> %127, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %129 = load <1 x bfloat>, ptr addrspace(3) %18, align 2
+ %130 = load <1 x bfloat>, ptr addrspace(3) %19, align 2
+ %131 = shufflevector <1 x bfloat> %130, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %132 = load <1 x bfloat>, ptr addrspace(3) %20, align 2
+ %133 = load <1 x bfloat>, ptr addrspace(3) %21, align 2
+ %134 = load <1 x bfloat>, ptr addrspace(3) %22, align 2
+ %135 = load <1 x bfloat>, ptr addrspace(3) %23, align 2
+ %136 = load <1 x bfloat>, ptr addrspace(3) %24, align 2
+ %137 = load <1 x bfloat>, ptr addrspace(3) %40, align 2
+ %138 = load <1 x bfloat>, ptr addrspace(3) %25, align 2
+ %139 = shufflevector <1 x bfloat> %138, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %140 = load <1 x bfloat>, ptr addrspace(3) %26, align 2
+ %141 = load <1 x bfloat>, ptr addrspace(3) %27, align 2
+ %142 = load <1 x bfloat>, ptr addrspace(3) %28, align 2
+ %143 = load <1 x bfloat>, ptr addrspace(3) %29, align 2
+ %144 = load <1 x bfloat>, ptr addrspace(3) %31, align 2
+ %145 = load <1 x bfloat>, ptr addrspace(3) %32, align 2
+ %146 = load <1 x bfloat>, ptr addrspace(3) %33, align 2
+ %147 = load <1 x bfloat>, ptr addrspace(3) %34, align 2
+ %148 = load <1 x bfloat>, ptr addrspace(3) %35, align 2
+ %149 = load <1 x bfloat>, ptr addrspace(3) %36, align 2
+ %150 = load <1 x bfloat>, ptr addrspace(3) %37, align 2
+ %151 = shufflevector <1 x bfloat> %150, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %152 = load <1 x bfloat>, ptr addrspace(3) %38, align 2
+ %153 = shufflevector <1 x bfloat> %152, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %154 = load <1 x bfloat>, ptr addrspace(3) %39, align 2
+ %155 = shufflevector <1 x bfloat> %154, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %156 = load <1 x bfloat>, ptr addrspace(3) %9, align 2
+ %157 = load <1 x bfloat>, ptr addrspace(3) %43, align 2
+ %158 = load <1 x bfloat>, ptr addrspace(3) %44, align 2
+ %159 = load <1 x bfloat>, ptr addrspace(3) %45, align 2
+ %160 = load <1 x bfloat>, ptr addrspace(3) %46, align 2
+ %161 = load <1 x bfloat>, ptr addrspace(3) %47, align 2
+ %162 = load <1 x bfloat>, ptr addrspace(3) %48, align 2
+ %163 = load <1 x bfloat>, ptr addrspace(3) %49, align 2
+ %164 = load <1 x bfloat>, ptr addrspace(3) %50, align 2
+ %165 = shufflevector <1 x bfloat> %164, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %166 = load <1 x bfloat>, ptr addrspace(3) %51, align 2
+ %167 = load <1 x bfloat>, ptr addrspace(3) %7, align 2
+ %168 = shufflevector <1 x bfloat> %167, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %169 = load <1 x bfloat>, ptr addrspace(3) %53, align 2
+ %170 = load <1 x bfloat>, ptr addrspace(3) %54, align 2
+ %171 = load <1 x bfloat>, ptr addrspace(3) %55, align 2
+ %172 = load <1 x bfloat>, ptr addrspace(3) %56, align 2
+ %173 = load <1 x bfloat>, ptr addrspace(3) %3, align 2
+ %174 = load <1 x bfloat>, ptr addrspace(3) %6, align 2
+ %175 = load <1 x bfloat>, ptr addrspace(3) %5, align 2
+ %176 = load <1 x bfloat>, ptr addrspace(3) %57, align 2
+ %177 = load <1 x bfloat>, ptr addrspace(3) %8, align 2
+ %178 = load <1 x bfloat>, ptr addrspace(3) %11, align 2
+ %179 = load <1 x bfloat>, ptr addrspace(3) %30, align 2
+ %180 = load <1 x bfloat>, ptr addrspace(3) %59, align 2
+ %181 = load <1 x bfloat>, ptr addrspace(3) %58, align 2
+ %182 = load <1 x bfloat>, ptr addrspace(3) %52, align 2
+ %183 = load <1 x bfloat>, ptr addrspace(3) %42, align 2
+ %184 = load <1 x bfloat>, ptr addrspace(3) %1, align 2
+ %185 = load <1 x bfloat>, ptr addrspace(3) %4, align 2
+ %186 = load <1 x bfloat>, ptr addrspace(3) %10, align 2
+ %187 = load <1 x bfloat>, ptr addrspace(3) %41, align 2
+ %188 = load <1 x bfloat>, ptr addrspace(3) %2, align 2
+ %189 = load <1 x bfloat>, ptr addrspace(3) %12, align 2
+ %190 = shufflevector <1 x bfloat> %74, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %191 = shufflevector <4 x bfloat> zeroinitializer, <4 x bfloat> %190, <4 x i32> <i32 0, i32 1, i32 2, i32 4>
+ %192 = extractvalue [4 x [2 x [1 x <4 x float>]]] %75, 3, 1, 0
+ %193 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16> zeroinitializer, <4 x i16> %67, <4 x float> %192, i32 0, i32 0, i32 0)
+ %194 = shufflevector <1 x bfloat> %123, <1 x bfloat> %125, <4 x i32> <i32 0, i32 1, i32 poison, i32 poison>
+ %195 = shufflevector <4 x bfloat> %194, <4 x bfloat> %128, <4 x i32> <i32 0, i32 1, i32 4, i32 poison>
+ %196 = shufflevector <1 x bfloat> %156, <1 x bfloat> %160, <4 x i32> <i32 0, i32 1, i32 poison, i32 poison>
+ %197 = shufflevector <4 x bfloat> %196, <4 x bfloat> %165, <4 x i32> <i32 0, i32 1, i32 4, i32 poison>
+ %198 = shufflevector <1 x bfloat> %170, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %199 = shufflevector <4 x bfloat> %195, <4 x bfloat> %131, <4 x i32> <i32 0, i32 1, i32 2, i32 4>
+ %200 = bitcast <4 x bfloat> %199 to <4 x i16>
+ %201 = shufflevector <4 x bfloat> %197, <4 x bfloat> %198, <4 x i32> <i32 0, i32 1, i32 2, i32 4>
+ %202 = bitcast <4 x bfloat> %201 to <4 x i16>
+ %203 = extractvalue [4 x [2 x [1 x <4 x float>]]] %75, 0, 0, 0
+ %204 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16> zeroinitializer, <4 x i16> %61, <4 x float> %203, i32 0, i32 0, i32 0)
+ %205 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16> %200, <4 x i16> %202, <4 x float> %204, i32 0, i32 0, i32 0)
+ %206 = shufflevector <1 x bfloat> %157, <1 x bfloat> %161, <4 x i32> <i32 0, i32 1, i32 poison, i32 poison>
+ %207 = shufflevector <1 x bfloat> %166, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %208 = shufflevector <4 x bfloat> %206, <4 x bfloat> %207, <4 x i32> <i32 0, i32 1, i32 4, i32 poison>
+ %209 = shufflevector <1 x bfloat> %171, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %210 = shufflevector <4 x bfloat> %208, <4 x bfloat> %209, <4 x i32> <i32 0, i32 1, i32 2, i32 4>
+ %211 = shufflevector <1 x bfloat> %158, <1 x bfloat> %162, <4 x i32> <i32 0, i32 1, i32 poison, i32 poison>
+ %212 = shufflevector <4 x bfloat> %211, <4 x bfloat> %168, <4 x i32> <i32 0, i32 1, i32 4, i32 poison>
+ %213 = shufflevector <1 x bfloat> %172, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %214 = shufflevector <1 x bfloat> %159, <1 x bfloat> %163, <4 x i32> <i32 0, i32 1, i32 poison, i32 poison>
+ %215 = shufflevector <1 x bfloat> %169, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %216 = shufflevector <4 x bfloat> %214, <4 x bfloat> %215, <4 x i32> <i32 0, i32 1, i32 4, i32 poison>
+ %217 = shufflevector <1 x bfloat> %173, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %218 = shufflevector <4 x bfloat> %216, <4 x bfloat> %217, <4 x i32> <i32 0, i32 1, i32 2, i32 4>
+ %219 = bitcast <4 x bfloat> %218 to <4 x i16>
+ %220 = shufflevector <1 x bfloat> %124, <1 x bfloat> %126, <4 x i32> <i32 0, i32 1, i32 poison, i32 poison>
+ %221 = shufflevector <1 x bfloat> %129, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %222 = bitcast <4 x bfloat> %210 to <4 x i16>
+ %223 = extractvalue [4 x [2 x [1 x <4 x float>]]] %75, 1, 1, 0
+ %224 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16> zeroinitializer, <4 x i16> %222, <4 x float> %223, i32 0, i32 0, i32 0)
+ %225 = shufflevector <4 x bfloat> %220, <4 x bfloat> %221, <4 x i32> <i32 0, i32 1, i32 4, i32 poison>
+ %226 = shufflevector <1 x bfloat> %132, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %227 = shufflevector <4 x bfloat> %225, <4 x bfloat> %226, <4 x i32> <i32 0, i32 1, i32 2, i32 4>
+ %228 = bitcast <4 x bfloat> %227 to <4 x i16>
+ %229 = shufflevector <4 x bfloat> %212, <4 x bfloat> %213, <4 x i32> <i32 0, i32 1, i32 2, i32 4>
+ %230 = bitcast <4 x bfloat> %229 to <4 x i16>
+ %231 = bitcast <4 x bfloat> %191 to <4 x i16>
+ %232 = extractvalue [4 x [2 x [1 x <4 x float>]]] %75, 2, 1, 0
+ %233 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16> %62, <4 x i16> %231, <4 x float> %232, i32 0, i32 0, i32 0)
+ %234 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16> %228, <4 x i16> %230, <4 x float> %233, i32 0, i32 0, i32 0)
+ %235 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16> zeroinitializer, <4 x i16> %219, <4 x float> %193, i32 0, i32 0, i32 0)
+ %236 = shufflevector <1 x bfloat> %133, <1 x bfloat> %135, <4 x i32> <i32 0, i32 1, i32 poison, i32 poison>
+ %237 = shufflevector <1 x bfloat> %137, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %238 = shufflevector <4 x bfloat> %236, <4 x bfloat> %237, <4 x i32> <i32 0, i32 1, i32 4, i32 poison>
+ %239 = shufflevector <4 x bfloat> %238, <4 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 1, i32 2, i32 4>
+ %240 = shufflevector <1 x bfloat> %174, <1 x bfloat> %178, <4 x i32> <i32 0, i32 1, i32 poison, i32 poison>
+ %241 = shufflevector <1 x bfloat> %182, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %242 = shufflevector <4 x bfloat> %240, <4 x bfloat> %241, <4 x i32> <i32 0, i32 1, i32 4, i32 poison>
+ %243 = shufflevector <1 x bfloat> %185, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %244 = shufflevector <4 x bfloat> %242, <4 x bfloat> %243, <4 x i32> <i32 0, i32 1, i32 2, i32 4>
+ %245 = bitcast <4 x bfloat> %239 to <4 x i16>
+ %246 = bitcast <4 x bfloat> %244 to <4 x i16>
+ %247 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16> %245, <4 x i16> %246, <4 x float> %205, i32 0, i32 0, i32 0)
+ %248 = shufflevector <1 x bfloat> %175, <1 x bfloat> %179, <4 x i32> <i32 0, i32 1, i32 poison, i32 poison>
+ %249 = shufflevector <1 x bfloat> %183, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %250 = shufflevector <4 x bfloat> %248, <4 x bfloat> %249, <4 x i32> <i32 0, i32 1, i32 4, i32 poison>
+ %251 = shufflevector <1 x bfloat> %186, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %252 = shufflevector <4 x bfloat> %250, <4 x bfloat> %251, <4 x i32> <i32 0, i32 1, i32 2, i32 4>
+ %253 = bitcast <4 x bfloat> %252 to <4 x i16>
+ %254 = shufflevector <1 x bfloat> %176, <1 x bfloat> %180, <4 x i32> <i32 0, i32 1, i32 poison, i32 poison>
+ %255 = shufflevector <4 x bfloat> %254, <4 x bfloat> %60, <4 x i32> <i32 0, i32 1, i32 4, i32 poison>
+ %256 = shufflevector <1 x bfloat> %187, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %257 = shufflevector <4 x bfloat> %255, <4 x bfloat> %256, <4 x i32> <i32 0, i32 1, i32 2, i32 4>
+ %258 = bitcast <4 x bfloat> %257 to <4 x i16>
+ %259 = shufflevector <1 x bfloat> %177, <1 x bfloat> %181, <4 x i32> <i32 0, i32 1, i32 poison, i32 poison>
+ %260 = shufflevector <1 x bfloat> %184, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %261 = shufflevector <4 x bfloat> %259, <4 x bfloat> %260, <4 x i32> <i32 0, i32 1, i32 4, i32 poison>
+ %262 = shufflevector <1 x bfloat> %188, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %263 = shufflevector <4 x bfloat> %261, <4 x bfloat> %262, <4 x i32> <i32 0, i32 1, i32 2, i32 4>
+ %264 = bitcast <4 x bfloat> %263 to <4 x i16>
+ %265 = shufflevector <1 x bfloat> %134, <1 x bfloat> %136, <4 x i32> <i32 0, i32 1, i32 poison, i32 poison>
+ %266 = shufflevector <4 x bfloat> %265, <4 x bfloat> %139, <4 x i32> <i32 0, i32 1, i32 2, i32 4>
+ %267 = bitcast <4 x bfloat> %266 to <4 x i16>
+ %268 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16> %267, <4 x i16> %253, <4 x float> %224, i32 0, i32 0, i32 0)
+ %269 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16> zeroinitializer, <4 x i16> %258, <4 x float> %234, i32 0, i32 0, i32 0)
+ %270 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16> zeroinitializer, <4 x i16> %264, <4 x float> %235, i32 0, i32 0, i32 0)
+ %271 = shufflevector <1 x bfloat> %140, <1 x bfloat> %141, <4 x i32> <i32 0, i32 1, i32 poison, i32 poison>
+ %272 = shufflevector <1 x bfloat> %142, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %273 = shufflevector <4 x bfloat> %271, <4 x bfloat> %272, <4 x i32> <i32 0, i32 1, i32 4, i32 poison>
+ %274 = shufflevector <1 x bfloat> %143, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %275 = shufflevector <4 x bfloat> %273, <4 x bfloat> %274, <4 x i32> <i32 0, i32 1, i32 2, i32 4>
+ %276 = shufflevector <1 x bfloat> splat (bfloat 0xR3F80), <1 x bfloat> %189, <4 x i32> <i32 0, i32 1, i32 poison, i32 poison>
+ %277 = bitcast <4 x bfloat> %275 to <4 x i16>
+ %278 = bitcast <4 x bfloat> %276 to <4 x i16>
+ %279 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16> %277, <4 x i16> %278, <4 x float> %247, i32 0, i32 0, i32 0)
+ %280 = shufflevector <1 x bfloat> %144, <1 x bfloat> %145, <4 x i32> <i32 0, i32 1, i32 poison, i32 poison>
+ %281 = shufflevector <1 x bfloat> %146, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %282 = shufflevector <4 x bfloat> %280, <4 x bfloat> %281, <4 x i32> <i32 0, i32 1, i32 4, i32 poison>
+ %283 = shufflevector <1 x bfloat> %147, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %284 = shufflevector <4 x bfloat> %282, <4 x bfloat> %283, <4 x i32> <i32 0, i32 1, i32 2, i32 4>
+ %285 = bitcast <4 x bfloat> %284 to <4 x i16>
+ %286 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16> %285, <4 x i16> %63, <4 x float> %279, i32 0, i32 0, i32 0)
+ %287 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16> zeroinitializer, <4 x i16> %69, <4 x float> %268, i32 0, i32 0, i32 0)
+ %288 = shufflevector <1 x bfloat> %148, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 1, i32 poison, i32 poison>
+ %289 = shufflevector <1 x bfloat> %149, <1 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
+ %290 = shufflevector <4 x bfloat> %288, <4 x bfloat> %289, <4 x i32> <i32 0, i32 1, i32 4, i32 poison>
+ %291 = shufflevector <4 x bfloat> %290, <4 x bfloat> %153, <4 x i32> <i32 0, i32 1, i32 2, i32 4>
+ %292 = bitcast <4 x bfloat> %291 to <4 x i16>
+ %293 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16> %292, <4 x i16> splat (i16 1), <4 x float> %286, i32 0, i32 0, i32 0)
+ %294 = shufflevector <4 x bfloat> %151, <4 x bfloat> %155, <4 x i32> <i32 0, i32 1, i32 2, i32 4>
+ %295 = bitcast <4 x bfloat> %294 to <4 x i16>
+ %296 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16> %70, <4 x i16> zeroinitializer, <4 x float> %293, i32 0, i32 0, i32 0)
+ %297 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16> %65, <4 x i16> zeroinitializer, <4 x float> %296, i32 0, i32 0, i32 0)
+ %298 = insertvalue [4 x [2 x [1 x <4 x float>]]] zeroinitializer, <4 x float> %297, 0, 0, 0
+ %299 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16> %295, <4 x i16> zeroinitializer, <4 x float> %287, i32 0, i32 0, i32 0)
+ %300 = insertvalue [4 x [2 x [1 x <4 x float>]]] %298, <4 x float> %299, 1, 1, 0
+ %301 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16> zeroinitializer, <4 x i16> %64, <4 x float> %269, i32 0, i32 0, i32 0)
+ %302 = insertvalue [4 x [2 x [1 x <4 x float>]]] %300, <4 x float> %301, 2, 1, 0
+ %303 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16> %68, <4 x i16> zeroinitializer, <4 x float> %270, i32 0, i32 0, i32 0)
+ %304 = insertvalue [4 x [2 x [1 x <4 x float>]]] %302, <4 x float> %303, 3, 1, 0
+ %305 = trunc i32 %81 to i1
+ %306 = select i1 %305, <8 x bfloat> %82, <8 x bfloat> zeroinitializer
+ tail call void @llvm.amdgcn.sched.barrier(i32 0)
+ %307 = and i32 %78, 56
+ %308 = getelementptr bfloat, ptr addrspace(3) null, i32 %307
+ store <8 x bfloat> %306, ptr addrspace(3) %308, align 8
+ %309 = trunc i32 %85 to i1
+ %310 = select i1 %309, <8 x bfloat> %86, <8 x bfloat> zeroinitializer
+ %311 = getelementptr bfloat, ptr addrspace(3) %1, i32 %307
+ store <8 x bfloat> %310, ptr addrspace(3) %311, align 8
+ %312 = trunc i32 %91 to i1
+ %313 = select i1 %312, <8 x bfloat> %92, <8 x bfloat> zeroinitializer
+ store <8 x bfloat> %313, ptr addrspace(3) null, align 8
+ %314 = trunc i32 %97 to i1
+ %315 = select i1 %314, <8 x bfloat> %98, <8 x bfloat> zeroinitializer
+ store <8 x bfloat> %315, ptr addrspace(3) %1, align 8
+ %316 = trunc i32 %100 to i1
+ %317 = select i1 %316, <8 x bfloat> %101, <8 x bfloat> zeroinitializer
+ %318 = and i32 %78, 1
+ %319 = getelementptr bfloat, ptr addrspace(3) %3, i32 %318
+ store <8 x bfloat> %317, ptr addrspace(3) %319, align 8
+ %320 = trunc i32 %105 to i1
+ %321 = select i1 %320, <8 x bfloat> %106, <8 x bfloat> zeroinitializer
+ %322 = getelementptr bfloat, ptr addrspace(3) %4, i32 %318
+ store <8 x bfloat> %321, ptr addrspace(3) %322, align 8
+ %323 = trunc i32 %107 to i1
+ %324 = select i1 %323, <8 x bfloat> %108, <8 x bfloat> zeroinitializer
+ store <8 x bfloat> %324, ptr addrspace(3) null, align 8
+ %325 = trunc i32 %109 to i1
+ %326 = select i1 %325, <8 x bfloat> %110, <8 x bfloat> zeroinitializer
+ %327 = getelementptr bfloat, ptr addrspace(3) null, i32 %318
+ store <8 x bfloat> %326, ptr addrspace(3) %327, align 8
+ %328 = trunc i32 %112 to i1
+ %329 = select i1 %328, <8 x bfloat> %113, <8 x bfloat> zeroinitializer
+ %330 = getelementptr bfloat, ptr addrspace(3) %5, i32 %318
+ store <8 x bfloat> %329, ptr addrspace(3) %330, align 8
+ %331 = trunc i32 %115 to i1
+ %332 = select i1 %331, <8 x bfloat> %116, <8 x bfloat> zeroinitializer
+ store <8 x bfloat> %332, ptr addrspace(3) %1, align 8
+ %333 = trunc i32 %118 to i1
+ %334 = select i1 %333, <8 x bfloat> %119, <8 x bfloat> zeroinitializer
+ store <8 x bfloat> %334, ptr addrspace(3) null, align 8
+ %335 = trunc i32 %121 to i1
+ %336 = select i1 %335, <8 x bfloat> splat (bfloat 0xR3F80), <8 x bfloat> zeroinitializer
+ store <8 x bfloat> %122, ptr addrspace(3) %1, align 8
+ %337 = bitcast <8 x bfloat> %336 to i128
+ %338 = trunc i128 %337 to i16
+ %339 = bitcast i16 %338 to <1 x bfloat>
+ %340 = icmp ult i32 %76, 1
+ br i1 %340, label %73, label %341
+
+341: ; preds = %341, %73
+ %342 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16> zeroinitializer, <4 x i16> zeroinitializer, <4 x float> splat (float 0x7FF8000000000000), i32 0, i32 0, i32 0)
+ %343 = shufflevector <4 x float> %342, <4 x float> zeroinitializer, <32 x i32> <i32 0, i32 1, i32 2, i32 3, i32 poison, i32 poison, i32 poison, i32 poison, i32 poison, i32 poison, i32 poison, i32 poison, i32 poison, i32 poison, i32 poison, i32 poison, i32 poison, i32 poison, i32 poison, i32 poison, i32 poison, i32 poison, i32 poison, i32 poison, i32 poison, i32 poison, i32 poison, i32 poison, i32 poison, i32 poison, i32 poison, i32 poison>
+ %344 = extractelement <32 x float> %343, i32 %0
+ %345 = insertelement <4 x float> zeroinitializer, float %344, i64 0
+ %346 = shufflevector <4 x float> %345, <4 x float> zeroinitializer, <8 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7>
+ %347 = fcmp uno <8 x float> %346, zeroinitializer
+ %348 = select <8 x i1> %347, <8 x bfloat> splat (bfloat 0xR3F80), <8 x bfloat> zeroinitializer
+ tail call void @llvm.masked.scatter.v8bf16.v8p7(<8 x bfloat> %348, <8 x ptr addrspace(7)> zeroinitializer, <8 x i1> %66)
+ br label %341
+}
+
+attributes #0 = { "amdgpu-flat-work-group-size"="256,256" }
+attributes #1 = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
----------------
shiltian wrote:
These arguments are not needed except `#0`.
https://github.com/llvm/llvm-project/pull/181047
More information about the llvm-commits
mailing list