authorgravatar for alex@alexrp.comAlex Rønne Petersen <alex@alexrp.com> 2026-10-01 00:47:42+02:00
committergravatar for alex@alexrp.comAlex Rønne Petersen <alex@alexrp.com> 2026-10-01 00:48:32+02:00
log92b739ac096dc65a243c5ddf51f94b43bf200650
tree34d8e7fc183ad3485d94580d528e46d5ccf3e463
parentcab5fea319542aa09e824f8957bbd26eb1d54657
signature Signed by SSH key SHA256:7B/LJ7bpR1eX8aCXSr4mtd5M45VMPKcx9zY8e95b5QM

llvm: fix @workGroupSize() on amdgcn to return u32 instead of u16

1 files changed, 2 insertions(+), 1 deletions(-)

src/codegen/llvm/FuncGen.zig+2-1
...@@ -6639,7 +6639,8 @@ fn airWorkGroupSize(self: *FuncGen, inst: Air.Inst.Index) Allocator.Error!Builde...@@ -6639,7 +6639,8 @@ fn airWorkGroupSize(self: *FuncGen, inst: Air.Inst.Index) Allocator.Error!Builde
6639 // Load the work_group_* member from the struct as u16.6639 // Load the work_group_* member from the struct as u16.
6640 // Just treat the dispatch pointer as an array of u16 to keep things simple.6640 // Just treat the dispatch pointer as an array of u16 to keep things simple.
6641 const workgroup_size_ptr = try self.ptraddConst(dispatch_ptr, (2 + dimension) * 2);6641 const workgroup_size_ptr = try self.ptraddConst(dispatch_ptr, (2 + dimension) * 2);
6642 return self.load(workgroup_size_ptr, .@"2", .u16, .normal);6642 const workgroup_size = try self.load(workgroup_size_ptr, .@"2", .u16, .normal);
6643 return try self.wip.cast(.zext, workgroup_size, .i32, "");
6643 },6644 },
6644 .nvptx, .nvptx64 => {6645 .nvptx, .nvptx64 => {
6645 return self.workIntrinsic(dimension, 1, "nvvm.read.ptx.sreg.ntid");6646 return self.workIntrinsic(dimension, 1, "nvvm.read.ptx.sreg.ntid");