Skip to content

Commit 15cb5b5

Browse files
author
Jun Wang
committed
Change the LLVM attribute name from amdgpu-max-num-work-groups
to amdgpu-max-num-workgroups; clang attribute name unchanged.
1 parent 4d0ab6e commit 15cb5b5

File tree

10 files changed

+170
-170
lines changed

10 files changed

+170
-170
lines changed

‎clang/lib/CodeGen/Targets/AMDGPU.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -377,7 +377,7 @@ void AMDGPUTargetCodeGenInfo::setFunctionDeclAttributes(
377377
llvm::raw_svector_ostream OS(AttrVal);
378378
OS << X << ',' << Y << ',' << Z;
379379

380-
F->addFnAttr("amdgpu-max-num-work-groups", AttrVal.str());
380+
F->addFnAttr("amdgpu-max-num-workgroups", AttrVal.str());
381381
}
382382
}
383383

‎clang/test/CodeGenCUDA/amdgpu-kernel-attrs.cu

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -86,7 +86,7 @@ template __global__ void template_32_4_a_max_num_work_groups<2>();
8686
// CHECK-DAG: attributes [[WAVES_PER_EU_2]] = {{.*}}"amdgpu-waves-per-eu"="2"
8787
// CHECK-DAG: attributes [[NUM_SGPR_32]] = {{.*}}"amdgpu-num-sgpr"="32"
8888
// CHECK-DAG: attributes [[NUM_VGPR_64]] = {{.*}}"amdgpu-num-vgpr"="64"
89-
// CHECK-DAG: attributes [[MAX_NUM_WORK_GROUPS_32_4_2]] = {{.*}}"amdgpu-max-num-work-groups"="32,4,2"
90-
// CHECK-DAG: attributes [[MAX_NUM_WORK_GROUPS_32_1_1]] = {{.*}}"amdgpu-max-num-work-groups"="32,1,1"
89+
// CHECK-DAG: attributes [[MAX_NUM_WORK_GROUPS_32_4_2]] = {{.*}}"amdgpu-max-num-workgroups"="32,4,2"
90+
// CHECK-DAG: attributes [[MAX_NUM_WORK_GROUPS_32_1_1]] = {{.*}}"amdgpu-max-num-workgroups"="32,1,1"
9191

9292
// NOUB-NOT: "uniform-work-group-size"="true"

