144 lines · plain
1; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx803 < %s | FileCheck --check-prefix=CHECK %s2; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx803 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefix=HSAMD %s3 4; CHECK-LABEL: {{^}}min_64_max_64:5; CHECK: SGPRBlocks: 06; CHECK: VGPRBlocks: 07; CHECK: NumSGPRsForWavesPerEU: 18; CHECK: NumVGPRsForWavesPerEU: 19define amdgpu_kernel void @min_64_max_64() #0 {10entry:11 ret void12}13attributes #0 = {"amdgpu-flat-work-group-size"="64,64"}14 15; CHECK-LABEL: {{^}}min_64_max_128:16; CHECK: SGPRBlocks: 017; CHECK: VGPRBlocks: 018; CHECK: NumSGPRsForWavesPerEU: 119; CHECK: NumVGPRsForWavesPerEU: 120define amdgpu_kernel void @min_64_max_128() #1 {21entry:22 ret void23}24attributes #1 = {"amdgpu-flat-work-group-size"="64,128"}25 26; CHECK-LABEL: {{^}}min_128_max_128:27; CHECK: SGPRBlocks: 828; CHECK: VGPRBlocks: 729; CHECK: NumSGPRsForWavesPerEU: 6530; CHECK: NumVGPRsForWavesPerEU: 2931define amdgpu_kernel void @min_128_max_128() #2 {32entry:33 ret void34}35attributes #2 = {"amdgpu-flat-work-group-size"="128,128"}36 37; CHECK-LABEL: {{^}}min_1024_max_102438; CHECK: SGPRBlocks: 839; CHECK: VGPRBlocks: 1040; CHECK: NumSGPRsForWavesPerEU: 6541; CHECK: NumVGPRsForWavesPerEU: 4342@var = addrspace(1) global float 0.043define amdgpu_kernel void @min_1024_max_1024() #3 {44 %val0 = load volatile float, ptr addrspace(1) @var45 %val1 = load volatile float, ptr addrspace(1) @var46 %val2 = load volatile float, ptr addrspace(1) @var47 %val3 = load volatile float, ptr addrspace(1) @var48 %val4 = load volatile float, ptr addrspace(1) @var49 %val5 = load volatile float, ptr addrspace(1) @var50 %val6 = load volatile float, ptr addrspace(1) @var51 %val7 = load volatile float, ptr addrspace(1) @var52 %val8 = load volatile float, ptr addrspace(1) @var53 %val9 = load volatile float, ptr addrspace(1) @var54 %val10 = load volatile float, ptr addrspace(1) @var55 %val11 = load volatile float, ptr addrspace(1) @var56 %val12 = load volatile float, ptr addrspace(1) @var57 %val13 = load volatile float, ptr addrspace(1) @var58 %val14 = load volatile float, ptr addrspace(1) @var59 %val15 = load volatile float, ptr addrspace(1) @var60 %val16 = load volatile float, ptr addrspace(1) @var61 %val17 = load volatile float, ptr addrspace(1) @var62 %val18 = load volatile float, ptr addrspace(1) @var63 %val19 = load volatile float, ptr addrspace(1) @var64 %val20 = load volatile float, ptr addrspace(1) @var65 %val21 = load volatile float, ptr addrspace(1) @var66 %val22 = load volatile float, ptr addrspace(1) @var67 %val23 = load volatile float, ptr addrspace(1) @var68 %val24 = load volatile float, ptr addrspace(1) @var69 %val25 = load volatile float, ptr addrspace(1) @var70 %val26 = load volatile float, ptr addrspace(1) @var71 %val27 = load volatile float, ptr addrspace(1) @var72 %val28 = load volatile float, ptr addrspace(1) @var73 %val29 = load volatile float, ptr addrspace(1) @var74 %val30 = load volatile float, ptr addrspace(1) @var75 %val31 = load volatile float, ptr addrspace(1) @var76 %val32 = load volatile float, ptr addrspace(1) @var77 %val33 = load volatile float, ptr addrspace(1) @var78 %val34 = load volatile float, ptr addrspace(1) @var79 %val35 = load volatile float, ptr addrspace(1) @var80 %val36 = load volatile float, ptr addrspace(1) @var81 %val37 = load volatile float, ptr addrspace(1) @var82 %val38 = load volatile float, ptr addrspace(1) @var83 %val39 = load volatile float, ptr addrspace(1) @var84 %val40 = load volatile float, ptr addrspace(1) @var85 86 store volatile float %val0, ptr addrspace(1) @var87 store volatile float %val1, ptr addrspace(1) @var88 store volatile float %val2, ptr addrspace(1) @var89 store volatile float %val3, ptr addrspace(1) @var90 store volatile float %val4, ptr addrspace(1) @var91 store volatile float %val5, ptr addrspace(1) @var92 store volatile float %val6, ptr addrspace(1) @var93 store volatile float %val7, ptr addrspace(1) @var94 store volatile float %val8, ptr addrspace(1) @var95 store volatile float %val9, ptr addrspace(1) @var96 store volatile float %val10, ptr addrspace(1) @var97 store volatile float %val11, ptr addrspace(1) @var98 store volatile float %val12, ptr addrspace(1) @var99 store volatile float %val13, ptr addrspace(1) @var100 store volatile float %val14, ptr addrspace(1) @var101 store volatile float %val15, ptr addrspace(1) @var102 store volatile float %val16, ptr addrspace(1) @var103 store volatile float %val17, ptr addrspace(1) @var104 store volatile float %val18, ptr addrspace(1) @var105 store volatile float %val19, ptr addrspace(1) @var106 store volatile float %val20, ptr addrspace(1) @var107 store volatile float %val21, ptr addrspace(1) @var108 store volatile float %val22, ptr addrspace(1) @var109 store volatile float %val23, ptr addrspace(1) @var110 store volatile float %val24, ptr addrspace(1) @var111 store volatile float %val25, ptr addrspace(1) @var112 store volatile float %val26, ptr addrspace(1) @var113 store volatile float %val27, ptr addrspace(1) @var114 store volatile float %val28, ptr addrspace(1) @var115 store volatile float %val29, ptr addrspace(1) @var116 store volatile float %val30, ptr addrspace(1) @var117 store volatile float %val31, ptr addrspace(1) @var118 store volatile float %val32, ptr addrspace(1) @var119 store volatile float %val33, ptr addrspace(1) @var120 store volatile float %val34, ptr addrspace(1) @var121 store volatile float %val35, ptr addrspace(1) @var122 store volatile float %val36, ptr addrspace(1) @var123 store volatile float %val37, ptr addrspace(1) @var124 store volatile float %val38, ptr addrspace(1) @var125 store volatile float %val39, ptr addrspace(1) @var126 store volatile float %val40, ptr addrspace(1) @var127 128 ret void129}130attributes #3 = {"amdgpu-flat-work-group-size"="1024,1024"}131 132!llvm.module.flags = !{!0}133!0 = !{i32 1, !"amdhsa_code_object_version", i32 400}134 135; HSAMD: amdhsa.kernels136; HSAMD: .max_flat_workgroup_size: 64137; HSAMD: .name: min_64_max_64138; HSAMD: .max_flat_workgroup_size: 128139; HSAMD: .name: min_64_max_128140; HSAMD: .max_flat_workgroup_size: 128141; HSAMD: .name: min_128_max_128142; HSAMD: .max_flat_workgroup_size: 1024143; HSAMD: .name: min_1024_max_1024144