1; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 < %s | FileCheck %s 2 3; Attribute not specified. 4; CHECK-LABEL: {{^}}empty_no_attribute: 5define amdgpu_kernel void @empty_no_attribute() { 6entry: 7 ret void 8} 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 void 15} 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 void 23} 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 void 31} 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 void 38} 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 void 45} 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 void 55} 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 void 63} 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 void 71} 72attributes #7 = {"amdgpu-max-num-workgroups"="1,2,4294967295"} 73 74 75; CHECK: .amdgpu_metadata 76; CHECK: - .args: 77; CHECK: .max_flat_workgroup_size: 1024 78; CHECK-NEXT: .name: empty_no_attribute 79; CHECK-NEXT: .private_segment_fixed_size: 0 80 81; CHECK: - .args: 82; CHECK: .max_flat_workgroup_size: 1024 83; CHECK-NEXT: .max_num_workgroups_y: 2 84; CHECK-NEXT: .max_num_workgroups_z: 3 85; CHECK-NEXT: .name: empty_max_num_workgroups_x0 86; CHECK-NEXT: .private_segment_fixed_size: 0 87 88; CHECK: - .args: 89; CHECK: .max_flat_workgroup_size: 1024 90; CHECK-NEXT: .max_num_workgroups_x: 1 91; CHECK-NEXT: .max_num_workgroups_z: 3 92; CHECK-NEXT: .name: empty_max_num_workgroups_y0 93; CHECK-NEXT: .private_segment_fixed_size: 0 94 95; CHECK: - .args: 96; CHECK: .max_flat_workgroup_size: 1024 97; CHECK-NEXT: .max_num_workgroups_x: 1 98; CHECK-NEXT: .max_num_workgroups_y: 2 99; CHECK-NEXT: .name: empty_max_num_workgroups_z0 100; CHECK-NEXT: .private_segment_fixed_size: 0 101 102; CHECK: - .args: 103; CHECK: .max_flat_workgroup_size: 1024 104; CHECK-NEXT: .max_num_workgroups_x: 1 105; CHECK-NEXT: .max_num_workgroups_y: 2 106; CHECK-NEXT: .max_num_workgroups_z: 3 107; CHECK-NEXT: .name: empty_max_num_workgroups_1_2_3 108; CHECK-NEXT: .private_segment_fixed_size: 0 109 110; CHECK: - .args: 111; CHECK: .max_flat_workgroup_size: 1024 112; CHECK-NEXT: .max_num_workgroups_x: 1024 113; CHECK-NEXT: .max_num_workgroups_y: 1024 114; CHECK-NEXT: .max_num_workgroups_z: 1024 115; CHECK-NEXT: .name: empty_max_num_workgroups_1024_1024_1024 116; CHECK-NEXT: .private_segment_fixed_size: 0 117 118 119; CHECK: - .args: 120; CHECK: .max_flat_workgroup_size: 1024 121; CHECK-NEXT: .max_num_workgroups_y: 2 122; CHECK-NEXT: .max_num_workgroups_z: 3 123; CHECK-NEXT: .name: empty_max_num_workgroups_x_max 124; CHECK-NEXT: .private_segment_fixed_size: 0 125 126; CHECK: - .args: 127; CHECK: .max_flat_workgroup_size: 1024 128; CHECK-NEXT: .max_num_workgroups_x: 1 129; CHECK-NEXT: .max_num_workgroups_z: 3 130; CHECK-NEXT: .name: empty_max_num_workgroups_y_max 131; CHECK-NEXT: .private_segment_fixed_size: 0 132 133; CHECK: - .args: 134; CHECK: .max_flat_workgroup_size: 1024 135; CHECK-NEXT: .max_num_workgroups_x: 1 136; CHECK-NEXT: .max_num_workgroups_y: 2 137; CHECK-NEXT: .name: empty_max_num_workgroups_z_max 138; CHECK-NEXT: .private_segment_fixed_size: 0 139