‎clang/test/CodeGenOpenCL/amdgpu-attrs.cl

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -229,12 +229,12 @@ kernel void default_kernel() {
229229
// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_NUM_SGPR_32_NUM_VGPR_64]] = {{.*}} "amdgpu-flat-work-group-size"="32,64" "amdgpu-num-sgpr"="32" "amdgpu-num-vgpr"="64" "amdgpu-waves-per-eu"="2"
230230
// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4_NUM_SGPR_32_NUM_VGPR_64]] = {{.*}} "amdgpu-flat-work-group-size"="32,64" "amdgpu-num-sgpr"="32" "amdgpu-num-vgpr"="64" "amdgpu-waves-per-eu"="2,4"
231231

232-
// CHECK-DAG: attributes [[MAX_NUM_WORK_GROUPS_1_1_1]] = {{.*}} "amdgpu-max-num-work-groups"="1,1,1"
233-
// CHECK-DAG: attributes [[MAX_NUM_WORK_GROUPS_32_1_1]] = {{.*}} "amdgpu-max-num-work-groups"="32,1,1"
234-
// CHECK-DAG: attributes [[MAX_NUM_WORK_GROUPS_32_8_1]] = {{.*}} "amdgpu-max-num-work-groups"="32,8,1"
235-
// CHECK-DAG: attributes [[MAX_NUM_WORK_GROUPS_1_1_32]] = {{.*}} "amdgpu-max-num-work-groups"="1,1,32"
236-
// CHECK-DAG: attributes [[MAX_NUM_WORK_GROUPS_1_8_32]] = {{.*}} "amdgpu-max-num-work-groups"="1,8,32"
237-
// CHECK-DAG: attributes [[MAX_NUM_WORK_GROUPS_4_8_32]] = {{.*}} "amdgpu-max-num-work-groups"="4,8,32"
232+
// CHECK-DAG: attributes [[MAX_NUM_WORK_GROUPS_1_1_1]] = {{.*}} "amdgpu-max-num-workgroups"="1,1,1"
233+
// CHECK-DAG: attributes [[MAX_NUM_WORK_GROUPS_32_1_1]] = {{.*}} "amdgpu-max-num-workgroups"="32,1,1"
234+
// CHECK-DAG: attributes [[MAX_NUM_WORK_GROUPS_32_8_1]] = {{.*}} "amdgpu-max-num-workgroups"="32,8,1"
235+
// CHECK-DAG: attributes [[MAX_NUM_WORK_GROUPS_1_1_32]] = {{.*}} "amdgpu-max-num-workgroups"="1,1,32"
236+
// CHECK-DAG: attributes [[MAX_NUM_WORK_GROUPS_1_8_32]] = {{.*}} "amdgpu-max-num-workgroups"="1,8,32"
237+
// CHECK-DAG: attributes [[MAX_NUM_WORK_GROUPS_4_8_32]] = {{.*}} "amdgpu-max-num-workgroups"="4,8,32"
238238

239239
// CHECK-DAG: attributes [[A_FUNCTION]] = {{.*}}
240240
// CHECK-DAG: attributes [[DEFAULT_KERNEL_ATTRS]] = {{.*}} "amdgpu-flat-work-group-size"="1,256"

‎llvm/docs/AMDGPUUsage.rst

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1436,7 +1436,7 @@ The AMDGPU backend supports the following LLVM IR attributes.
14361436
the frame. This is an internal detail of how LDS variables are lowered,
14371437
language front ends should not set this attribute.
14381438

1439-
"amdgpu-max-num-work-groups"="x,y,z" Specify the maximum number of work groups for the kernel dispatch in the
1439+
"amdgpu-max-num-workgroups"="x,y,z" Specify the maximum number of work groups for the kernel dispatch in the
14401440
X, Y, and Z dimensions. Generated by the ``amdgpu_max_num_work_groups``
14411441
CLANG attribute [CLANG-ATTR]_. Clang only emits this attribute when all
14421442
the three numbers are >= 1.

‎llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -498,9 +498,9 @@ MetadataStreamerMsgPackV4::getHSAKernelProps(const MachineFunction &MF,
498498
unsigned NumWGY = MFI.getMaxNumWorkGroupsY();
499499
unsigned NumWGZ = MFI.getMaxNumWorkGroupsZ();
500500
if (NumWGX != 0 && NumWGY != 0 && NumWGZ != 0) {
501-
Kern[".max_num_work_groups_x"] = Kern.getDocument()->getNode(NumWGX);
502-
Kern[".max_num_work_groups_y"] = Kern.getDocument()->getNode(NumWGY);
503-
Kern[".max_num_work_groups_z"] = Kern.getDocument()->getNode(NumWGZ);
501+
Kern[".max_num_workgroups_x"] = Kern.getDocument()->getNode(NumWGX);
502+
Kern[".max_num_workgroups_y"] = Kern.getDocument()->getNode(NumWGY);
503+
Kern[".max_num_workgroups_z"] = Kern.getDocument()->getNode(NumWGZ);
504504
}
505505
Kern[".sgpr_spill_count"] =
506506
Kern.getDocument()->getNode(MFI.getNumSpilledSGPRs());

‎llvm/lib/Target/AMDGPU/AMDGPUSubtarget.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -432,7 +432,7 @@ std::pair<unsigned, unsigned> AMDGPUSubtarget::getEffectiveWavesPerEU(
432432
std::pair<unsigned, unsigned> Default(1, getMaxWavesPerEU());
433433

434434
// If minimum/maximum flat work group sizes were explicitly requested using
435-
// "amdgpu-flat-work-group-size" attribute, then set default minimum/maximum
435+
// "amdgpu-flat-workgroup-size" attribute, then set default minimum/maximum
436436
// number of waves per execution unit to values implied by requested
437437
// minimum/maximum flat work group sizes.
438438
unsigned MinImpliedByFlatWorkGroupSize =
@@ -1111,5 +1111,5 @@ unsigned GCNUserSGPRUsageInfo::getNumFreeUserSGPRs() {
11111111

11121112
SmallVector<unsigned>
11131113
AMDGPUSubtarget::getMaxNumWorkGroups(const Function &F) const {
1114-
return AMDGPU::getIntegerVecAttribute(F, "amdgpu-max-num-work-groups", 3);
1114+
return AMDGPU::getIntegerVecAttribute(F, "amdgpu-max-num-workgroups", 3);
11151115
}

‎llvm/test/CodeGen/AMDGPU/attr-amdgpu-num-work-groups.ll

Lines changed: 0 additions & 84 deletions
This file was deleted.

‎llvm/test/CodeGen/AMDGPU/attr-amdgpu-num-work-groups_error_check.ll

Lines changed: 0 additions & 71 deletions
This file was deleted.

‎llvm/test/CodeGen/AMDGPU/attr-amdgpu-num-workgroups.ll

Lines changed: 84 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,84 @@
1+
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 < %s | FileCheck %s
2+
3+
; Attribute not specified.
4+
; CHECK-LABEL: {{^}}empty_no_attribute:
5+
define amdgpu_kernel void @empty_no_attribute() {
6+
entry:
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:
12+
define amdgpu_kernel void @empty_max_num_workgroups_x0() #0 {
13+
entry:
14+
ret void
15+
}
16+
attributes #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:
20+
define amdgpu_kernel void @empty_max_num_workgroups_y0() #1 {
21+
entry:
22+
ret void
23+
}
24+
attributes #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:
28+
define amdgpu_kernel void @empty_max_num_workgroups_z0() #2 {
29+
entry:
30+
ret void
31+
}
32+
attributes #2 = {"amdgpu-max-num-workgroups"="1,2,0"}
33+
34+
; CHECK-LABEL: {{^}}empty_max_num_workgroups_1_2_3:
35+
define amdgpu_kernel void @empty_max_num_workgroups_1_2_3() #3 {
36+
entry:
37+
ret void
38+
}
39+
attributes #3 = {"amdgpu-max-num-workgroups"="1,2,3"}
40+
41+
; CHECK-LABEL: {{^}}empty_max_num_workgroups_1024_1024_1024:
42+
define amdgpu_kernel void @empty_max_num_workgroups_1024_1024_1024() #4 {
43+
entry:
44+
ret void
45+
}
46+
attributes #4 = {"amdgpu-max-num-workgroups"="1024,1024,1024"}
47+
48+
49+
; CHECK: .amdgpu_metadata
50+
; CHECK: - .args:
51+
; CHECK: .max_flat_workgroup_size: 1024
52+
; CHECK-NEXT: .name: empty_no_attribute
53+
; CHECK-NEXT: .private_segment_fixed_size: 0
54+
55+
; CHECK: - .args:
56+
; CHECK: .max_flat_workgroup_size: 1024
57+
; CHECK-NEXT: .name: empty_max_num_workgroups_x0
58+
; CHECK-NEXT: .private_segment_fixed_size: 0
59+
60+
; CHECK: - .args:
61+
; CHECK: .max_flat_workgroup_size: 1024
62+
; CHECK-NEXT: .name: empty_max_num_workgroups_y0
63+
; CHECK-NEXT: .private_segment_fixed_size: 0
64+
65+
; CHECK: - .args:
66+
; CHECK: .max_flat_workgroup_size: 1024
67+
; CHECK-NEXT: .name: empty_max_num_workgroups_z0
68+
; CHECK-NEXT: .private_segment_fixed_size: 0
69+
70+
; CHECK: - .args:
71+
; CHECK: .max_flat_workgroup_size: 1024
72+
; CHECK-NEXT: .max_num_workgroups_x: 1
73+
; CHECK-NEXT: .max_num_workgroups_y: 2
74+
; CHECK-NEXT: .max_num_workgroups_z: 3
75+
; CHECK-NEXT: .name: empty_max_num_workgroups_1_2_3
76+
; CHECK-NEXT: .private_segment_fixed_size: 0
77+
78+
; CHECK: - .args:
79+
; CHECK: .max_flat_workgroup_size: 1024
80+
; CHECK-NEXT: .max_num_workgroups_x: 1024
81+
; CHECK-NEXT: .max_num_workgroups_y: 1024
82+
; CHECK-NEXT: .max_num_workgroups_z: 1024
83+
; CHECK-NEXT: .name: empty_max_num_workgroups_1024_1024_1024
84+
; CHECK-NEXT: .private_segment_fixed_size: 0

