authorgravatar for robin@voetter.nlRobin Voetter <robin@voetter.nl> 2024-08-11 12:22:31+02:00
committergravatar for robin@voetter.nlRobin Voetter <robin@voetter.nl> 2024-08-13 20:07:23+02:00
logbcfc7cf13cd3eb16f4b864efac5269d68200b070
treed570452b4551f314f5ea435b471587d320e661c6
parent76f062690c3373b24170e3e1d7dd62735eb9685e
signaturebadge-check Signed by SSH key SHA256:ZS52FNyUv2WUXvO4njmVaFVO46RHojFuOrxRc4LuKzg

nvptx: add implementations for GPU builtins

3 files changed, 113 insertions(+), 20 deletions(-)

src/Sema.zig+1-1
......@@ -26703,7 +26703,7 @@ fn zirWorkItem(
2670326703
2670426704 switch (target.cpu.arch) {
2670526705 // TODO: Allow for other GPU targets.
26706 .amdgcn, .spirv, .spirv64, .spirv32 => {},
26706 .amdgcn, .spirv, .spirv64, .spirv32, .nvptx, .nvptx64 => {},
2670726707 else => {
2670826708 return sema.fail(block, builtin_src, "builtin only available on GPU targets; targeted architecture is {s}", .{@tagName(target.cpu.arch)});
2670926709 },
src/codegen/llvm.zig+35-19
......@@ -10286,7 +10286,7 @@ pub const FuncGen = struct {
1028610286 return self.wip.cast(.addrspacecast, operand, try o.lowerType(inst_ty), "");
1028710287 }
1028810288
10289 fn amdgcnWorkIntrinsic(
10289 fn workIntrinsic(
1029010290 self: *FuncGen,
1029110291 dimension: u32,
1029210292 default: u32,
......@@ -10303,44 +10303,60 @@ pub const FuncGen = struct {
1030310303 fn airWorkItemId(self: *FuncGen, inst: Air.Inst.Index) !Builder.Value {
1030410304 const o = self.ng.object;
1030510305 const target = o.pt.zcu.getTarget();
10306 assert(target.cpu.arch == .amdgcn); // TODO is to port this function to other GPU architectures
1030710306
1030810307 const pl_op = self.air.instructions.items(.data)[@intFromEnum(inst)].pl_op;
1030910308 const dimension = pl_op.payload;
10310 return self.amdgcnWorkIntrinsic(dimension, 0, "amdgcn.workitem.id");
10309
10310 return switch (target.cpu.arch) {
10311 .amdgcn => self.workIntrinsic(dimension, 0, "amdgcn.workitem.id"),
10312 .nvptx, .nvptx64 => self.workIntrinsic(dimension, 0, "nvvm.read.ptx.sreg.tid"),
10313 else => unreachable,
10314 };
1031110315 }
1031210316
1031310317 fn airWorkGroupSize(self: *FuncGen, inst: Air.Inst.Index) !Builder.Value {
1031410318 const o = self.ng.object;
1031510319 const target = o.pt.zcu.getTarget();
10316 assert(target.cpu.arch == .amdgcn); // TODO is to port this function to other GPU architectures
1031710320
1031810321 const pl_op = self.air.instructions.items(.data)[@intFromEnum(inst)].pl_op;
1031910322 const dimension = pl_op.payload;
10320 if (dimension >= 3) return .@"1";
10321
10322 // Fetch the dispatch pointer, which points to this structure:
10323 // https://github.com/RadeonOpenCompute/ROCR-Runtime/blob/adae6c61e10d371f7cbc3d0e94ae2c070cab18a4/src/inc/hsa.h#L2913
10324 const dispatch_ptr =
10325 try self.wip.callIntrinsic(.normal, .none, .@"amdgcn.dispatch.ptr", &.{}, &.{}, "");
1032610323
10327 // Load the work_group_* member from the struct as u16.
10328 // Just treat the dispatch pointer as an array of u16 to keep things simple.
10329 const workgroup_size_ptr = try self.wip.gep(.inbounds, .i16, dispatch_ptr, &.{
10330 try o.builder.intValue(try o.lowerType(Type.usize), 2 + dimension),
10331 }, "");
10332 const workgroup_size_alignment = comptime Builder.Alignment.fromByteUnits(2);
10333 return self.wip.load(.normal, .i16, workgroup_size_ptr, workgroup_size_alignment, "");
10324 switch (target.cpu.arch) {
10325 .amdgcn => {
10326 if (dimension >= 3) return .@"1";
10327
10328 // Fetch the dispatch pointer, which points to this structure:
10329 // https://github.com/RadeonOpenCompute/ROCR-Runtime/blob/adae6c61e10d371f7cbc3d0e94ae2c070cab18a4/src/inc/hsa.h#L2913
10330 const dispatch_ptr =
10331 try self.wip.callIntrinsic(.normal, .none, .@"amdgcn.dispatch.ptr", &.{}, &.{}, "");
10332
10333 // Load the work_group_* member from the struct as u16.
10334 // Just treat the dispatch pointer as an array of u16 to keep things simple.
10335 const workgroup_size_ptr = try self.wip.gep(.inbounds, .i16, dispatch_ptr, &.{
10336 try o.builder.intValue(try o.lowerType(Type.usize), 2 + dimension),
10337 }, "");
10338 const workgroup_size_alignment = comptime Builder.Alignment.fromByteUnits(2);
10339 return self.wip.load(.normal, .i16, workgroup_size_ptr, workgroup_size_alignment, "");
10340 },
10341 .nvptx, .nvptx64 => {
10342 return self.workIntrinsic(dimension, 1, "nvvm.read.ptx.sreg.ntid");
10343 },
10344 else => unreachable,
10345 }
1033410346 }
1033510347
1033610348 fn airWorkGroupId(self: *FuncGen, inst: Air.Inst.Index) !Builder.Value {
1033710349 const o = self.ng.object;
1033810350 const target = o.pt.zcu.getTarget();
10339 assert(target.cpu.arch == .amdgcn); // TODO is to port this function to other GPU architectures
1034010351
1034110352 const pl_op = self.air.instructions.items(.data)[@intFromEnum(inst)].pl_op;
1034210353 const dimension = pl_op.payload;
10343 return self.amdgcnWorkIntrinsic(dimension, 0, "amdgcn.workgroup.id");
10354
10355 return switch (target.cpu.arch) {
10356 .amdgcn => self.workIntrinsic(dimension, 0, "amdgcn.workgroup.id"),
10357 .nvptx, .nvptx64 => self.workIntrinsic(dimension, 0, "nvvm.read.ptx.sreg.ctaid"),
10358 else => unreachable,
10359 };
1034410360 }
1034510361
1034610362 fn getErrorNameTable(self: *FuncGen) Allocator.Error!Builder.Variable.Index {
src/codegen/llvm/Builder.zig+77
......@@ -2729,6 +2729,17 @@ pub const Intrinsic = enum {
27292729 @"amdgcn.workgroup.id.z",
27302730 @"amdgcn.dispatch.ptr",
27312731
2732 // NVPTX
2733 @"nvvm.read.ptx.sreg.tid.x",
2734 @"nvvm.read.ptx.sreg.tid.y",
2735 @"nvvm.read.ptx.sreg.tid.z",
2736 @"nvvm.read.ptx.sreg.ntid.x",
2737 @"nvvm.read.ptx.sreg.ntid.y",
2738 @"nvvm.read.ptx.sreg.ntid.z",
2739 @"nvvm.read.ptx.sreg.ctaid.x",
2740 @"nvvm.read.ptx.sreg.ctaid.y",
2741 @"nvvm.read.ptx.sreg.ctaid.z",
2742
27322743 // WebAssembly
27332744 @"wasm.memory.size",
27342745 @"wasm.memory.grow",
......@@ -3886,6 +3897,72 @@ pub const Intrinsic = enum {
38863897 .attrs = &.{ .nocallback, .nofree, .nosync, .nounwind, .speculatable, .willreturn, .{ .memory = Attribute.Memory.all(.none) } },
38873898 },
38883899
3900 .@"nvvm.read.ptx.sreg.tid.x" = .{
3901 .ret_len = 1,
3902 .params = &.{
3903 .{ .kind = .{ .type = .i32 } },
3904 },
3905 .attrs = &.{ .nounwind, .readnone },
3906 },
3907 .@"nvvm.read.ptx.sreg.tid.y" = .{
3908 .ret_len = 1,
3909 .params = &.{
3910 .{ .kind = .{ .type = .i32 } },
3911 },
3912 .attrs = &.{ .nounwind, .readnone },
3913 },
3914 .@"nvvm.read.ptx.sreg.tid.z" = .{
3915 .ret_len = 1,
3916 .params = &.{
3917 .{ .kind = .{ .type = .i32 } },
3918 },
3919 .attrs = &.{ .nounwind, .readnone },
3920 },
3921
3922 .@"nvvm.read.ptx.sreg.ntid.x" = .{
3923 .ret_len = 1,
3924 .params = &.{
3925 .{ .kind = .{ .type = .i32 } },
3926 },
3927 .attrs = &.{ .nounwind, .readnone },
3928 },
3929 .@"nvvm.read.ptx.sreg.ntid.y" = .{
3930 .ret_len = 1,
3931 .params = &.{
3932 .{ .kind = .{ .type = .i32 } },
3933 },
3934 .attrs = &.{ .nounwind, .readnone },
3935 },
3936 .@"nvvm.read.ptx.sreg.ntid.z" = .{
3937 .ret_len = 1,
3938 .params = &.{
3939 .{ .kind = .{ .type = .i32 } },
3940 },
3941 .attrs = &.{ .nounwind, .readnone },
3942 },
3943
3944 .@"nvvm.read.ptx.sreg.ctaid.x" = .{
3945 .ret_len = 1,
3946 .params = &.{
3947 .{ .kind = .{ .type = .i32 } },
3948 },
3949 .attrs = &.{ .nounwind, .readnone },
3950 },
3951 .@"nvvm.read.ptx.sreg.ctaid.y" = .{
3952 .ret_len = 1,
3953 .params = &.{
3954 .{ .kind = .{ .type = .i32 } },
3955 },
3956 .attrs = &.{ .nounwind, .readnone },
3957 },
3958 .@"nvvm.read.ptx.sreg.ctaid.z" = .{
3959 .ret_len = 1,
3960 .params = &.{
3961 .{ .kind = .{ .type = .i32 } },
3962 },
3963 .attrs = &.{ .nounwind, .readnone },
3964 },
3965
38893966 .@"wasm.memory.size" = .{
38903967 .ret_len = 1,
38913968 .params = &.{