From 92b739ac096dc65a243c5ddf51f94b43bf200650 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Alex=20R=C3=B8nne=20Petersen?= Date: Thu, 1 Oct 2026 00:47:42 +0200 Subject: [PATCH] llvm: fix @workGroupSize() on amdgcn to return u32 instead of u16 --- src/codegen/llvm/FuncGen.zig | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/src/codegen/llvm/FuncGen.zig b/src/codegen/llvm/FuncGen.zig index 8d098e26bd8968b709df6415404b1552b2d34721..4a776752e8342365bf07ea87595083048aff0903 100644 --- a/src/codegen/llvm/FuncGen.zig +++ b/src/codegen/llvm/FuncGen.zig @@ -6639,7 +6639,8 @@ fn airWorkGroupSize(self: *FuncGen, inst: Air.Inst.Index) Allocator.Error!Builde // Load the work_group_* member from the struct as u16. // Just treat the dispatch pointer as an array of u16 to keep things simple. const workgroup_size_ptr = try self.ptraddConst(dispatch_ptr, (2 + dimension) * 2); - return self.load(workgroup_size_ptr, .@"2", .u16, .normal); + const workgroup_size = try self.load(workgroup_size_ptr, .@"2", .u16, .normal); + return try self.wip.cast(.zext, workgroup_size, .i32, ""); }, .nvptx, .nvptx64 => { return self.workIntrinsic(dimension, 1, "nvvm.read.ptx.sreg.ntid"); -- 2.54.0