Skip to content

Commit fada8ca

Browse files
committed
Avoid invalid 0 case
1 parent e15235f commit fada8ca

File tree

2 files changed

+6
-8
lines changed

2 files changed

+6
-8
lines changed

llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp

Lines changed: 6 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -505,16 +505,18 @@ MetadataStreamerMsgPackV4::getHSAKernelProps(const MachineFunction &MF,
505505
Kern[".max_flat_workgroup_size"] =
506506
Kern.getDocument()->getNode(MFI.getMaxFlatWorkGroupSize());
507507

508-
uint32_t NumWGX = MFI.getMaxNumWorkGroupsX();
509508
uint32_t NumWGY = MFI.getMaxNumWorkGroupsY();
510509
uint32_t NumWGZ = MFI.getMaxNumWorkGroupsZ();
511-
if (NumWGX != std::numeric_limits<uint32_t>::max())
510+
uint32_t NumWGX = MFI.getMaxNumWorkGroupsX();
511+
512+
// TODO: Should consider 0 invalid and reject in IR verifier.
513+
if (NumWGX != std::numeric_limits<uint32_t>::max() && NumWGX != 0)
512514
Kern[".max_num_workgroups_x"] = Kern.getDocument()->getNode(NumWGX);
513515

514-
if (NumWGY != std::numeric_limits<uint32_t>::max())
516+
if (NumWGY != std::numeric_limits<uint32_t>::max() && NumWGY != 0)
515517
Kern[".max_num_workgroups_y"] = Kern.getDocument()->getNode(NumWGY);
516518

517-
if (NumWGZ != std::numeric_limits<uint32_t>::max())
519+
if (NumWGZ != std::numeric_limits<uint32_t>::max() && NumWGZ != 0)
518520
Kern[".max_num_workgroups_z"] = Kern.getDocument()->getNode(NumWGZ);
519521

520522
Kern[".sgpr_spill_count"] =

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

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -72,7 +72,6 @@ entry:
7272
attributes #7 = {"amdgpu-max-num-workgroups"="1,2,4294967295"}
7373

7474

75-
7675
; CHECK: .amdgpu_metadata
7776
; CHECK: - .args:
7877
; CHECK: .max_flat_workgroup_size: 1024
@@ -81,7 +80,6 @@ attributes #7 = {"amdgpu-max-num-workgroups"="1,2,4294967295"}
8180

8281
; CHECK: - .args:
8382
; CHECK: .max_flat_workgroup_size: 1024
84-
; CHECK-NEXT: .max_num_workgroups_x: 0
8583
; CHECK-NEXT: .max_num_workgroups_y: 2
8684
; CHECK-NEXT: .max_num_workgroups_z: 3
8785
; CHECK-NEXT: .name: empty_max_num_workgroups_x0
@@ -90,7 +88,6 @@ attributes #7 = {"amdgpu-max-num-workgroups"="1,2,4294967295"}
9088
; CHECK: - .args:
9189
; CHECK: .max_flat_workgroup_size: 1024
9290
; CHECK-NEXT: .max_num_workgroups_x: 1
93-
; CHECK-NEXT: .max_num_workgroups_y: 0
9491
; CHECK-NEXT: .max_num_workgroups_z: 3
9592
; CHECK-NEXT: .name: empty_max_num_workgroups_y0
9693
; CHECK-NEXT: .private_segment_fixed_size: 0
@@ -99,7 +96,6 @@ attributes #7 = {"amdgpu-max-num-workgroups"="1,2,4294967295"}
9996
; CHECK: .max_flat_workgroup_size: 1024
10097
; CHECK-NEXT: .max_num_workgroups_x: 1
10198
; CHECK-NEXT: .max_num_workgroups_y: 2
102-
; CHECK-NEXT: .max_num_workgroups_z: 0
10399
; CHECK-NEXT: .name: empty_max_num_workgroups_z0
104100
; CHECK-NEXT: .private_segment_fixed_size: 0
105101

0 commit comments

Comments
 (0)