Avoid invalid 0 case

Author: arsenm

llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp

@@ -505,16 +505,18 @@ MetadataStreamerMsgPackV4::getHSAKernelProps(const MachineFunction &MF,
   Kern[".max_flat_workgroup_size"] =
       Kern.getDocument()->getNode(MFI.getMaxFlatWorkGroupSize());
 
-  uint32_t NumWGX = MFI.getMaxNumWorkGroupsX();
   uint32_t NumWGY = MFI.getMaxNumWorkGroupsY();
   uint32_t NumWGZ = MFI.getMaxNumWorkGroupsZ();
-  if (NumWGX != std::numeric_limits<uint32_t>::max())
+  uint32_t NumWGX = MFI.getMaxNumWorkGroupsX();
+
+  // TODO: Should consider 0 invalid and reject in IR verifier.
+  if (NumWGX != std::numeric_limits<uint32_t>::max() && NumWGX != 0)
     Kern[".max_num_workgroups_x"] = Kern.getDocument()->getNode(NumWGX);
 
-  if (NumWGY != std::numeric_limits<uint32_t>::max())
+  if (NumWGY != std::numeric_limits<uint32_t>::max() && NumWGY != 0)
     Kern[".max_num_workgroups_y"] = Kern.getDocument()->getNode(NumWGY);
 
-  if (NumWGZ != std::numeric_limits<uint32_t>::max())
+  if (NumWGZ != std::numeric_limits<uint32_t>::max() && NumWGZ != 0)
     Kern[".max_num_workgroups_z"] = Kern.getDocument()->getNode(NumWGZ);
 
   Kern[".sgpr_spill_count"] =

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

@@ -72,7 +72,6 @@ entry:
 attributes #7 = {"amdgpu-max-num-workgroups"="1,2,4294967295"}
 
 
-
 ; CHECK: .amdgpu_metadata
 ; CHECK: - .args:
 ; CHECK:        .max_flat_workgroup_size: 1024
@@ -81,7 +80,6 @@ attributes #7 = {"amdgpu-max-num-workgroups"="1,2,4294967295"}
 
 ; CHECK: - .args:
 ; CHECK:        .max_flat_workgroup_size: 1024
-; CHECK-NEXT:   .max_num_workgroups_x: 0
 ; CHECK-NEXT:   .max_num_workgroups_y: 2
 ; CHECK-NEXT:   .max_num_workgroups_z: 3
 ; CHECK-NEXT:   .name:           empty_max_num_workgroups_x0
@@ -90,7 +88,6 @@ attributes #7 = {"amdgpu-max-num-workgroups"="1,2,4294967295"}
 ; CHECK: - .args:
 ; CHECK:        .max_flat_workgroup_size: 1024
 ; CHECK-NEXT:   .max_num_workgroups_x: 1
-; CHECK-NEXT:   .max_num_workgroups_y: 0
 ; CHECK-NEXT:   .max_num_workgroups_z: 3
 ; CHECK-NEXT:   .name:           empty_max_num_workgroups_y0
 ; CHECK-NEXT:   .private_segment_fixed_size: 0
@@ -99,7 +96,6 @@ attributes #7 = {"amdgpu-max-num-workgroups"="1,2,4294967295"}
 ; CHECK:        .max_flat_workgroup_size: 1024
 ; CHECK-NEXT:   .max_num_workgroups_x: 1
 ; CHECK-NEXT:   .max_num_workgroups_y: 2
-; CHECK-NEXT:   .max_num_workgroups_z: 0
 ; CHECK-NEXT:   .name:           empty_max_num_workgroups_z0
 ; CHECK-NEXT:   .private_segment_fixed_size: 0