‎llvm/test/CodeGen/AMDGPU/attr-amdgpu-num-workgroups_error_check.ll

Lines changed: 71 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,71 @@
1+
; RUN: not llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 < %s 2>&1 | FileCheck --check-prefix=ERROR %s
2+
3+
; ERROR: error: can't parse integer attribute -1 in amdgpu-max-num-workgroups
4+
define amdgpu_kernel void @empty_max_num_workgroups_neg_num1() #21 {
5+
entry:
6+
ret void
7+
}
8+
attributes #21 = {"amdgpu-max-num-workgroups"="-1,2,3"}
9+
10+
; ERROR: error: can't parse integer attribute -2 in amdgpu-max-num-workgroups
11+
define amdgpu_kernel void @empty_max_num_workgroups_neg_num2() #22 {
12+
entry:
13+
ret void
14+
}
15+
attributes #22 = {"amdgpu-max-num-workgroups"="1,-2,3"}
16+
17+
; ERROR: error: can't parse integer attribute -3 in amdgpu-max-num-workgroups
18+
define amdgpu_kernel void @empty_max_num_workgroups_neg_num3() #23 {
19+
entry:
20+
ret void
21+
}
22+
attributes #23 = {"amdgpu-max-num-workgroups"="1,2,-3"}
23+
24+
; ERROR: error: can't parse integer attribute 1.0 in amdgpu-max-num-workgroups
25+
define amdgpu_kernel void @empty_max_num_workgroups_non_int1() #31 {
26+
entry:
27+
ret void
28+
}
29+
attributes #31 = {"amdgpu-max-num-workgroups"="1.0,2,3"}
30+
31+
; ERROR: error: can't parse integer attribute 2.0 in amdgpu-max-num-workgroups
32+
define amdgpu_kernel void @empty_max_num_workgroups_non_int2() #32 {
33+
entry:
34+
ret void
35+
}
36+
attributes #32 = {"amdgpu-max-num-workgroups"="1,2.0,3"}
37+
38+
; ERROR: error: can't parse integer attribute 3.0 in amdgpu-max-num-workgroups
39+
define amdgpu_kernel void @empty_max_num_workgroups_non_int3() #33 {
40+
entry:
41+
ret void
42+
}
43+
attributes #33 = {"amdgpu-max-num-workgroups"="1,2,3.0"}
44+
45+
; ERROR: error: can't parse integer attribute 10000000000 in amdgpu-max-num-workgroups
46+
define amdgpu_kernel void @empty_max_num_workgroups_too_large() #41 {
47+
entry:
48+
ret void
49+
}
50+
attributes #41 = {"amdgpu-max-num-workgroups"="10000000000,2,3"}
51+
52+
; ERROR: error: attribute amdgpu-max-num-workgroups has incorrect number of integers; expected 3
53+
define amdgpu_kernel void @empty_max_num_workgroups_1_arg() #51 {
54+
entry:
55+
ret void
56+
}
57+
attributes #51 = {"amdgpu-max-num-workgroups"="1"}
58+
59+
; ERROR: error: attribute amdgpu-max-num-workgroups has incorrect number of integers; expected 3
60+
define amdgpu_kernel void @empty_max_num_workgroups_2_args() #52 {
61+
entry:
62+
ret void
63+
}
64+
attributes #52 = {"amdgpu-max-num-workgroups"="1,2"}
65+
66+
; ERROR: error: attribute amdgpu-max-num-workgroups has incorrect number of integers; expected 3
67+
define amdgpu_kernel void @empty_max_num_workgroups_4_args() #53 {
68+
entry:
69+
ret void
70+
}
71+
attributes #53 = {"amdgpu-max-num-workgroups"="1,2,3,4"}

0 commit comments

Comments
 (0)