191 lines · plain
1; RUN: llc -mtriple=amdgcn -mcpu=gfx908 < %s | FileCheck -check-prefixes=GCN,GFX908 %s2 3; GCN-LABEL: {{^}}accvgpr_write_read:4; GFX908: v_accvgpr_write [[AREG:a[0-9]+]], 15; GFX908: v_accvgpr_read [[VREG:v[0-9]+]], [[AREG]]6; GFX908: global_store_dword v{{[0-9]+}}, [[VREG]], s{{\[[0-9]+:[0-9]+\]}}7define amdgpu_kernel void @accvgpr_write_read(ptr addrspace(1) %arg) {8bb:9 %in.1 = load float, ptr addrspace(1) %arg10 %init = tail call float asm "v_accvgpr_write $0, 1", "=a"()11 %read = tail call float asm "v_accvgpr_read $0, $1", "=v,a"(float %init)12 store float %read, ptr addrspace(1) %arg13 ret void14}15 16; GCN-LABEL: {{^}}v_mfma_f32_4x4x1f32_avva17; GFX908: v_accvgpr_write_b3218; GFX908: v_accvgpr_write_b3219; GFX908: v_accvgpr_write_b3220; GFX908: v_accvgpr_write_b3221; GFX908: v_mfma_f32_4x4x1f32 a[{{[0-9:]+}}], v{{[0-9]+}}, v{{[0-9]+}}, a[{{[0-9:]+}}]22; GFX908: v_accvgpr_read_b3223; GFX908: v_accvgpr_read_b3224; GFX908: v_accvgpr_read_b3225; GFX908: v_accvgpr_read_b3226define amdgpu_kernel void @v_mfma_f32_4x4x1f32_avva(ptr addrspace(1) %arg) {27bb:28 %in.1 = load <4 x float>, ptr addrspace(1) %arg29 %mai.1 = tail call <4 x float> asm "v_mfma_f32_4x4x1f32 $0, $1, $2, $3", "=a,v,v,a"(float 1.0, float 2.0, <4 x float> %in.1)30 store <4 x float> %mai.1, ptr addrspace(1) %arg31 ret void32}33 34; GCN-LABEL: {{^}}v_mfma_f32_4x4x1f32_aaaa35; GFX908: v_accvgpr_write_b3236; GFX908: v_accvgpr_write_b3237; GFX908: v_accvgpr_write_b3238; GFX908: v_accvgpr_write_b3239; GFX908: v_mfma_f32_4x4x1f32 a[{{[0-9:]+}}], a{{[0-9]+}}, a{{[0-9]+}}, a[{{[0-9:]+}}]40; GFX908: v_accvgpr_read_b3241; GFX908: v_accvgpr_read_b3242; GFX908: v_accvgpr_read_b3243; GFX908: v_accvgpr_read_b3244define amdgpu_kernel void @v_mfma_f32_4x4x1f32_aaaa(ptr addrspace(1) %arg) {45bb:46 %in.1 = load <4 x float>, ptr addrspace(1) %arg47 %mai.1 = tail call <4 x float> asm "v_mfma_f32_4x4x1f32 $0, $1, $2, $3", "=a,a,a,a"(float 1.0, float 2.0, <4 x float> %in.1)48 store <4 x float> %mai.1, ptr addrspace(1) %arg49 ret void50}51 52; GCN-LABEL: {{^}}v_mfma_f32_4x4x4f16_aaaa53; GFX908: v_accvgpr_write_b3254; GFX908: v_accvgpr_write_b3255; GFX908: v_accvgpr_write_b3256; GFX908: v_accvgpr_write_b3257; GFX908: v_accvgpr_write_b3258; GFX908: v_accvgpr_write_b3259; GFX908: v_accvgpr_write_b3260; GFX908: v_accvgpr_write_b3261; GFX908: v_mfma_f32_4x4x4f16 a[{{[0-9:]+}}], a[{{[0-9]+:[0-9]+}}], a[{{[0-9]+:[0-9]+}}], a[{{[0-9:]+}}]62; GFX908: v_accvgpr_read_b3263; GFX908: v_accvgpr_read_b3264; GFX908: v_accvgpr_read_b3265; GFX908: v_accvgpr_read_b3266define amdgpu_kernel void @v_mfma_f32_4x4x4f16_aaaa(ptr addrspace(1) %arg) {67bb:68 %in.1 = load <4 x float>, ptr addrspace(1) %arg69 %mai.1 = tail call <4 x float> asm "v_mfma_f32_4x4x4f16 $0, $1, $2, $3", "=a,a,a,a"(<4 x half> <half 0xH3800, half 0xH3800, half 0xH3800, half 0xH3800>, <4 x half> <half 0xH03FF, half 0xH03FF, half 0xH03FF, half 0xH03FF>, <4 x float> %in.1)70 store <4 x float> %mai.1, ptr addrspace(1) %arg71 ret void72}73 74; GCN-LABEL: {{^}}v_mfma_f32_16x16x1f32_avaa75; GFX908: v_accvgpr_write_b3276; GFX908: v_accvgpr_write_b3277; GFX908: v_accvgpr_write_b3278; GFX908: v_accvgpr_write_b3279; GFX908: v_accvgpr_write_b3280; GFX908: v_accvgpr_write_b3281; GFX908: v_accvgpr_write_b3282; GFX908: v_accvgpr_write_b3283; GFX908: v_accvgpr_write_b3284; GFX908: v_accvgpr_write_b3285; GFX908: v_accvgpr_write_b3286; GFX908: v_accvgpr_write_b3287; GFX908: v_accvgpr_write_b3288; GFX908: v_accvgpr_write_b3289; GFX908: v_accvgpr_write_b3290; GFX908: v_accvgpr_write_b3291; GFX908: v_accvgpr_write_b3292; GFX908: v_mfma_f32_16x16x1f32 a[{{[0-9:]+}}], v{{[0-9]+}}, a{{[0-9]+}}, a[{{[0-9:]+}}]93; GFX908: v_accvgpr_read_b3294; GFX908: v_accvgpr_read_b3295; GFX908: v_accvgpr_read_b3296; GFX908: v_accvgpr_read_b3297; GFX908: v_accvgpr_read_b3298; GFX908: v_accvgpr_read_b3299; GFX908: v_accvgpr_read_b32100; GFX908: v_accvgpr_read_b32101; GFX908: v_accvgpr_read_b32102; GFX908: v_accvgpr_read_b32103; GFX908: v_accvgpr_read_b32104; GFX908: v_accvgpr_read_b32105; GFX908: v_accvgpr_read_b32106; GFX908: v_accvgpr_read_b32107; GFX908: v_accvgpr_read_b32108; GFX908: v_accvgpr_read_b32109define amdgpu_kernel void @v_mfma_f32_16x16x1f32_avaa(ptr addrspace(1) %arg) {110bb:111 %in.1 = load <16 x float>, ptr addrspace(1) %arg112 %mai.1 = tail call <16 x float> asm "v_mfma_f32_16x16x1f32 $0, $1, $2, $3", "=a,v,a,a"(float 1.0, float 2.0, <16 x float> %in.1)113 store <16 x float> %mai.1, ptr addrspace(1) %arg114 ret void115}116 117; GCN-LABEL: {{^}}v_mfma_f32_32x32x1f32_avaa118; GFX908: v_accvgpr_write_b32119; GFX908: v_accvgpr_write_b32120; GFX908: v_accvgpr_write_b32121; GFX908: v_accvgpr_write_b32122; GFX908: v_accvgpr_write_b32123; GFX908: v_accvgpr_write_b32124; GFX908: v_accvgpr_write_b32125; GFX908: v_accvgpr_write_b32126; GFX908: v_accvgpr_write_b32127; GFX908: v_accvgpr_write_b32128; GFX908: v_accvgpr_write_b32129; GFX908: v_accvgpr_write_b32130; GFX908: v_accvgpr_write_b32131; GFX908: v_accvgpr_write_b32132; GFX908: v_accvgpr_write_b32133; GFX908: v_accvgpr_write_b32134; GFX908: v_accvgpr_write_b32135; GFX908: v_accvgpr_write_b32136; GFX908: v_accvgpr_write_b32137; GFX908: v_accvgpr_write_b32138; GFX908: v_accvgpr_write_b32139; GFX908: v_accvgpr_write_b32140; GFX908: v_accvgpr_write_b32141; GFX908: v_accvgpr_write_b32142; GFX908: v_accvgpr_write_b32143; GFX908: v_accvgpr_write_b32144; GFX908: v_accvgpr_write_b32145; GFX908: v_accvgpr_write_b32146; GFX908: v_accvgpr_write_b32147; GFX908: v_accvgpr_write_b32148; GFX908: v_accvgpr_write_b32149; GFX908: v_accvgpr_write_b32150; GFX908: v_accvgpr_write_b32151; GFX908: v_mfma_f32_32x32x1f32 a[{{[0-9:]+}}], v{{[0-9]+}}, a{{[0-9]+}}, a[{{[0-9:]+}}]152; GFX908: v_accvgpr_read_b32153; GFX908: v_accvgpr_read_b32154; GFX908: v_accvgpr_read_b32155; GFX908: v_accvgpr_read_b32156; GFX908: v_accvgpr_read_b32157; GFX908: v_accvgpr_read_b32158; GFX908: v_accvgpr_read_b32159; GFX908: v_accvgpr_read_b32160; GFX908: v_accvgpr_read_b32161; GFX908: v_accvgpr_read_b32162; GFX908: v_accvgpr_read_b32163; GFX908: v_accvgpr_read_b32164; GFX908: v_accvgpr_read_b32165; GFX908: v_accvgpr_read_b32166; GFX908: v_accvgpr_read_b32167; GFX908: v_accvgpr_read_b32168; GFX908: v_accvgpr_read_b32169; GFX908: v_accvgpr_read_b32170; GFX908: v_accvgpr_read_b32171; GFX908: v_accvgpr_read_b32172; GFX908: v_accvgpr_read_b32173; GFX908: v_accvgpr_read_b32174; GFX908: v_accvgpr_read_b32175; GFX908: v_accvgpr_read_b32176; GFX908: v_accvgpr_read_b32177; GFX908: v_accvgpr_read_b32178; GFX908: v_accvgpr_read_b32179; GFX908: v_accvgpr_read_b32180; GFX908: v_accvgpr_read_b32181; GFX908: v_accvgpr_read_b32182; GFX908: v_accvgpr_read_b32183; GFX908: v_accvgpr_read_b32184define amdgpu_kernel void @v_mfma_f32_32x32x1f32_avaa(ptr addrspace(1) %arg) {185bb:186 %in.1 = load <32 x i32>, ptr addrspace(1) %arg187 %mai.1 = tail call <32 x i32> asm "v_mfma_f32_32x32x1f32 $0, $1, $2, $3", "=a,v,a,a"(float 1.0, float 2.0, <32 x i32> %in.1)188 store <32 x i32> %mai.1, ptr addrspace(1) %arg189 ret void190}191