139 lines · plain
1; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 < %s | FileCheck %s2 3; Attribute not specified.4; CHECK-LABEL: {{^}}empty_no_attribute:5define amdgpu_kernel void @empty_no_attribute() {6entry:7 ret void8}9 10; Ignore if number of work groups for x dimension is 0.11; CHECK-LABEL: {{^}}empty_max_num_workgroups_x0:12define amdgpu_kernel void @empty_max_num_workgroups_x0() #0 {13entry:14 ret void15}16attributes #0 = {"amdgpu-max-num-workgroups"="0,2,3"}17 18; Ignore if number of work groups for y dimension is 0.19; CHECK-LABEL: {{^}}empty_max_num_workgroups_y0:20define amdgpu_kernel void @empty_max_num_workgroups_y0() #1 {21entry:22 ret void23}24attributes #1 = {"amdgpu-max-num-workgroups"="1,0,3"}25 26; Ignore if number of work groups for z dimension is 0.27; CHECK-LABEL: {{^}}empty_max_num_workgroups_z0:28define amdgpu_kernel void @empty_max_num_workgroups_z0() #2 {29entry:30 ret void31}32attributes #2 = {"amdgpu-max-num-workgroups"="1,2,0"}33 34; CHECK-LABEL: {{^}}empty_max_num_workgroups_1_2_3:35define amdgpu_kernel void @empty_max_num_workgroups_1_2_3() #3 {36entry:37 ret void38}39attributes #3 = {"amdgpu-max-num-workgroups"="1,2,3"}40 41; CHECK-LABEL: {{^}}empty_max_num_workgroups_1024_1024_1024:42define amdgpu_kernel void @empty_max_num_workgroups_1024_1024_1024() #4 {43entry:44 ret void45}46attributes #4 = {"amdgpu-max-num-workgroups"="1024,1024,1024"}47 48 49 50; Ignore if number of work groups for x dimension is 0.51; CHECK-LABEL: {{^}}empty_max_num_workgroups_x_max:52define amdgpu_kernel void @empty_max_num_workgroups_x_max() #5 {53entry:54 ret void55}56attributes #5 = {"amdgpu-max-num-workgroups"="4294967295,2,3"}57 58; Ignore if number of work groups for y dimension is 0.59; CHECK-LABEL: {{^}}empty_max_num_workgroups_y_max:60define amdgpu_kernel void @empty_max_num_workgroups_y_max() #6 {61entry:62 ret void63}64attributes #6 = {"amdgpu-max-num-workgroups"="1,4294967295,3"}65 66; Ignore if number of work groups for z dimension is 0.67; CHECK-LABEL: {{^}}empty_max_num_workgroups_z_max:68define amdgpu_kernel void @empty_max_num_workgroups_z_max() #7 {69entry:70 ret void71}72attributes #7 = {"amdgpu-max-num-workgroups"="1,2,4294967295"}73 74 75; CHECK: .amdgpu_metadata76; CHECK: - .args:77; CHECK: .max_flat_workgroup_size: 102478; CHECK-NEXT: .name: empty_no_attribute79; CHECK-NEXT: .private_segment_fixed_size: 080 81; CHECK: - .args:82; CHECK: .max_flat_workgroup_size: 102483; CHECK-NEXT: .max_num_workgroups_y: 284; CHECK-NEXT: .max_num_workgroups_z: 385; CHECK-NEXT: .name: empty_max_num_workgroups_x086; CHECK-NEXT: .private_segment_fixed_size: 087 88; CHECK: - .args:89; CHECK: .max_flat_workgroup_size: 102490; CHECK-NEXT: .max_num_workgroups_x: 191; CHECK-NEXT: .max_num_workgroups_z: 392; CHECK-NEXT: .name: empty_max_num_workgroups_y093; CHECK-NEXT: .private_segment_fixed_size: 094 95; CHECK: - .args:96; CHECK: .max_flat_workgroup_size: 102497; CHECK-NEXT: .max_num_workgroups_x: 198; CHECK-NEXT: .max_num_workgroups_y: 299; CHECK-NEXT: .name: empty_max_num_workgroups_z0100; CHECK-NEXT: .private_segment_fixed_size: 0101 102; CHECK: - .args:103; CHECK: .max_flat_workgroup_size: 1024104; CHECK-NEXT: .max_num_workgroups_x: 1105; CHECK-NEXT: .max_num_workgroups_y: 2106; CHECK-NEXT: .max_num_workgroups_z: 3107; CHECK-NEXT: .name: empty_max_num_workgroups_1_2_3108; CHECK-NEXT: .private_segment_fixed_size: 0109 110; CHECK: - .args:111; CHECK: .max_flat_workgroup_size: 1024112; CHECK-NEXT: .max_num_workgroups_x: 1024113; CHECK-NEXT: .max_num_workgroups_y: 1024114; CHECK-NEXT: .max_num_workgroups_z: 1024115; CHECK-NEXT: .name: empty_max_num_workgroups_1024_1024_1024116; CHECK-NEXT: .private_segment_fixed_size: 0117 118 119; CHECK: - .args:120; CHECK: .max_flat_workgroup_size: 1024121; CHECK-NEXT: .max_num_workgroups_y: 2122; CHECK-NEXT: .max_num_workgroups_z: 3123; CHECK-NEXT: .name: empty_max_num_workgroups_x_max124; CHECK-NEXT: .private_segment_fixed_size: 0125 126; CHECK: - .args:127; CHECK: .max_flat_workgroup_size: 1024128; CHECK-NEXT: .max_num_workgroups_x: 1129; CHECK-NEXT: .max_num_workgroups_z: 3130; CHECK-NEXT: .name: empty_max_num_workgroups_y_max131; CHECK-NEXT: .private_segment_fixed_size: 0132 133; CHECK: - .args:134; CHECK: .max_flat_workgroup_size: 1024135; CHECK-NEXT: .max_num_workgroups_x: 1136; CHECK-NEXT: .max_num_workgroups_y: 2137; CHECK-NEXT: .name: empty_max_num_workgroups_z_max138; CHECK-NEXT: .private_segment_fixed_size: 0139