forked from OSchip/llvm-project
317 lines
15 KiB
LLVM
317 lines
15 KiB
LLVM
; RUN: llc -march=amdgcn -mcpu=gfx90a -verify-machineinstrs < %s | FileCheck -enable-var-scope --check-prefix=GCN %s
|
|
|
|
declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float, float, <32 x float>, i32, i32, i32)
|
|
declare <4 x i32> @llvm.amdgcn.mfma.i32.4x4x4i8(i32, i32, <4 x i32>, i32, i32, i32)
|
|
declare i32 @llvm.amdgcn.workitem.id.x()
|
|
|
|
; GCN-LABEL: {{^}}test_load_mfma_store16:
|
|
; GCN-COUNT-8: global_load_dwordx4 a[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}]
|
|
; GCN-NOT: v_accvgpr_write
|
|
; GCN: v_mfma_f32_32x32x1f32
|
|
; GCN-NEXT: s_nop 7
|
|
; GCN-NEXT: s_nop 7
|
|
; GCN-NEXT: s_nop 2
|
|
; GCN-NOT: v_accvgpr_read
|
|
; GCN-COUNT-8: global_store_dwordx4 v{{[0-9:]+}}, a[{{[0-9:]+}}], s[{{[0-9:]+}}]
|
|
define amdgpu_kernel void @test_load_mfma_store16(<32 x float> addrspace(1)* %arg) {
|
|
bb:
|
|
%tid = call i32 @llvm.amdgcn.workitem.id.x()
|
|
%gep = getelementptr inbounds <32 x float>, <32 x float> addrspace(1)* %arg, i32 %tid
|
|
%in.1 = load <32 x float>, <32 x float> addrspace(1)* %gep
|
|
%mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %in.1, i32 1, i32 2, i32 3)
|
|
store <32 x float> %mai.1, <32 x float> addrspace(1)* %gep
|
|
ret void
|
|
}
|
|
|
|
; GCN-LABEL: {{^}}test_load1_mfma_store1:
|
|
; GCN: global_load_dword a{{[0-9]+}}, v{{[0-9:]+}}, s[{{[0-9:]+}}]
|
|
; GCN-NOT: v_accvgpr_read
|
|
; GCN: v_mfma_f32_32x32x1f32 a{{\[}}[[N:[0-9]+]]:
|
|
; GCN-NEXT: s_nop 7
|
|
; GCN-NEXT: s_nop 7
|
|
; GCN-NEXT: s_nop 2
|
|
; GCN-NOT: v_accvgpr_read
|
|
; GCN-NEXT: global_store_dword v{{[0-9:]+}}, a[[N]], s[{{[0-9:]+}}]
|
|
define amdgpu_kernel void @test_load1_mfma_store1(float addrspace(1)* %arg) {
|
|
bb:
|
|
%tid = call i32 @llvm.amdgcn.workitem.id.x()
|
|
%gep = getelementptr inbounds float, float addrspace(1)* %arg, i32 %tid
|
|
%in.1 = load float, float addrspace(1)* %gep
|
|
%init = insertelement <32 x float> zeroinitializer, float %in.1, i32 0
|
|
%mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %init, i32 1, i32 2, i32 3)
|
|
%elt = extractelement <32 x float> %mai.1, i32 0
|
|
store float %elt, float addrspace(1)* %gep
|
|
ret void
|
|
}
|
|
|
|
; GCN-LABEL: {{^}}test_load4_mfma_store4:
|
|
; GCN: global_load_dwordx4 a[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}]
|
|
; GCN-NOT: v_accvgpr_write
|
|
; GCN: v_mfma_i32_4x4x4i8 [[A:a\[[0-9:]+\]]]
|
|
; GCN-NEXT: s_nop 4
|
|
; GCN-NOT: v_accvgpr_read
|
|
; GCN-NEXT: global_store_dwordx4 v{{[0-9:]+}}, [[A]], s[{{[0-9:]+}}]
|
|
define amdgpu_kernel void @test_load4_mfma_store4(<4 x i32> addrspace(1)* %arg) {
|
|
bb:
|
|
%tid = call i32 @llvm.amdgcn.workitem.id.x()
|
|
%gep = getelementptr inbounds <4 x i32>, <4 x i32> addrspace(1)* %arg, i32 %tid
|
|
%in.1 = load <4 x i32>, <4 x i32> addrspace(1)* %gep
|
|
%mai.1 = tail call <4 x i32> @llvm.amdgcn.mfma.i32.4x4x4i8(i32 1, i32 2, <4 x i32> %in.1, i32 0, i32 0, i32 0)
|
|
store <4 x i32> %mai.1, <4 x i32> addrspace(1)* %gep
|
|
ret void
|
|
}
|
|
|
|
; GCN-LABEL: {{^}}test_load_store:
|
|
; GCN-COUNT-8: global_load_dwordx4 v[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}]
|
|
; GCN-NOT: v_accvgpr
|
|
; GCN-COUNT-8: global_store_dwordx4 v[{{[0-9:]+}}], v[{{[0-9:]+}}]
|
|
define amdgpu_kernel void @test_load_store(<32 x float> addrspace(1)* %arg) {
|
|
bb:
|
|
%tid = call i32 @llvm.amdgcn.workitem.id.x()
|
|
%gep.1 = getelementptr inbounds <32 x float>, <32 x float> addrspace(1)* %arg, i32 %tid
|
|
%gep.2 = getelementptr inbounds <32 x float>, <32 x float> addrspace(1)* %gep.1, i32 32
|
|
%in.1 = load <32 x float>, <32 x float> addrspace(1)* %gep.1
|
|
store <32 x float> %in.1, <32 x float> addrspace(1)* %gep.2
|
|
ret void
|
|
}
|
|
|
|
; GCN-LABEL: {{^}}test_load_add_mfma_store:
|
|
; GCN-COUNT-8: global_load_dwordx4 v[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}]
|
|
; GCN-COUNT-32: v_accvgpr_write
|
|
; GCN: v_mfma_f32_32x32x1f32
|
|
; GCN-NEXT: s_nop 7
|
|
; GCN-NEXT: s_nop 7
|
|
; GCN-NEXT: s_nop 2
|
|
; GCN-NOT: v_accvgpr_read
|
|
; GCN-COUNT-8: global_store_dwordx4 v{{[0-9:]+}}, a[{{[0-9:]+}}]
|
|
define amdgpu_kernel void @test_load_add_mfma_store(<32 x float> addrspace(1)* %arg) {
|
|
bb:
|
|
%tid = call i32 @llvm.amdgcn.workitem.id.x()
|
|
%gep = getelementptr inbounds <32 x float>, <32 x float> addrspace(1)* %arg, i32 %tid
|
|
%in.1 = load <32 x float>, <32 x float> addrspace(1)* %gep
|
|
%add.1 = fadd <32 x float> %in.1, %in.1
|
|
%mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %add.1, i32 1, i32 2, i32 3)
|
|
store <32 x float> %mai.1, <32 x float> addrspace(1)* %gep
|
|
ret void
|
|
}
|
|
|
|
; GCN-LABEL: {{^}}test_load_add_store:
|
|
; GCN-COUNT-8: global_load_dwordx4 v[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}]
|
|
; GCN-NOT: v_accvgpr
|
|
; GCN-COUNT-16: v_pk_add_f32
|
|
; GCN-NOT: v_accvgpr
|
|
; GCN-COUNT-8: global_store_dwordx4 v{{[0-9:]+}}, v[{{[0-9:]+}}]
|
|
define amdgpu_kernel void @test_load_add_store(<32 x float> addrspace(1)* %arg) {
|
|
bb:
|
|
%tid = call i32 @llvm.amdgcn.workitem.id.x()
|
|
%gep = getelementptr inbounds <32 x float>, <32 x float> addrspace(1)* %arg, i32 %tid
|
|
%in.1 = load <32 x float>, <32 x float> addrspace(1)* %gep
|
|
%add.1 = fadd <32 x float> %in.1, %in.1
|
|
store <32 x float> %add.1, <32 x float> addrspace(1)* %gep
|
|
ret void
|
|
}
|
|
|
|
; GCN-LABEL: {{^}}test_load_mfma_add_store:
|
|
; GCN-COUNT-8: global_load_dwordx4 v[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}]
|
|
; GCN-COUNT-32: v_accvgpr_write
|
|
; GCN: v_mfma_f32_32x32x1f32
|
|
; GCN-COUNT-32: v_accvgpr_read
|
|
; GCN: v_pk_add_f32
|
|
; GCN-COUNT-8: global_store_dwordx4 v{{[0-9:]+}}, v[{{[0-9:]+}}]
|
|
define amdgpu_kernel void @test_load_mfma_add_store(<32 x float> addrspace(1)* %arg) {
|
|
bb:
|
|
%tid = call i32 @llvm.amdgcn.workitem.id.x()
|
|
%gep = getelementptr inbounds <32 x float>, <32 x float> addrspace(1)* %arg, i32 %tid
|
|
%in.1 = load <32 x float>, <32 x float> addrspace(1)* %gep
|
|
%mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %in.1, i32 1, i32 2, i32 3)
|
|
%add.1 = fadd <32 x float> %mai.1, %in.1
|
|
store <32 x float> %add.1, <32 x float> addrspace(1)* %gep
|
|
ret void
|
|
}
|
|
|
|
; GCN-LABEL: {{^}}test_load_add_mfma_mul_store:
|
|
; GCN-COUNT-8: global_load_dwordx4 v[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}]
|
|
; GCN: v_pk_add_f32
|
|
; GCN-COUNT-32: v_accvgpr_write
|
|
; GCN: v_mfma_f32_32x32x1f32
|
|
; GCN-COUNT-32: v_accvgpr_read
|
|
; GCN: v_pk_mul_f32
|
|
; GCN-COUNT-8: global_store_dwordx4 v{{[0-9:]+}}, v[{{[0-9:]+}}]
|
|
define amdgpu_kernel void @test_load_add_mfma_mul_store(<32 x float> addrspace(1)* %arg) {
|
|
bb:
|
|
%tid = call i32 @llvm.amdgcn.workitem.id.x()
|
|
%gep = getelementptr inbounds <32 x float>, <32 x float> addrspace(1)* %arg, i32 %tid
|
|
%in.1 = load <32 x float>, <32 x float> addrspace(1)* %gep
|
|
%add.1 = fadd <32 x float> %in.1, %in.1
|
|
%mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %add.1, i32 1, i32 2, i32 3)
|
|
%mul.1 = fmul <32 x float> %mai.1, %mai.1
|
|
store <32 x float> %mul.1, <32 x float> addrspace(1)* %gep
|
|
ret void
|
|
}
|
|
|
|
; GCN-LABEL: {{^}}test_mixeduse_load_add_mfma_mul_store:
|
|
; GCN-COUNT-8: global_load_dwordx4 v[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}]
|
|
; GCN-COUNT-32: v_accvgpr_write
|
|
; GCN: v_mfma_f32_32x32x1f32
|
|
; GCN-COUNT-32: v_accvgpr_read
|
|
; GCN: v_pk_mul_f32
|
|
; GCN-COUNT-8: global_store_dwordx4 v{{[0-9:]+}}, v[{{[0-9:]+}}]
|
|
define amdgpu_kernel void @test_mixeduse_load_add_mfma_mul_store(<32 x float> addrspace(1)* %arg) {
|
|
bb:
|
|
%tid = call i32 @llvm.amdgcn.workitem.id.x()
|
|
%gep = getelementptr inbounds <32 x float>, <32 x float> addrspace(1)* %arg, i32 %tid
|
|
%in.1 = load <32 x float>, <32 x float> addrspace(1)* %gep
|
|
%add.1 = fadd <32 x float> %in.1, %in.1
|
|
%mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %add.1, i32 1, i32 2, i32 3)
|
|
%mul.1 = fmul <32 x float> %mai.1, %in.1
|
|
store <32 x float> %mul.1, <32 x float> addrspace(1)* %gep
|
|
ret void
|
|
}
|
|
|
|
; GCN-LABEL: {{^}}test_multiuse_load_mfma_mfma_store:
|
|
; GCN-COUNT-8: global_load_dwordx4 a[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}]
|
|
; GCN-NOT: v_accvgpr_write
|
|
; GCN: v_mfma_f32_32x32x1f32
|
|
; GCN-NOT: v_accvgpr_read
|
|
; GCN-COUNT-8: global_store_dwordx4 v[{{[0-9:]+}}], a[{{[0-9:]+}}]
|
|
define amdgpu_kernel void @test_multiuse_load_mfma_mfma_store(<32 x float> addrspace(1)* %arg) {
|
|
bb:
|
|
%tid = call i32 @llvm.amdgcn.workitem.id.x()
|
|
%gep.1 = getelementptr inbounds <32 x float>, <32 x float> addrspace(1)* %arg, i32 %tid
|
|
%gep.2 = getelementptr inbounds <32 x float>, <32 x float> addrspace(1)* %gep.1, i32 32
|
|
%in.1 = load <32 x float>, <32 x float> addrspace(1)* %gep.1
|
|
%mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %in.1, i32 1, i32 2, i32 3)
|
|
%mai.2 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %in.1, i32 0, i32 0, i32 0)
|
|
store <32 x float> %mai.1, <32 x float> addrspace(1)* %gep.1
|
|
store <32 x float> %mai.2, <32 x float> addrspace(1)* %gep.2
|
|
ret void
|
|
}
|
|
|
|
; NB: for atomics both vdata and vdst shall be either VGPR or AGPR
|
|
; GCN-LABEL: {{^}}test_atomic_mfma_4xi32_atomic_store:
|
|
; GCN: global_atomic_sub [[IN:v[0-9]+]], v{{[0-9:]+}}, v{{[0-9]+}}, s[{{[0-9:]+}}] glc
|
|
; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, [[IN]]
|
|
; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
|
|
; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
|
|
; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
|
|
; GCN: v_mfma_i32_4x4x4i8 a{{\[}}[[N:[0-9]+]]:
|
|
; GCN: v_accvgpr_read_b32 [[V:v[0-9]+]], a[[N]]{{$}}
|
|
; GCN: global_atomic_add v{{[0-9]+}}, v{{[0-9:]+}}, [[V]], s[{{[0-9:]+}}] glc
|
|
; GCN: global_store_dword v{{[0-9]+}}, v{{[0-9]+}},
|
|
define amdgpu_kernel void @test_atomic_mfma_4xi32_atomic_store(i32 addrspace(1)* %arg) {
|
|
bb:
|
|
%tid = call i32 @llvm.amdgcn.workitem.id.x()
|
|
%gep = getelementptr inbounds i32, i32 addrspace(1)* %arg, i32 %tid
|
|
%in.1 = atomicrmw volatile sub i32 addrspace(1)* %gep, i32 1 seq_cst
|
|
%tmp0 = insertelement <4 x i32> undef, i32 %in.1, i32 0
|
|
%tmp1 = insertelement <4 x i32> %tmp0, i32 0, i32 1
|
|
%tmp2 = insertelement <4 x i32> %tmp1, i32 0, i32 2
|
|
%tmp3 = insertelement <4 x i32> %tmp2, i32 0, i32 3
|
|
%mai.1 = tail call <4 x i32> @llvm.amdgcn.mfma.i32.4x4x4i8(i32 1, i32 2, <4 x i32> %tmp3, i32 0, i32 0, i32 0)
|
|
%elt = extractelement <4 x i32> %mai.1, i32 0
|
|
%val = atomicrmw volatile add i32 addrspace(1)* %gep, i32 %elt seq_cst
|
|
store i32 %val, i32 addrspace(1)* %arg
|
|
ret void
|
|
}
|
|
|
|
; GCN-LABEL: {{^}}test_atomic_mfma_4xi32_atomic64_store:
|
|
; GCN: global_atomic_sub_x2 v[{{[0-9:]+}}], v{{[0-9:]+}}, v[{{[0-9:]+}}], s[{{[0-9:]+}}] glc
|
|
; GCN-COUNT-4: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
|
|
; GCN: v_mfma_i32_4x4x4i8 a{{\[}}[[N:[0-9]+]]:
|
|
; GCN: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
|
|
; GCN: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
|
|
; GCN: global_atomic_add_x2 v[{{[0-9:]+}}], v{{[0-9:]+}}, v[{{[0-9:]+}}], s[{{[0-9:]+}}] glc
|
|
define amdgpu_kernel void @test_atomic_mfma_4xi32_atomic64_store(i64 addrspace(1)* %arg) {
|
|
bb:
|
|
%tid = call i32 @llvm.amdgcn.workitem.id.x()
|
|
%gep = getelementptr inbounds i64, i64 addrspace(1)* %arg, i32 %tid
|
|
%in.1 = atomicrmw volatile sub i64 addrspace(1)* %gep, i64 1 seq_cst
|
|
%tmp0 = insertelement <2 x i64> undef, i64 %in.1, i32 0
|
|
%tmp1 = insertelement <2 x i64> %tmp0, i64 0, i32 1
|
|
%tmp2 = bitcast <2 x i64> %tmp0 to <4 x i32>
|
|
%mai.1 = tail call <4 x i32> @llvm.amdgcn.mfma.i32.4x4x4i8(i32 1, i32 2, <4 x i32> %tmp2, i32 0, i32 0, i32 0)
|
|
%elt.1 = extractelement <4 x i32> %mai.1, i32 0
|
|
%elt.2 = extractelement <4 x i32> %mai.1, i32 1
|
|
%v2.1 = insertelement <2 x i32> undef, i32 %elt.1, i32 0
|
|
%v2.2 = insertelement <2 x i32> %v2.1, i32 %elt.2, i32 1
|
|
%v2 = bitcast <2 x i32> %v2.2 to i64
|
|
%val = atomicrmw volatile add i64 addrspace(1)* %gep, i64 %v2 seq_cst
|
|
store i64 %val, i64 addrspace(1)* %arg
|
|
ret void
|
|
}
|
|
|
|
; NB: both data operands should be VGPR or AGPR
|
|
; GCN-LABEL: {{^}}test_load_mfma_ds2_store:
|
|
; GCN-DAG: ds_read_b128 [[IN:a\[[0-9:]+\]]], v{{[0-9:]+}}
|
|
; GCN-NOT: v_accvgpr_write
|
|
; GCN-DAG: v_mfma_i32_4x4x4i8 a{{\[}}[[N:[0-9]+]]:{{[0-9]+}}], v{{[0-9:]+}}, v{{[0-9:]+}}, [[IN]]
|
|
; GCN-DAG: ds_write_b32 v{{[0-9]+}}, v{{[0-9]+}}
|
|
; GCN-NOT: v_accvgpr_read
|
|
; GCN: ds_write_b32 v{{[0-9]+}}, a[[N]] offset:128
|
|
define amdgpu_kernel void @test_load_mfma_ds2_store(<4 x i32> addrspace(3)* %arg) {
|
|
bb:
|
|
%tid = call i32 @llvm.amdgcn.workitem.id.x()
|
|
%gep.1 = getelementptr inbounds <4 x i32>, <4 x i32> addrspace(3)* %arg, i32 %tid
|
|
%in.1 = load <4 x i32>, <4 x i32> addrspace(3)* %gep.1
|
|
%mai.1 = tail call <4 x i32> @llvm.amdgcn.mfma.i32.4x4x4i8(i32 1, i32 2, <4 x i32> %in.1, i32 0, i32 0, i32 0)
|
|
%elt = extractelement <4 x i32> %mai.1, i32 0
|
|
%ptr = bitcast <4 x i32> addrspace(3)* %arg to i32 addrspace(3)*
|
|
%gep.2 = getelementptr inbounds i32, i32 addrspace(3)* %ptr, i32 32
|
|
store i32 1, i32 addrspace(3)* %ptr
|
|
store i32 %elt, i32 addrspace(3)* %gep.2
|
|
ret void
|
|
}
|
|
|
|
; GCN-LABEL: {{^}}test_mfma_loop_4xi32:
|
|
; GCN: global_load_dwordx4 [[IN:a\[[0-9:]+\]]], v{{[0-9:]+}}, s[{{[0-9:]+}}]
|
|
; GCN-NOT: v_accvgpr_write
|
|
; GCN: v_mfma_i32_4x4x4i8 [[RES:a\[[0-9:]+\]]], v{{[0-9:]+}}, v{{[0-9:]+}}, [[IN]]
|
|
; GCN-NOT: v_accvgpr_read
|
|
; GCN: global_store_dwordx4 v[{{[0-9:]+}}], [[RES]],
|
|
define amdgpu_kernel void @test_mfma_loop_4xi32(<4 x i32> addrspace(1)* %arg) {
|
|
entry:
|
|
%tid = call i32 @llvm.amdgcn.workitem.id.x()
|
|
%gep = getelementptr inbounds <4 x i32>, <4 x i32> addrspace(1)* %arg, i32 %tid
|
|
%in = load <4 x i32>, <4 x i32> addrspace(1)* %gep
|
|
br label %for.cond.preheader
|
|
|
|
for.cond.preheader:
|
|
%phi = phi <4 x i32> [ %in, %entry ], [ %mai.1, %for.cond.preheader ]
|
|
%c = phi i32 [ 0, %entry ], [ %inc, %for.cond.preheader ]
|
|
%mai.1 = tail call <4 x i32> @llvm.amdgcn.mfma.i32.4x4x4i8(i32 1, i32 2, <4 x i32> %phi, i32 0, i32 0, i32 0)
|
|
%inc = add nuw nsw i32 %c, 1
|
|
%cc = icmp eq i32 %inc, 16
|
|
br i1 %cc, label %exit, label %for.cond.preheader
|
|
|
|
exit:
|
|
store <4 x i32> %mai.1, <4 x i32> addrspace(1)* %gep
|
|
ret void
|
|
}
|
|
|
|
; GCN-LABEL: {{^}}test_mfma_loop_32xfloat:
|
|
; GCN-COUNT-8: global_load_dwordx4 a[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}]
|
|
; GCN-NOT: v_accvgpr_write
|
|
; GCN: v_mfma_f32_32x32x1f32
|
|
; GCN-NOT: v_accvgpr_read
|
|
; GCN-COUNT-8: global_store_dwordx4 v[{{[0-9:]+}}], a[{{[0-9:]+}}],
|
|
; GCN: s_endpgm
|
|
define amdgpu_kernel void @test_mfma_loop_32xfloat(<32 x float> addrspace(1)* %arg) {
|
|
entry:
|
|
%tid = call i32 @llvm.amdgcn.workitem.id.x()
|
|
%gep = getelementptr inbounds <32 x float>, <32 x float> addrspace(1)* %arg, i32 %tid
|
|
%in = load <32 x float>, <32 x float> addrspace(1)* %gep
|
|
br label %for.cond.preheader
|
|
|
|
for.cond.preheader:
|
|
%phi = phi <32 x float> [ %in, %entry ], [ %mai.1, %for.cond.preheader ]
|
|
%c = phi i32 [ 0, %entry ], [ %inc, %for.cond.preheader ]
|
|
%mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %phi, i32 0, i32 0, i32 0)
|
|
%inc = add nuw nsw i32 %c, 1
|
|
%cc = icmp eq i32 %inc, 16
|
|
br i1 %cc, label %exit, label %for.cond.preheader
|
|
|
|
exit:
|
|
store <32 x float> %mai.1, <32 x float> addrspace(1)* %gep
|
|
ret void
|
|
}
|