lib/accy/src/target/nvptx/test.zig

daab053ee43316e1809a84551d573ddd1e5bf3d2

   1 const std = @import("std");
   2 const gpu = @import("gpu");
   3 const choir = @import("choir");
   4 const choir_abi = @import("choir_abi");
   5 const accy = @import("../../root.zig");
   6 const pretty = @import("pretty");
   7 const kernel = accy.kernel;
   8 const Allocator = std.mem.Allocator;
   9 const emitPtx = choir.backends.gpu.nvptx.ptx.emitPtx;
  10 
  11 const scalar_kernel_name = "accy_cuda_scalar_words";
  12 const scalar_output_count = 8;
  13 
  14 fn scalarParameterGraph(allocator: std.mem.Allocator) !kernel.Graph {
  15     const params = [_]kernel.Param{
  16         kernel.buffer(.u64, scalar_output_count),
  17         kernel.scalar(.i64),
  18         kernel.scalar(.u32),
  19         kernel.scalar(.u64),
  20         kernel.scalar(.i32),
  21     };
  22     var builder = try kernel.Builder.init(
  23         allocator,
  24         kernel.Builder.Limits.testing,
  25         scalar_kernel_name,
  26         &params,
  27     );
  28     errdefer builder.deinit();
  29     try builder.bind(try builder.axis("i", 1), .thread_x);
  30     const signed = builder.argument(1);
  31     const unsigned = builder.argument(3);
  32     const bits = try builder.bitcast(signed, .u64);
  33     const narrow_signed = try builder.cast(builder.argument(4), .i64);
  34     const signed_sum = try builder.add(signed, narrow_signed);
  35     const signed_less = try builder.compare(.lt, signed, try builder.bitcast(unsigned, .i64));
  36     const unsigned_less = try builder.compare(.lt, unsigned, bits);
  37     const one = try builder.constantInt(.u64, 1);
  38     const zero = try builder.constantInt(.u64, 0);
  39     const values = [_]kernel.Value{
  40         bits,
  41         unsigned,
  42         try builder.mul(bits, unsigned),
  43         try builder.bitcast(signed_sum, .u64),
  44         try builder.select(signed_less, one, zero),
  45         try builder.select(unsigned_less, one, zero),
  46         try builder.cast(builder.argument(2), .u64),
  47         try builder.bitcast(narrow_signed, .u64),
  48     };
  49     const base = try builder.globalId(.x);
  50     for (values, 0..) |value, offset| {
  51         const index = try builder.add(base, try builder.constantIndex(@intCast(offset)));
  52         try builder.store(value, builder.argument(0), index);
  53     }
  54     try builder.return_();
  55     return builder.finish();
  56 }
  57 
  58 test "CUDA 64-bit scalar parameters compile from a sealed source graph" {
  59     var graph = try scalarParameterGraph(std.testing.allocator);
  60     defer graph.deinit();
  61     const before = try graph.bodyFingerprint(std.testing.allocator);
  62     try std.testing.expect(graph.kernelModule().getContext().isFrozen());
  63     const payload = (try accy.target.compileKernelForArtifactFormat(
  64         std.testing.allocator,
  65         .cuda_ptx,
  66         scalar_kernel_name,
  67         graph.kernelModule(),
  68         .{},
  69     )).payload;
  70     const text = payload.text;
  71     defer std.testing.allocator.free(@constCast(text));
  72     try std.testing.expectEqual(before, try graph.bodyFingerprint(std.testing.allocator));
  73     try std.testing.expect(std.mem.indexOf(u8, text, ".param .u64 param1,") != null);
  74     try std.testing.expect(std.mem.indexOf(u8, text, ".param .u32 param2,") != null);
  75     try std.testing.expect(std.mem.indexOf(u8, text, ".param .u64 param3,") != null);
  76     try std.testing.expect(std.mem.indexOf(u8, text, "setp.lt.s64") != null);
  77     try std.testing.expect(std.mem.indexOf(u8, text, "setp.lt.u64") != null);
  78 }
  79 
  80 const ScalarArguments = struct {
  81     signed: i64,
  82     narrow_unsigned: u32,
  83     unsigned: u64,
  84     narrow_signed: i32,
  85 
  86     fn launch(self: ScalarArguments) [4]choir_abi.ScalarArgument {
  87         return .{
  88             .{ .i64 = self.signed },   .{ .u32 = self.narrow_unsigned },
  89             .{ .u64 = self.unsigned }, .{ .i32 = self.narrow_signed },
  90         };
  91     }
  92 
  93     fn reference(self: ScalarArguments) [scalar_output_count]u64 {
  94         const bits: u64 = @bitCast(self.signed);
  95         const product: u128 = @as(u128, bits) * self.unsigned;
  96         const sum: i128 = @as(i128, self.signed) + self.narrow_signed;
  97         return .{
  98             bits,
  99             self.unsigned,
 100             @truncate(product),
 101             @truncate(@as(u128, @bitCast(sum))),
 102             @intFromBool(self.signed < @as(i64, @bitCast(self.unsigned))),
 103             @intFromBool(self.unsigned < bits),
 104             self.narrow_unsigned,
 105             @bitCast(@as(i64, self.narrow_signed)),
 106         };
 107     }
 108 };
 109 
 110 test "CUDA 64-bit scalar parameters preserve every bit on the device" {
 111     var state = try cudaDevice();
 112     defer state.deinit();
 113     const handle = state.handle();
 114     var graph = try scalarParameterGraph(std.testing.allocator);
 115     defer graph.deinit();
 116     var artifact = try kernel.createKernelArtifact(std.testing.allocator, handle, &graph, .{
 117         .artifact_format = .cuda_ptx,
 118         .authored_kernel_diagnostic_id = "accy/cuda/scalar-parameters",
 119     });
 120     defer artifact.deinit();
 121     try std.testing.expectEqual(@as(u32, 5), artifact.argument_count);
 122     try std.testing.expectEqual(@as(u32, 4), artifact.scalar_argument_count);
 123     const loaded = try handle.loadArtifact(&artifact);
 124     defer handle.destroyObject(loaded.id);
 125     const output = try handle.allocateBuffer(.{
 126         .byte_size = scalar_output_count * @sizeOf(u64),
 127         .alignment = 256,
 128         .dtype = .u64,
 129         .element_count = scalar_output_count,
 130     });
 131     defer handle.destroyObject(output.id);
 132     const words = [_]u64{
 133         0,                     1,                     0xffff_ffff,           0x1_0000_0000,
 134         0x7fff_ffff_ffff_ffff, 0x8000_0000_0000_0000, 0xffff_ffff_ffff_ffff,
 135     };
 136     const narrow_words = [_]u32{ 0, 1, 0xffff_ffff, 0x8000_0000, 0x7fff_ffff, 0x1234_5678 };
 137     for (words, 0..) |signed, i| {
 138         for (words, 0..) |unsigned, j| {
 139             const args = ScalarArguments{
 140                 .signed = @bitCast(signed),
 141                 .unsigned = unsigned,
 142                 .narrow_unsigned = narrow_words[(i + j) % narrow_words.len],
 143                 .narrow_signed = @bitCast(narrow_words[(i + j + 1) % narrow_words.len]),
 144             };
 145             try checkScalarLaunch(handle, &artifact, loaded, output, args);
 146         }
 147     }
 148 }
 149 
 150 fn cudaDevice() !gpu.cuda.State {
 151     if (!@import("builtin").link_libc) try accy.validation.gating.skipIfBuildFlagDisabled(.cuda);
 152     return gpu.cuda.State.initDevice(std.testing.allocator, 0) catch |err| {
 153         if (@import("builtin").link_libc or err != error.RuntimeUnavailable) return err;
 154         return accy.validation.gating.skip(.cuda, .cuda_runtime_unavailable);
 155     };
 156 }
 157 
 158 fn narrowGraph(allocator: std.mem.Allocator) !kernel.Graph {
 159     var builder = try kernel.Builder.init(
 160         allocator,
 161         kernel.Builder.Limits.testing,
 162         "accy_cuda_narrow",
 163         &.{ kernel.buffer(.u64, 4), kernel.scalar(.u64) },
 164     );
 165     errdefer builder.deinit();
 166     try builder.bind(try builder.axis("i", 1), .thread_x);
 167     const kinds = [_]choir_abi.DType{ .u8, .i8, .u16, .i16 };
 168     for (kinds, 0..) |kind, index| {
 169         const input = try builder.cast(builder.argument(1), kind);
 170         const sum = try builder.add(input, try builder.constantInt(kind, 1));
 171         const wide = try builder.cast(sum, .i64);
 172         const bits = try builder.bitcast(wide, .u64);
 173         try builder.store(bits, builder.argument(0), try builder.constantIndex(@intCast(index)));
 174     }
 175     try builder.return_();
 176     return builder.finish();
 177 }
 178 
 179 test "CUDA narrow arithmetic wraps and sign extends before widening" {
 180     var state = try cudaDevice();
 181     defer state.deinit();
 182     const handle = state.handle();
 183     var graph = try narrowGraph(std.testing.allocator);
 184     defer graph.deinit();
 185     var artifact = try kernel.createKernelArtifact(std.testing.allocator, handle, &graph, .{
 186         .artifact_format = .cuda_ptx,
 187         .authored_kernel_diagnostic_id = "accy/cuda/narrow",
 188     });
 189     defer artifact.deinit();
 190     const loaded = try handle.loadArtifact(&artifact);
 191     defer handle.destroyObject(loaded.id);
 192     const output = try handle.allocateBuffer(.{
 193         .byte_size = 32,
 194         .alignment = 256,
 195         .dtype = .u64,
 196         .element_count = 4,
 197     });
 198     defer handle.destroyObject(output.id);
 199     const words = [_]u64{
 200         0, 126, 127, 128, 254, 255, 256, 32767, 32768, 65535, 65536, 0xffff_ffff_ffff_ffff,
 201     };
 202     for (words) |word| {
 203         try handle.launch(.{
 204             .artifact = &artifact,
 205             .loaded_artifact = loaded,
 206             .buffers = &.{.{
 207                 .handle = output,
 208                 .access = .read_write,
 209                 .ownership = output.ownership,
 210                 .byte_size = output.byte_size,
 211             }},
 212             .scalar_arguments = &.{.{ .u64 = word }},
 213             .geometry = .{ .grid = .{ 1, 1, 1 }, .threadgroup = .{ 1, 1, 1 } },
 214         });
 215         try handle.synchronize(.{ .scope = .device });
 216         var observed: [4]u64 = undefined;
 217         try handle.readBuffer(.{ .handle = output, .bytes = std.mem.asBytes(&observed) });
 218         const byte: i64 = @intCast((@as(u128, word) + 1) % 256);
 219         const half: i64 = @intCast((@as(u128, word) + 1) % 65536);
 220         const expected = [_]u64{
 221             @intCast(byte),
 222             @bitCast(if (byte < 128) byte else byte - 256),
 223             @intCast(half),
 224             @bitCast(if (half < 32768) half else half - 65536),
 225         };
 226         if (!std.mem.eql(u64, &expected, &observed)) {
 227             var arena = std.heap.ArenaAllocator.init(std.testing.allocator);
 228             defer arena.deinit();
 229             var report = try pretty.diagnostic.Report.init(
 230                 arena.allocator(),
 231                 "CUDA narrow mismatch",
 232             );
 233             defer report.deinit();
 234             try report.field("input", "{x}", .{word});
 235             try report.field("expected", "{any}", .{expected});
 236             try report.field("observed", "{any}", .{observed});
 237             pretty.diagnostic.writeStderr(&report, .{ .width = 100 });
 238             return error.NarrowMismatch;
 239         }
 240     }
 241 }
 242 
 243 fn checkScalarLaunch(
 244     handle: gpu.BackendHandle,
 245     artifact: *const gpu.KernelArtifact,
 246     loaded: gpu.LoadedArtifact,
 247     output: gpu.BufferHandle,
 248     args: ScalarArguments,
 249 ) !void {
 250     std.debug.assert(output.byte_size == scalar_output_count * @sizeOf(u64));
 251     std.debug.assert(artifact.format == .cuda_ptx);
 252     const scalar_args = args.launch();
 253     const binding = gpu.BufferBinding{
 254         .handle = output,
 255         .access = .read_write,
 256         .ownership = output.ownership,
 257         .byte_size = output.byte_size,
 258     };
 259     var observed: [scalar_output_count]u64 = @splat(0xa5a5_a5a5_a5a5_a5a5);
 260     try handle.writeBuffer(.{ .handle = output, .bytes = std.mem.sliceAsBytes(&observed) });
 261     try handle.launch(.{
 262         .artifact = artifact,
 263         .loaded_artifact = loaded,
 264         .buffers = &.{binding},
 265         .scalar_arguments = &scalar_args,
 266         .geometry = .{ .grid = .{ 1, 1, 1 }, .threadgroup = .{ 1, 1, 1 } },
 267     });
 268     try handle.synchronize(.{ .scope = .device });
 269     try handle.readBuffer(.{ .handle = output, .bytes = std.mem.sliceAsBytes(&observed) });
 270     const expected = args.reference();
 271     if (!std.mem.eql(u64, &expected, &observed)) {
 272         try reportScalarMismatch(artifact, args, expected, observed);
 273         return error.ScalarParameterMismatch;
 274     }
 275 }
 276 
 277 fn reportScalarMismatch(
 278     artifact: *const gpu.KernelArtifact,
 279     args: ScalarArguments,
 280     expected: [scalar_output_count]u64,
 281     observed: [scalar_output_count]u64,
 282 ) !void {
 283     var arena = std.heap.ArenaAllocator.init(std.testing.allocator);
 284     defer arena.deinit();
 285     var report = try pretty.diagnostic.Report.init(
 286         arena.allocator(),
 287         "CUDA scalar parameter mismatch",
 288     );
 289     defer report.deinit();
 290     try report.field("backend", "CUDA device 0", .{});
 291     try report.field("kernel", "{s}", .{scalar_kernel_name});
 292     try report.field("artifact", "{x}", .{std.hash.Wyhash.hash(0, artifact.payload.text)});
 293     try report.field("arguments", "{any}", .{args});
 294     try report.field("expected", "{any}", .{expected});
 295     try report.field("observed", "{any}", .{observed});
 296     pretty.diagnostic.writeStderr(&report, .{ .width = 100 });
 297 }
 298 
 299 const loop_kernel_name = "accy_cuda_integer_loop";
 300 const loop_output_count = 4;
 301 const loop_first: u64 = 0x1234_5678_9abc_def0;
 302 const loop_second: u64 = 0x6edc_ba98_7654_3210;
 303 
 304 const Counter = enum {
 305     i32,
 306     index,
 307     u32,
 308     i64,
 309     u64,
 310 
 311     fn parameter(self: Counter) kernel.Param {
 312         return switch (self) {
 313             .i32 => kernel.scalar(.i32),
 314             .index, .u32 => kernel.scalar(.u32),
 315             .i64 => kernel.scalar(.i64),
 316             .u64 => kernel.scalar(.u64),
 317         };
 318     }
 319 
 320     fn argument(self: Counter, value: i128) choir_abi.ScalarArgument {
 321         return switch (self) {
 322             .i32 => .{ .i32 = @intCast(value) },
 323             .index, .u32 => .{ .u32 = @intCast(value) },
 324             .i64 => .{ .i64 = @intCast(value) },
 325             .u64 => .{ .u64 = @intCast(value) },
 326         };
 327     }
 328 
 329     fn starts(self: Counter) []const i128 {
 330         return switch (self) {
 331             .i32 => &.{ -1, std.math.minInt(i32), -8, -4, 0, 1, std.math.maxInt(i32) - 16 },
 332             .index, .u32 => &.{ 0, 1, 0x7fff_fffc, 0x8000_0000, std.math.maxInt(u32) - 16 },
 333             .i64 => &.{
 334                 -1,          std.math.minInt(i64), -8,                        -4, 0, 1,
 335                 0xffff_fffc, 0x1_0000_0000,        std.math.maxInt(i64) - 16,
 336             },
 337             .u64 => &.{
 338                 0,
 339                 1,
 340                 0xffff_fffc,
 341                 0x1_0000_0000,
 342                 0x7fff_ffff_ffff_fffc,
 343                 0x8000_0000_0000_0000,
 344                 std.math.maxInt(u64) - 16,
 345             },
 346         };
 347     }
 348 };
 349 
 350 const LoopCase = struct {
 351     counter: Counter,
 352     lower: i128,
 353     upper: i128,
 354     step: u8,
 355 
 356     fn arguments(self: LoopCase) [3]choir_abi.ScalarArgument {
 357         std.debug.assert(self.step > 0);
 358         return .{
 359             self.counter.argument(self.lower), self.counter.argument(self.upper),
 360             self.counter.argument(self.step),
 361         };
 362     }
 363 
 364     fn reference(self: LoopCase) [loop_output_count]u64 {
 365         std.debug.assert(self.step > 0);
 366         const distance: i128 = @max(0, self.upper - self.lower);
 367         const count: i128 = @divTrunc(distance + self.step - 1, self.step);
 368         std.debug.assert(count <= 5);
 369         const triangle: i128 = @divTrunc(count * (count - 1), 2);
 370         const sum: i128 = 3 * count * self.lower + 3 * self.step * triangle + 3 * count;
 371         return .{
 372             @truncate(@as(u128, @bitCast(sum))),
 373             if (@mod(count, 2) == 0) loop_first else loop_second,
 374             if (@mod(count, 2) == 0) loop_second else loop_first,
 375             @intCast(count),
 376         };
 377     }
 378 };
 379 
 380 fn loopGraph(allocator: std.mem.Allocator, counter: Counter) !kernel.Graph {
 381     const parameter = counter.parameter();
 382     var builder = try kernel.Builder.init(
 383         allocator,
 384         kernel.Builder.Limits.testing,
 385         loop_kernel_name,
 386         &.{ kernel.buffer(.u64, loop_output_count), parameter, parameter, parameter },
 387     );
 388     errdefer builder.deinit();
 389     try builder.bind(try builder.axis("i", 1), .thread_x);
 390     var bounds: [3]kernel.Value = undefined;
 391     for (&bounds, 0..) |*bound, index| {
 392         const argument = builder.argument(index + 1);
 393         bound.* = if (counter == .index) try builder.castIndex(argument) else argument;
 394     }
 395     const zero = try builder.constantInt(.u64, 0);
 396     const first = try builder.constantInt(.u64, @bitCast(loop_first));
 397     const second = try builder.constantInt(.u64, @bitCast(loop_second));
 398     const word = zero.valueType();
 399     var scope = try builder.forScope(
 400         bounds[0],
 401         bounds[1],
 402         bounds[2],
 403         &.{ zero, first, second, zero },
 404         &.{ word, word, word, word },
 405     );
 406     errdefer scope.abort();
 407     const nested = try nestedLoopValue(&builder, scope.inductionVar());
 408     const sum = try builder.add(scope.iterArg(0).?, nested);
 409     const count = try builder.add(scope.iterArg(3).?, try builder.constantInt(.u64, 1));
 410     try scope.leave(&.{ sum, scope.iterArg(2).?, scope.iterArg(1).?, count });
 411     for (0..loop_output_count) |index| {
 412         try builder.store(
 413             scope.result(index).?,
 414             builder.argument(0),
 415             try builder.constantIndex(@intCast(index)),
 416         );
 417     }
 418     try builder.return_();
 419     return builder.finish();
 420 }
 421 
 422 fn nestedLoopValue(builder: *kernel.Builder, outer: kernel.Value) !kernel.Value {
 423     const value = try builder.cast(outer, .u64);
 424     const zero = try builder.constantInt(.u64, 0);
 425     var scope = try builder.forScope(
 426         try builder.constantIndex(0),
 427         try builder.constantIndex(3),
 428         try builder.constantIndex(1),
 429         &.{zero},
 430         &.{zero.valueType()},
 431     );
 432     errdefer scope.abort();
 433     const inner = try builder.cast(scope.inductionVar(), .u64);
 434     const sum = try builder.add(try builder.add(scope.iterArg(0).?, value), inner);
 435     try scope.leave(&.{sum});
 436     return scope.result(0).?;
 437 }
 438 
 439 test "CUDA structured loops preserve source graphs and counter widths" {
 440     for (std.enums.values(Counter)) |counter| {
 441         var graph = try loopGraph(std.testing.allocator, counter);
 442         defer graph.deinit();
 443         const before = try graph.bodyFingerprint(std.testing.allocator);
 444         try std.testing.expect(graph.kernelModule().getContext().isFrozen());
 445         const payload = (try accy.target.compileKernelForArtifactFormat(
 446             std.testing.allocator,
 447             .cuda_ptx,
 448             loop_kernel_name,
 449             graph.kernelModule(),
 450             .{},
 451         )).payload;
 452         defer std.testing.allocator.free(@constCast(payload.text));
 453         try std.testing.expectEqual(before, try graph.bodyFingerprint(std.testing.allocator));
 454         const comparison = switch (counter) {
 455             .i32 => "setp.ge.s32",
 456             .index, .u32 => "setp.ge.u32",
 457             .i64 => "setp.ge.s64",
 458             .u64 => "setp.ge.u64",
 459         };
 460         try std.testing.expect(std.mem.indexOf(u8, payload.text, comparison) != null);
 461     }
 462 }
 463 
 464 test "CUDA structured loops match independent sums and simultaneous swaps on the device" {
 465     var state = try cudaDevice();
 466     defer state.deinit();
 467     const handle = state.handle();
 468     for (std.enums.values(Counter)) |counter| try qualifyLoops(handle, counter);
 469 }
 470 
 471 fn qualifyLoops(handle: gpu.BackendHandle, counter: Counter) !void {
 472     var graph = try loopGraph(std.testing.allocator, counter);
 473     defer graph.deinit();
 474     var artifact = try kernel.createKernelArtifact(std.testing.allocator, handle, &graph, .{
 475         .artifact_format = .cuda_ptx,
 476         .authored_kernel_diagnostic_id = "accy/cuda/integer-loops",
 477     });
 478     defer artifact.deinit();
 479     try std.testing.expectEqual(@as(u32, 4), artifact.argument_count);
 480     try std.testing.expectEqual(@as(u32, 3), artifact.scalar_argument_count);
 481     const loaded = try handle.loadArtifact(&artifact);
 482     defer handle.destroyObject(loaded.id);
 483     const output = try handle.allocateBuffer(.{
 484         .byte_size = loop_output_count * @sizeOf(u64),
 485         .alignment = 256,
 486         .dtype = .u64,
 487         .element_count = loop_output_count,
 488     });
 489     defer handle.destroyObject(output.id);
 490     for (counter.starts()) |start| {
 491         for (0..5) |trips| {
 492             for (1..4) |step| {
 493                 try checkLoopLaunch(handle, &artifact, loaded, output, .{
 494                     .counter = counter,
 495                     .lower = start,
 496                     .upper = start + @as(i128, trips * step - (trips % 2) * (step - 1)),
 497                     .step = @intCast(step),
 498                 });
 499             }
 500         }
 501     }
 502     try checkLoopLaunch(handle, &artifact, loaded, output, .{
 503         .counter = counter,
 504         .lower = 3,
 505         .upper = 0,
 506         .step = 1,
 507     });
 508 }
 509 
 510 fn checkLoopLaunch(
 511     handle: gpu.BackendHandle,
 512     artifact: *const gpu.KernelArtifact,
 513     loaded: gpu.LoadedArtifact,
 514     output: gpu.BufferHandle,
 515     case: LoopCase,
 516 ) !void {
 517     std.debug.assert(output.byte_size == loop_output_count * @sizeOf(u64));
 518     std.debug.assert(artifact.format == .cuda_ptx);
 519     const arguments = case.arguments();
 520     const binding = gpu.BufferBinding{
 521         .handle = output,
 522         .access = .read_write,
 523         .ownership = output.ownership,
 524         .byte_size = output.byte_size,
 525     };
 526     var observed: [loop_output_count]u64 = @splat(0xa5a5_a5a5_a5a5_a5a5);
 527     try handle.writeBuffer(.{ .handle = output, .bytes = std.mem.sliceAsBytes(&observed) });
 528     try handle.launch(.{
 529         .artifact = artifact,
 530         .loaded_artifact = loaded,
 531         .buffers = &.{binding},
 532         .scalar_arguments = &arguments,
 533         .geometry = .{ .grid = .{ 1, 1, 1 }, .threadgroup = .{ 1, 1, 1 } },
 534     });
 535     try handle.synchronize(.{ .scope = .device });
 536     try handle.readBuffer(.{ .handle = output, .bytes = std.mem.sliceAsBytes(&observed) });
 537     const expected = case.reference();
 538     if (!std.mem.eql(u64, &expected, &observed)) {
 539         try reportLoopMismatch(artifact, case, expected, observed);
 540         return error.IntegerLoopMismatch;
 541     }
 542 }
 543 
 544 fn reportLoopMismatch(
 545     artifact: *const gpu.KernelArtifact,
 546     case: LoopCase,
 547     expected: [loop_output_count]u64,
 548     observed: [loop_output_count]u64,
 549 ) !void {
 550     var arena = std.heap.ArenaAllocator.init(std.testing.allocator);
 551     defer arena.deinit();
 552     var report = try pretty.diagnostic.Report.init(arena.allocator(), "CUDA integer loop mismatch");
 553     defer report.deinit();
 554     try report.field("backend", "CUDA device 0", .{});
 555     try report.field("kernel", "{s}", .{loop_kernel_name});
 556     try report.field("artifact", "{x}", .{std.hash.Wyhash.hash(0, artifact.payload.text)});
 557     try report.field("case", "{any}", .{case});
 558     try report.field("expected", "{any}", .{expected});
 559     try report.field("observed", "{any}", .{observed});
 560     pretty.diagnostic.writeStderr(&report, .{ .width = 100 });
 561 }
 562 
 563 fn emitLoweredPtx(
 564     allocator: Allocator,
 565     entry_name: []const u8,
 566     program: *accy.kernel.Graph,
 567 ) gpu.BackendError![]u8 {
 568     const ctx = program.kernelModule().getContext();
 569     const handler = ctx.registerDiagnosticHandler(.{ .handle = reportTestDiagnostic }) catch
 570         return error.OutOfMemory;
 571     defer ctx.eraseDiagnosticHandler(handler);
 572     const payload = (try accy.target.compileKernelForArtifactFormat(
 573         allocator,
 574         .cuda_ptx,
 575         entry_name,
 576         program.kernelModule(),
 577         .{},
 578     )).payload;
 579     return switch (payload) {
 580         .text => |text| @constCast(text),
 581         else => error.InvalidArtifact,
 582     };
 583 }
 584 
 585 fn reportTestDiagnostic(
 586     _: ?*anyopaque,
 587     diagnostic: *const choir.diagnostics.Diagnostic,
 588 ) !choir.diagnostics.HandlerResult {
 589     var arena = std.heap.ArenaAllocator.init(std.testing.allocator);
 590     defer arena.deinit();
 591     var report = try pretty.diagnostic.Report.init(arena.allocator(), diagnostic.message);
 592     defer report.deinit();
 593     if (diagnostic.error_name) |name| try report.field("error", "{s}", .{name});
 594     for (diagnostic.metadata) |entry| try report.field(entry.name, "{s}", .{entry.value});
 595     pretty.diagnostic.writeStderr(&report, .{ .width = 100 });
 596     return .consumed;
 597 }
 598 
 599 test "cuda choir ptx emitter requires nvptx-lowered gpu ops" {
 600     const testing = std.testing;
 601     const builder_mod = accy.kernel;
 602     const params = [_]builder_mod.Param{
 603         builder_mod.dynamicBuffer(.i32),
 604         builder_mod.dynamicBuffer(.i32),
 605     };
 606 
 607     var builder = try builder_mod.Builder.init(testing.allocator, builder_mod.Builder.Limits.testing, "accy_choir_test_raw_gpu", &params);
 608     errdefer builder.deinit();
 609     const out = builder.argument(0);
 610     const input = builder.argument(1);
 611     const index = try builder.globalId(.x);
 612     const value = try builder.load(input, index);
 613     try builder.store(value, out, index);
 614     try builder.return_();
 615 
 616     var program = try builder.finish();
 617     defer program.deinit();
 618 
 619     try testing.expectError(
 620         error.UnsupportedOperation,
 621         emitPtx(testing.allocator, "accy_choir_test_raw_gpu", program.kernelModule()),
 622     );
 623 }
 624 
 625 test "cuda choir ptx emitter lowers simple elementwise kernel" {
 626     const testing = std.testing;
 627     const builder_mod = accy.kernel;
 628     const params = [_]builder_mod.Param{
 629         builder_mod.dynamicBuffer(.f32),
 630         builder_mod.dynamicBuffer(.f32),
 631         builder_mod.dynamicBuffer(.f32),
 632         builder_mod.dynamicBuffer(.i32),
 633     };
 634 
 635     var builder = try builder_mod.Builder.init(testing.allocator, builder_mod.Builder.Limits.testing, "accy_choir_test_add", &params);
 636     errdefer builder.deinit();
 637     const out = builder.argument(0);
 638     const lhs = builder.argument(1);
 639     const rhs = builder.argument(2);
 640     const count_buffer = builder.argument(3);
 641     const index = try builder.globalId(.x);
 642     const zero = try builder.constantIndex(0);
 643     const count = try builder.load(count_buffer, zero);
 644     const index_i32 = try builder.cast(index, .i32);
 645     const in_bounds = try builder.compare(.lt, index_i32, count);
 646     var if_op = try builder.if_(in_bounds, &.{});
 647     const previous = builder.insertionBlock();
 648     builder.setInsertionBlock(if_op.thenBlock());
 649     const left = try builder.load(lhs, index);
 650     const right = try builder.load(rhs, index);
 651     const sum = try builder.add(left, right);
 652     try builder.store(sum, out, index);
 653     try builder.yield_(&.{});
 654     builder.setInsertionBlock(if_op.elseBlock().?);
 655     try builder.yield_(&.{});
 656     builder.setInsertionBlock(previous);
 657     try builder.return_();
 658 
 659     var program = try builder.finish();
 660     defer program.deinit();
 661     const text = try emitLoweredPtx(testing.allocator, "accy_choir_test_add", &program);
 662     defer testing.allocator.free(text);
 663 
 664     try testing.expect(std.mem.indexOf(u8, text, ".visible .entry accy_choir_test_add") != null);
 665     try testing.expect(std.mem.indexOf(u8, text, "ld.global.f32") != null);
 666     try testing.expect(std.mem.indexOf(u8, text, "add.f32") != null);
 667     try testing.expect(std.mem.indexOf(u8, text, "st.global.f32") != null);
 668     try testing.expect(std.mem.endsWith(u8, text, "}\n"));
 669     try testing.expectEqual(std.mem.count(u8, text, "{"), std.mem.count(u8, text, "}"));
 670 }
 671 
 672 test "cuda choir ptx emitter lowers while loops with carried exits" {
 673     const testing = std.testing;
 674     const builder_mod = accy.kernel;
 675     const params = [_]builder_mod.Param{
 676         builder_mod.dynamicBuffer(.i32),
 677         builder_mod.dynamicBuffer(.i32),
 678     };
 679 
 680     var builder = try builder_mod.Builder.init(testing.allocator, builder_mod.Builder.Limits.testing, "accy_choir_test_while", &params);
 681     errdefer builder.deinit();
 682     const src = builder.argument(0);
 683     const dst = builder.argument(1);
 684     const index = try builder.globalId(.x);
 685     const start = try builder.load(src, index);
 686     const zero = try builder.constantInt(.i32, 0);
 687     const one = try builder.constantInt(.i32, 1);
 688 
 689     var scope = try builder.whileScope(&.{ start, zero }, &.{ start.valueType(), zero.valueType() });
 690     errdefer scope.abort();
 691     const remaining = scope.beforeArg(0).?;
 692     const total = scope.beforeArg(1).?;
 693     const proceed = try builder.compare(.gt, remaining, zero);
 694     try scope.condition(proceed, &.{ remaining, total });
 695     const after_remaining = scope.afterArg(0).?;
 696     const after_total = scope.afterArg(1).?;
 697     const next_total = try builder.add(after_total, after_remaining);
 698     const next_remaining = try builder.sub(after_remaining, one);
 699     try scope.leave(&.{ next_remaining, next_total });
 700 
 701     try builder.store(scope.result(1).?, dst, index);
 702     try builder.return_();
 703 
 704     var program = try builder.finish();
 705     defer program.deinit();
 706     const text = try emitLoweredPtx(testing.allocator, "accy_choir_test_while", &program);
 707     defer testing.allocator.free(text);
 708 
 709     try testing.expect(std.mem.indexOf(u8, text, "LWHILE_HEAD_") != null);
 710     try testing.expect(std.mem.indexOf(u8, text, "LWHILE_DONE_") != null);
 711     try testing.expect(std.mem.indexOf(u8, text, "bra        LWHILE_DONE_") != null);
 712     try testing.expect(std.mem.indexOf(u8, text, "bra                LWHILE_HEAD_") != null);
 713     try testing.expect(std.mem.indexOf(u8, text, "setp.gt.s32") != null);
 714     try testing.expectEqual(std.mem.count(u8, text, "{"), std.mem.count(u8, text, "}"));
 715 }
 716 
 717 test "cuda choir ptx emitter lowers value producing if arms" {
 718     const testing = std.testing;
 719     const builder_mod = accy.kernel;
 720     const params = [_]builder_mod.Param{
 721         builder_mod.dynamicBuffer(.f32),
 722         builder_mod.dynamicBuffer(.f32),
 723     };
 724 
 725     var builder = try builder_mod.Builder.init(testing.allocator, builder_mod.Builder.Limits.testing, "accy_choir_test_if_value", &params);
 726     errdefer builder.deinit();
 727     const dst = builder.argument(0);
 728     const src = builder.argument(1);
 729     const index = try builder.globalId(.x);
 730     const value = try builder.load(src, index);
 731     const zero = try builder.constantFloat(.f32, 0.0);
 732     const positive = try builder.compare(.gt, value, zero);
 733 
 734     var if_op = try builder.if_(positive, &.{value.valueType()});
 735     const previous = builder.insertionBlock();
 736     builder.setInsertionBlock(if_op.thenBlock());
 737     const doubled = try builder.add(value, value);
 738     try builder.yield_(&.{doubled});
 739     builder.setInsertionBlock(if_op.elseBlock().?);
 740     try builder.yield_(&.{zero});
 741     builder.setInsertionBlock(previous);
 742 
 743     try builder.store(if_op.result(0).?, dst, index);
 744     try builder.return_();
 745 
 746     var program = try builder.finish();
 747     defer program.deinit();
 748     const text = try emitLoweredPtx(testing.allocator, "accy_choir_test_if_value", &program);
 749     defer testing.allocator.free(text);
 750 
 751     try testing.expect(std.mem.indexOf(u8, text, "LIF_ELSE_") != null);
 752     try testing.expect(std.mem.indexOf(u8, text, "LIF_DONE_") != null);
 753     try testing.expect(std.mem.indexOf(u8, text, "st.global.f32") != null);
 754     try testing.expectEqual(std.mem.count(u8, text, "{"), std.mem.count(u8, text, "}"));
 755 }
 756 
 757 test "cuda choir ptx emitter lowers bool global load store as bytes" {
 758     const testing = std.testing;
 759     const builder_mod = accy.kernel;
 760     const params = [_]builder_mod.Param{
 761         builder_mod.dynamicBuffer(.i1),
 762         builder_mod.dynamicBuffer(.i1),
 763     };
 764 
 765     var builder = try builder_mod.Builder.init(testing.allocator, builder_mod.Builder.Limits.testing, "accy_choir_test_bool_copy", &params);
 766     errdefer builder.deinit();
 767     const out = builder.argument(0);
 768     const input = builder.argument(1);
 769     const index = try builder.globalId(.x);
 770     const value = try builder.load(input, index);
 771     try builder.store(value, out, index);
 772     try builder.return_();
 773 
 774     var program = try builder.finish();
 775     defer program.deinit();
 776     const text = try emitLoweredPtx(testing.allocator, "accy_choir_test_bool_copy", &program);
 777     defer testing.allocator.free(text);
 778 
 779     try testing.expect(std.mem.indexOf(u8, text, "ld.global.u8") != null);
 780     try testing.expect(std.mem.indexOf(u8, text, "setp.ne.u32") != null);
 781     try testing.expect(std.mem.indexOf(u8, text, "selp.u32") != null);
 782     try testing.expect(std.mem.indexOf(u8, text, "st.global.u8") != null);
 783 }
 784 
 785 test "cuda choir ptx emitter lowers bool select as predicates" {
 786     const testing = std.testing;
 787     const builder_mod = accy.kernel;
 788     const params = [_]builder_mod.Param{
 789         builder_mod.dynamicBuffer(.i1),
 790         builder_mod.dynamicBuffer(.i1),
 791         builder_mod.dynamicBuffer(.i1),
 792         builder_mod.dynamicBuffer(.i1),
 793     };
 794 
 795     var builder = try builder_mod.Builder.init(testing.allocator, builder_mod.Builder.Limits.testing, "accy_choir_test_bool_select", &params);
 796     errdefer builder.deinit();
 797     const out = builder.argument(0);
 798     const mask = builder.argument(1);
 799     const lhs = builder.argument(2);
 800     const rhs = builder.argument(3);
 801     const index = try builder.globalId(.x);
 802     const cond = try builder.load(mask, index);
 803     const left = try builder.load(lhs, index);
 804     const right = try builder.load(rhs, index);
 805     const selected = try builder.select(cond, left, right);
 806     try builder.store(selected, out, index);
 807     try builder.return_();
 808 
 809     var program = try builder.finish();
 810     defer program.deinit();
 811     const text = try emitLoweredPtx(testing.allocator, "accy_choir_test_bool_select", &program);
 812     defer testing.allocator.free(text);
 813 
 814     try testing.expect(std.mem.indexOf(u8, text, "not.pred") != null);
 815     try testing.expect(std.mem.indexOf(u8, text, "and.pred") != null);
 816     try testing.expect(std.mem.indexOf(u8, text, "or.pred") != null);
 817     try testing.expect(std.mem.indexOf(u8, text, "st.global.u8") != null);
 818 }
 819 
 820 test "cuda choir ptx emitter lowers signed integer compare as signed" {
 821     const testing = std.testing;
 822     const builder_mod = accy.kernel;
 823     const params = [_]builder_mod.Param{
 824         builder_mod.dynamicBuffer(.i32),
 825         builder_mod.dynamicBuffer(.i32),
 826     };
 827 
 828     var builder = try builder_mod.Builder.init(testing.allocator, builder_mod.Builder.Limits.testing, "accy_choir_test_signed_compare", &params);
 829     errdefer builder.deinit();
 830     const out = builder.argument(0);
 831     const input = builder.argument(1);
 832     const index = try builder.globalId(.x);
 833     const value = try builder.load(input, index);
 834     const zero = try builder.constantInt(.i32, 0);
 835     const one = try builder.constantInt(.i32, 1);
 836     const pred = try builder.compare(.gt, value, zero);
 837     const selected = try builder.select(pred, one, zero);
 838     try builder.store(selected, out, index);
 839     try builder.return_();
 840 
 841     var program = try builder.finish();
 842     defer program.deinit();
 843     const text = try emitLoweredPtx(testing.allocator, "accy_choir_test_signed_compare", &program);
 844     defer testing.allocator.free(text);
 845 
 846     try testing.expect(std.mem.indexOf(u8, text, "setp.gt.s32") != null);
 847     try testing.expect(std.mem.indexOf(u8, text, "setp.gt.u32") == null);
 848 }
 849 
 850 test "cuda choir ptx emitter lowers narrow integer storage with 32-bit compute" {
 851     const testing = std.testing;
 852     const builder_mod = accy.kernel;
 853     const params = [_]builder_mod.Param{
 854         builder_mod.dynamicBuffer(.i8),
 855         builder_mod.dynamicBuffer(.i8),
 856         builder_mod.dynamicBuffer(.u16),
 857         builder_mod.dynamicBuffer(.u16),
 858     };
 859 
 860     var builder = try builder_mod.Builder.init(testing.allocator, builder_mod.Builder.Limits.testing, "accy_choir_test_narrow_int_ops", &params);
 861     errdefer builder.deinit();
 862     const i8_out = builder.argument(0);
 863     const i8_input = builder.argument(1);
 864     const u16_out = builder.argument(2);
 865     const u16_input = builder.argument(3);
 866     const index = try builder.globalId(.x);
 867 
 868     const signed = try builder.load(i8_input, index);
 869     const signed_zero = try builder.constantInt(.i8, 0);
 870     const signed_two = try builder.constantInt(.i8, 2);
 871     const signed_sum = try builder.add(signed, signed_two);
 872     const signed_bounded = try builder.max(try builder.min(signed_sum, signed_two), signed_zero);
 873     const signed_positive = try builder.compare(.gt, signed, signed_zero);
 874     const signed_selected = try builder.select(signed_positive, signed_bounded, signed);
 875     try builder.store(signed_selected, i8_out, index);
 876 
 877     const unsigned = try builder.load(u16_input, index);
 878     const unsigned_one = try builder.constantInt(.u16, 1);
 879     const unsigned_two = try builder.constantInt(.u16, 2);
 880     const unsigned_sum = try builder.add(unsigned, unsigned_one);
 881     const unsigned_bounded = try builder.min(try builder.max(unsigned_sum, unsigned_one), unsigned_two);
 882     const unsigned_positive = try builder.compare(.gt, unsigned, unsigned_one);
 883     const unsigned_selected = try builder.select(unsigned_positive, unsigned_bounded, unsigned);
 884     try builder.store(unsigned_selected, u16_out, index);
 885     try builder.return_();
 886 
 887     var program = try builder.finish();
 888     defer program.deinit();
 889     const text = try emitLoweredPtx(testing.allocator, "accy_choir_test_narrow_int_ops", &program);
 890     defer testing.allocator.free(text);
 891 
 892     try testing.expect(std.mem.indexOf(u8, text, "ld.global.s8") != null);
 893     try testing.expect(std.mem.indexOf(u8, text, "st.global.u8") != null);
 894     try testing.expect(std.mem.indexOf(u8, text, "min.s32") != null);
 895     try testing.expect(std.mem.indexOf(u8, text, "max.s32") != null);
 896     try testing.expect(std.mem.indexOf(u8, text, "setp.gt.s32") != null);
 897     try testing.expect(std.mem.indexOf(u8, text, "ld.global.u16") != null);
 898     try testing.expect(std.mem.indexOf(u8, text, "st.global.u16") != null);
 899     try testing.expect(std.mem.indexOf(u8, text, "min.u32") != null);
 900     try testing.expect(std.mem.indexOf(u8, text, "max.u32") != null);
 901     try testing.expect(std.mem.indexOf(u8, text, "setp.gt.u32") != null);
 902 }
 903 
 904 test "cuda choir ptx emitter clamps signed indices before index cast" {
 905     const testing = std.testing;
 906     const builder_mod = accy.kernel;
 907     const params = [_]builder_mod.Param{
 908         builder_mod.dynamicBuffer(.f32),
 909         builder_mod.dynamicBuffer(.f32),
 910         builder_mod.dynamicBuffer(.i32),
 911     };
 912 
 913     var builder = try builder_mod.Builder.init(testing.allocator, builder_mod.Builder.Limits.testing, "accy_choir_test_signed_index_clamp", &params);
 914     errdefer builder.deinit();
 915     const out = builder.argument(0);
 916     const data = builder.argument(1);
 917     const indices = builder.argument(2);
 918     const index = try builder.globalId(.x);
 919     const loaded = try builder.load(indices, index);
 920     const zero = try builder.constantInt(.i32, 0);
 921     const upper = try builder.constantInt(.i32, 5);
 922     const bounded = try builder.min(try builder.max(loaded, zero), upper);
 923     const source = try builder.castIndex(bounded);
 924     const value = try builder.load(data, source);
 925     try builder.store(value, out, index);
 926     try builder.return_();
 927 
 928     var program = try builder.finish();
 929     defer program.deinit();
 930     const text = try emitLoweredPtx(testing.allocator, "accy_choir_test_signed_index_clamp", &program);
 931     defer testing.allocator.free(text);
 932 
 933     try testing.expect(std.mem.indexOf(u8, text, "max.s32") != null);
 934     try testing.expect(std.mem.indexOf(u8, text, "min.s32") != null);
 935     try testing.expect(std.mem.indexOf(u8, text, "max.u32") == null);
 936     try testing.expect(std.mem.indexOf(u8, text, "min.u32") == null);
 937 }
 938 
 939 test "cuda choir ptx emitter lowers i64 scalar operations" {
 940     const testing = std.testing;
 941     const builder_mod = accy.kernel;
 942     const params = [_]builder_mod.Param{
 943         builder_mod.dynamicBuffer(.i64),
 944         builder_mod.dynamicBuffer(.i64),
 945     };
 946 
 947     var builder = try builder_mod.Builder.init(testing.allocator, builder_mod.Builder.Limits.testing, "accy_choir_test_i64_ops", &params);
 948     errdefer builder.deinit();
 949     const out = builder.argument(0);
 950     const input = builder.argument(1);
 951     const index = try builder.globalId(.x);
 952 
 953     const x = try builder.load(input, index);
 954     const zero = try builder.constantInt(.i64, 0);
 955     const two = try builder.constantInt(.i64, 2);
 956     const index_i64 = try builder.cast(index, .i64);
 957     const absolute = try builder.abs(x);
 958     const negated = try builder.neg(absolute);
 959     const product = try builder.mul(absolute, two);
 960     const quotient = try builder.div(product, two);
 961     const bounded = try builder.max(try builder.min(quotient, product), negated);
 962     const shifted = try builder.add(bounded, index_i64);
 963     const positive = try builder.compare(.gt, x, zero);
 964     const selected = try builder.select(positive, shifted, absolute);
 965     try builder.store(selected, out, index);
 966     try builder.return_();
 967 
 968     var program = try builder.finish();
 969     defer program.deinit();
 970     const text = try emitLoweredPtx(testing.allocator, "accy_choir_test_i64_ops", &program);
 971     defer testing.allocator.free(text);
 972 
 973     try testing.expect(std.mem.indexOf(u8, text, ".reg .b64") != null);
 974     try testing.expect(std.mem.indexOf(u8, text, "ld.global.u64") != null);
 975     try testing.expect(std.mem.indexOf(u8, text, "mov.u64") != null);
 976     try testing.expect(std.mem.indexOf(u8, text, "cvt.u64.u32") != null);
 977     try testing.expect(std.mem.indexOf(u8, text, "abs.s64") != null);
 978     try testing.expect(std.mem.indexOf(u8, text, "neg.s64") != null);
 979     try testing.expect(std.mem.indexOf(u8, text, "mul.lo.u64") != null);
 980     try testing.expect(std.mem.indexOf(u8, text, "div.s64") != null);
 981     try testing.expect(std.mem.indexOf(u8, text, "min.s64") != null);
 982     try testing.expect(std.mem.indexOf(u8, text, "max.s64") != null);
 983     try testing.expect(std.mem.indexOf(u8, text, "setp.gt.s64") != null);
 984     try testing.expect(std.mem.indexOf(u8, text, "selp.b64") != null);
 985     try testing.expect(std.mem.indexOf(u8, text, "st.global.u64") != null);
 986 }
 987 
 988 test "cuda choir ptx emitter lowers u64 scalar operations" {
 989     const testing = std.testing;
 990     const builder_mod = accy.kernel;
 991     const params = [_]builder_mod.Param{
 992         builder_mod.dynamicBuffer(.u64),
 993         builder_mod.dynamicBuffer(.u64),
 994     };
 995 
 996     var builder = try builder_mod.Builder.init(testing.allocator, builder_mod.Builder.Limits.testing, "accy_choir_test_u64_ops", &params);
 997     errdefer builder.deinit();
 998     const out = builder.argument(0);
 999     const input = builder.argument(1);
1000     const index = try builder.globalId(.x);
1001 
1002     const x = try builder.load(input, index);
1003     const zero = try builder.constantInt(.u64, 0);
1004     const two = try builder.constantInt(.u64, 2);
1005     const index_u64 = try builder.cast(index, .u64);
1006     const absolute = try builder.abs(x);
1007     const negated = try builder.neg(absolute);
1008     const product = try builder.mul(absolute, two);
1009     const quotient = try builder.div(product, two);
1010     const bounded = try builder.max(try builder.min(quotient, product), negated);
1011     const shifted = try builder.add(bounded, index_u64);
1012     const positive = try builder.compare(.gt, x, zero);
1013     const selected = try builder.select(positive, shifted, absolute);
1014     try builder.store(selected, out, index);
1015     try builder.return_();
1016 
1017     var program = try builder.finish();
1018     defer program.deinit();
1019     const text = try emitLoweredPtx(testing.allocator, "accy_choir_test_u64_ops", &program);
1020     defer testing.allocator.free(text);
1021 
1022     try testing.expect(std.mem.indexOf(u8, text, ".reg .b64") != null);
1023     try testing.expect(std.mem.indexOf(u8, text, "ld.global.u64") != null);
1024     try testing.expect(std.mem.indexOf(u8, text, "mov.u64") != null);
1025     try testing.expect(std.mem.indexOf(u8, text, "cvt.u64.u32") != null);
1026     try testing.expect(std.mem.indexOf(u8, text, "neg.s64") != null);
1027     try testing.expect(std.mem.indexOf(u8, text, "mul.lo.u64") != null);
1028     try testing.expect(std.mem.indexOf(u8, text, "div.u64") != null);
1029     try testing.expect(std.mem.indexOf(u8, text, "min.u64") != null);
1030     try testing.expect(std.mem.indexOf(u8, text, "max.u64") != null);
1031     try testing.expect(std.mem.indexOf(u8, text, "setp.gt.u64") != null);
1032     try testing.expect(std.mem.indexOf(u8, text, "selp.b64") != null);
1033     try testing.expect(std.mem.indexOf(u8, text, "st.global.u64") != null);
1034 }
1035 
1036 test "cuda choir ptx emitter lowers f64 scalar operations" {
1037     const testing = std.testing;
1038     const builder_mod = accy.kernel;
1039     const params = [_]builder_mod.Param{
1040         builder_mod.dynamicBuffer(.f64),
1041         builder_mod.dynamicBuffer(.f64),
1042     };
1043 
1044     var builder = try builder_mod.Builder.init(testing.allocator, builder_mod.Builder.Limits.testing, "accy_choir_test_f64_ops", &params);
1045     errdefer builder.deinit();
1046     const out = builder.argument(0);
1047     const input = builder.argument(1);
1048     const index = try builder.globalId(.x);
1049 
1050     const x = try builder.load(input, index);
1051     const zero = try builder.constantFloat(.f64, 0.0);
1052     const two = try builder.constantFloat(.f64, 2.0);
1053     const index_f64 = try builder.cast(index, .f64);
1054     const absolute = try builder.abs(x);
1055     const negated = try builder.neg(absolute);
1056     const product = try builder.mul(absolute, two);
1057     const quotient = try builder.div(product, two);
1058     const bounded = try builder.max(try builder.min(quotient, product), negated);
1059     const shifted = try builder.add(bounded, index_f64);
1060     const positive = try builder.compare(.gt, x, zero);
1061     const selected = try builder.select(positive, shifted, absolute);
1062     try builder.store(selected, out, index);
1063     try builder.return_();
1064 
1065     var program = try builder.finish();
1066     defer program.deinit();
1067     const text = try emitLoweredPtx(testing.allocator, "accy_choir_test_f64_ops", &program);
1068     defer testing.allocator.free(text);
1069 
1070     try testing.expect(std.mem.indexOf(u8, text, ".reg .f64") != null);
1071     try testing.expect(std.mem.indexOf(u8, text, "ld.global.f64") != null);
1072     try testing.expect(std.mem.indexOf(u8, text, "0d4000000000000000") != null);
1073     try testing.expect(std.mem.indexOf(u8, text, "cvt.rn.f64.u32") != null);
1074     try testing.expect(std.mem.indexOf(u8, text, "abs.f64") != null);
1075     try testing.expect(std.mem.indexOf(u8, text, "neg.f64") != null);
1076     try testing.expect(std.mem.indexOf(u8, text, "mul.f64") != null);
1077     try testing.expect(std.mem.indexOf(u8, text, "div.rn.f64") != null);
1078     try testing.expect(std.mem.indexOf(u8, text, "min.f64") != null);
1079     try testing.expect(std.mem.indexOf(u8, text, "max.f64") != null);
1080     try testing.expect(std.mem.indexOf(u8, text, "setp.gt.f64") != null);
1081     try testing.expect(std.mem.indexOf(u8, text, "selp.f64") != null);
1082     try testing.expect(std.mem.indexOf(u8, text, "st.global.f64") != null);
1083 }
1084 
1085 test "cuda choir ptx emitter lowers bf16 storage with f32 compute" {
1086     const testing = std.testing;
1087     const builder_mod = accy.kernel;
1088     const params = [_]builder_mod.Param{
1089         builder_mod.dynamicBuffer(.bf16),
1090         builder_mod.dynamicBuffer(.bf16),
1091     };
1092 
1093     var builder = try builder_mod.Builder.init(testing.allocator, builder_mod.Builder.Limits.testing, "accy_choir_test_bf16_ops", &params);
1094     errdefer builder.deinit();
1095     const out = builder.argument(0);
1096     const input = builder.argument(1);
1097     const index = try builder.globalId(.x);
1098 
1099     const x = try builder.load(input, index);
1100     const zero = try builder.constantFloat(.bf16, 0.0);
1101     const two = try builder.constantFloat(.bf16, 2.0);
1102     const index_bf16 = try builder.cast(index, .bf16);
1103     const absolute = try builder.abs(x);
1104     const negated = try builder.neg(absolute);
1105     const product = try builder.mul(absolute, two);
1106     const quotient = try builder.div(product, two);
1107     const bounded = try builder.max(try builder.min(quotient, product), negated);
1108     const shifted = try builder.add(bounded, index_bf16);
1109     const positive = try builder.compare(.gt, x, zero);
1110     const selected = try builder.select(positive, shifted, absolute);
1111     try builder.store(selected, out, index);
1112     try builder.return_();
1113 
1114     var program = try builder.finish();
1115     defer program.deinit();
1116     const text = try emitLoweredPtx(testing.allocator, "accy_choir_test_bf16_ops", &program);
1117     defer testing.allocator.free(text);
1118 
1119     try testing.expect(std.mem.indexOf(u8, text, "ld.global.b16      %r") != null);
1120     try testing.expect(std.mem.indexOf(u8, text, "shl.b32") != null);
1121     try testing.expect(std.mem.indexOf(u8, text, "mov.b32            %f") != null);
1122     try testing.expect(std.mem.indexOf(u8, text, "0f40000000") != null);
1123     try testing.expect(std.mem.indexOf(u8, text, "cvt.rn.f32.u32") != null);
1124     try testing.expect(std.mem.indexOf(u8, text, "abs.f32") != null);
1125     try testing.expect(std.mem.indexOf(u8, text, "neg.f32") != null);
1126     try testing.expect(std.mem.indexOf(u8, text, "mul.f32") != null);
1127     try testing.expect(std.mem.indexOf(u8, text, "div.approx.f32") != null);
1128     try testing.expect(std.mem.indexOf(u8, text, "min.f32") != null);
1129     try testing.expect(std.mem.indexOf(u8, text, "max.f32") != null);
1130     try testing.expect(std.mem.indexOf(u8, text, "setp.gt.f32") != null);
1131     try testing.expect(std.mem.indexOf(u8, text, "selp.f32") != null);
1132     try testing.expect(std.mem.indexOf(u8, text, "shr.u32") != null);
1133     try testing.expect(std.mem.indexOf(u8, text, "st.global.b16") != null);
1134 }
1135 
1136 test "cuda choir ptx emitter lowers scalar math and bit operations" {
1137     const testing = std.testing;
1138     const builder_mod = accy.kernel;
1139     const params = [_]builder_mod.Param{
1140         builder_mod.dynamicBuffer(.f32),
1141         builder_mod.dynamicBuffer(.f32),
1142         builder_mod.dynamicBuffer(.i32),
1143         builder_mod.dynamicBuffer(.i32),
1144     };
1145 
1146     var builder = try builder_mod.Builder.init(testing.allocator, builder_mod.Builder.Limits.testing, "accy_choir_test_scalar_ops", &params);
1147     errdefer builder.deinit();
1148     const f_out = builder.argument(0);
1149     const f_in = builder.argument(1);
1150     const i_out = builder.argument(2);
1151     const i_in = builder.argument(3);
1152     const index = try builder.globalId(.x);
1153 
1154     const x = try builder.load(f_in, index);
1155     const zero = try builder.constantFloat(.f32, 0.0);
1156     const two = try builder.constantFloat(.f32, 2.0);
1157     const abs_x = try builder.abs(x);
1158     const root = try builder.sqrt(abs_x);
1159     const s = try builder.sin(x);
1160     const c = try builder.cos(x);
1161     const t = try builder.tan(x);
1162     const floored = try builder.floor(x);
1163     const rounded = try builder.round(x);
1164     const truncated = try builder.trunc(x);
1165     const p = try builder.pow(abs_x, two);
1166     const angle = try builder.atan2(x, root);
1167     const fma = try builder.fma(s, c, root);
1168     const divided = try builder.div(fma, two);
1169     const sum = try builder.add(try builder.add(try builder.add(try builder.add(try builder.add(try builder.add(divided, t), p), angle), floored), rounded), truncated);
1170     const pred = try builder.compare(.gt, x, zero);
1171     const selected = try builder.select(pred, sum, abs_x);
1172     const bits = try builder.bitcast(selected, .i32);
1173     const roundtrip = try builder.bitcast(bits, .f32);
1174     try builder.store(roundtrip, f_out, index);
1175 
1176     const int_value = try builder.load(i_in, index);
1177     const one = try builder.constantInt(.i32, 1);
1178     const shl = try builder.shl(int_value, one);
1179     const shr = try builder.shr(int_value, one);
1180     const ushr = try builder.ushr(int_value, one);
1181     const anded = try builder.and_(shl, ushr);
1182     const ored = try builder.or_(anded, one);
1183     const xored = try builder.xor(ored, shr);
1184     const inverted = try builder.not(xored);
1185     const high = try builder.umulhi(inverted, int_value);
1186 
1187     const wide = try builder.cast(high, .i64);
1188     const wide_mask = try builder.constantInt(.i64, 0xFFFFFFFF);
1189     const wide_masked = try builder.and_(wide, wide_mask);
1190     const wide_square = try builder.mul(wide_masked, wide_masked);
1191     const wide_sum = try builder.add(wide_square, wide_masked);
1192     const wide_half = try builder.constantInt(.i64, 32);
1193     const wide_rot = try builder.or_(try builder.ushr(wide_sum, wide_half), try builder.shl(wide_sum, wide_half));
1194     const wide_high = try builder.umulhi(wide_rot, wide_square);
1195     const narrowed = try builder.cast(wide_high, .i32);
1196     try builder.store(narrowed, i_out, index);
1197     try builder.return_();
1198 
1199     var program = try builder.finish();
1200     defer program.deinit();
1201     const text = try emitLoweredPtx(testing.allocator, "accy_choir_test_scalar_ops", &program);
1202     defer testing.allocator.free(text);
1203 
1204     try testing.expect(std.mem.indexOf(u8, text, "div.approx.f32") != null);
1205     try testing.expect(std.mem.indexOf(u8, text, "abs.f32") != null);
1206     try testing.expect(std.mem.indexOf(u8, text, "sqrt.approx.f32") != null);
1207     try testing.expect(std.mem.indexOf(u8, text, "sin.approx.f32") != null);
1208     try testing.expect(std.mem.indexOf(u8, text, "cos.approx.f32") != null);
1209     try testing.expect(std.mem.indexOf(u8, text, "cvt.rmi.f32.f32") != null);
1210     try testing.expect(std.mem.indexOf(u8, text, "0f3F000000") != null);
1211     try testing.expect(std.mem.indexOf(u8, text, "cvt.rzi.f32.f32") != null);
1212     try testing.expect(std.mem.indexOf(u8, text, "lg2.approx.f32") != null);
1213     try testing.expect(std.mem.indexOf(u8, text, "ex2.approx.f32") != null);
1214     try testing.expect(std.mem.indexOf(u8, text, "0f3F490FDB") != null);
1215     try testing.expect(std.mem.indexOf(u8, text, "0f3FC90FDB") != null);
1216     try testing.expect(std.mem.indexOf(u8, text, "fma.rn.f32") != null);
1217     try testing.expect(std.mem.indexOf(u8, text, "selp.f32") != null);
1218     try testing.expect(std.mem.indexOf(u8, text, "mov.b32") != null);
1219     try testing.expect(std.mem.indexOf(u8, text, "shl.b32") != null);
1220     try testing.expect(std.mem.indexOf(u8, text, "shr.s32") != null);
1221     try testing.expect(std.mem.indexOf(u8, text, "shr.u32") != null);
1222     try testing.expect(std.mem.indexOf(u8, text, "and.b32") != null);
1223     try testing.expect(std.mem.indexOf(u8, text, "or.b32") != null);
1224     try testing.expect(std.mem.indexOf(u8, text, "xor.b32") != null);
1225     try testing.expect(std.mem.indexOf(u8, text, "not.b32") != null);
1226     try testing.expect(std.mem.indexOf(u8, text, "mul.hi.u32") != null);
1227     try testing.expect(std.mem.indexOf(u8, text, "cvt.s64.s32") != null);
1228     try testing.expect(std.mem.indexOf(u8, text, "and.b64") != null);
1229     try testing.expect(std.mem.indexOf(u8, text, "mul.lo.u64") != null);
1230     try testing.expect(std.mem.indexOf(u8, text, "add.u64") != null);
1231     try testing.expect(std.mem.indexOf(u8, text, "shl.b64") != null);
1232     try testing.expect(std.mem.indexOf(u8, text, "shr.u64") != null);
1233     try testing.expect(std.mem.indexOf(u8, text, "or.b64") != null);
1234     try testing.expect(std.mem.indexOf(u8, text, "mul.hi.u64") != null);
1235     try testing.expect(std.mem.indexOf(u8, text, "cvt.u32.u64") != null);
1236 }
1237 
1238 test "cuda choir ptx emitter lowers warp control intrinsics" {
1239     const testing = std.testing;
1240     const builder_mod = accy.kernel;
1241     const params = [_]builder_mod.Param{
1242         builder_mod.dynamicBuffer(.i32),
1243         builder_mod.dynamicBuffer(.f32),
1244         builder_mod.dynamicBuffer(.f32),
1245     };
1246 
1247     var builder = try builder_mod.Builder.init(testing.allocator, builder_mod.Builder.Limits.testing, "accy_choir_test_warp_ops", &params);
1248     errdefer builder.deinit();
1249     const i_out = builder.argument(0);
1250     const f_out = builder.argument(1);
1251     const f_in = builder.argument(2);
1252     const index = try builder.globalId(.x);
1253 
1254     const lane = try builder.laneId();
1255     _ = try builder.warpId();
1256     _ = try builder.activeMask();
1257     const lane_i32 = try builder.cast(lane, .i32);
1258     const zero = try builder.constantInt(.i32, 0);
1259     const one = try builder.constantInt(.i32, 1);
1260     const is_lane_zero = try builder.compare(.eq, lane_i32, zero);
1261     const all = try builder.allSync(is_lane_zero);
1262     const any = try builder.anySync(is_lane_zero);
1263     const ballot = try builder.ballotSync(is_lane_zero);
1264     const shuffled_lane = try builder.shuffleSync(.xor, lane_i32, one);
1265     const all_i32 = try builder.cast(all, .i32);
1266     const any_i32 = try builder.cast(any, .i32);
1267     const folded_vote = try builder.xor(try builder.xor(ballot, all_i32), any_i32);
1268     const folded = try builder.xor(folded_vote, shuffled_lane);
1269     try builder.store(folded, i_out, index);
1270 
1271     const x = try builder.load(f_in, index);
1272     const shuffled_x = try builder.shuffleSync(.down, x, one);
1273     try builder.store(shuffled_x, f_out, index);
1274     try builder.syncWarp();
1275     try builder.barrier(.warp);
1276     try builder.barrier(.block);
1277     try builder.return_();
1278 
1279     var program = try builder.finish();
1280     defer program.deinit();
1281     const text = try emitLoweredPtx(testing.allocator, "accy_choir_test_warp_ops", &program);
1282     defer testing.allocator.free(text);
1283 
1284     try testing.expect(std.mem.indexOf(u8, text, ".version 6.2") != null);
1285     try testing.expect(std.mem.indexOf(u8, text, "%laneid") != null);
1286     try testing.expect(std.mem.indexOf(u8, text, "%warpid") == null);
1287     try testing.expect(std.mem.indexOf(u8, text, "%tid.z") != null);
1288     try testing.expect(std.mem.indexOf(u8, text, "shr.u32") != null);
1289     try testing.expect(std.mem.indexOf(u8, text, "activemask.b32") != null);
1290     try testing.expect(std.mem.indexOf(u8, text, "vote.sync.all.pred") != null);
1291     try testing.expect(std.mem.indexOf(u8, text, "vote.sync.any.pred") != null);
1292     try testing.expect(std.mem.indexOf(u8, text, "vote.sync.ballot.b32") != null);
1293     try testing.expect(std.mem.indexOf(u8, text, "selp.u32") != null);
1294     try testing.expect(std.mem.indexOf(u8, text, "shfl.sync.bfly.b32") != null);
1295     try testing.expect(std.mem.indexOf(u8, text, "shfl.sync.down.b32") != null);
1296     try testing.expect(std.mem.indexOf(u8, text, "bar.warp.sync") != null);
1297     try testing.expect(std.mem.indexOf(u8, text, "bar.sync") != null);
1298 }
1299 
1300 test "cuda choir ptx emitter lowers warp reductions and scans" {
1301     const testing = std.testing;
1302     const builder_mod = accy.kernel;
1303     const params = [_]builder_mod.Param{
1304         builder_mod.dynamicBuffer(.i32),
1305         builder_mod.dynamicBuffer(.i32),
1306         builder_mod.dynamicBuffer(.f32),
1307         builder_mod.dynamicBuffer(.f32),
1308         builder_mod.dynamicBuffer(.f64),
1309         builder_mod.dynamicBuffer(.f64),
1310     };
1311 
1312     var builder = try builder_mod.Builder.init(testing.allocator, builder_mod.Builder.Limits.testing, "accy_choir_test_warp_collectives", &params);
1313     errdefer builder.deinit();
1314     const i_out = builder.argument(0);
1315     const i_in = builder.argument(1);
1316     const f_out = builder.argument(2);
1317     const f_in = builder.argument(3);
1318     const d_out = builder.argument(4);
1319     const d_in = builder.argument(5);
1320     const index = try builder.globalId(.x);
1321 
1322     const int_value = try builder.load(i_in, index);
1323     const int_reduce = try builder.warpReduce(.add, int_value);
1324     const int_scan = try builder.warpScan(.xor, .inclusive, int_value);
1325     const int_folded = try builder.xor(int_reduce, int_scan);
1326     try builder.store(int_folded, i_out, index);
1327 
1328     const float_value = try builder.load(f_in, index);
1329     const float_reduce = try builder.warpReduce(.max, float_value);
1330     const float_scan = try builder.warpScan(.add, .exclusive, float_value);
1331     const float_folded = try builder.add(float_reduce, float_scan);
1332     try builder.store(float_folded, f_out, index);
1333 
1334     const double_value = try builder.load(d_in, index);
1335     const double_reduce = try builder.warpReduce(.add, double_value);
1336     try builder.store(double_reduce, d_out, index);
1337     try builder.return_();
1338 
1339     var program = try builder.finish();
1340     defer program.deinit();
1341     const text = try emitLoweredPtx(testing.allocator, "accy_choir_test_warp_collectives", &program);
1342     defer testing.allocator.free(text);
1343 
1344     try testing.expect(std.mem.indexOf(u8, text, "shfl.sync.bfly.b32") != null);
1345     try testing.expect(std.mem.indexOf(u8, text, "shfl.sync.up.b32") != null);
1346     try testing.expect(std.mem.indexOf(u8, text, "0x0, %r") != null);
1347     try testing.expect(std.mem.indexOf(u8, text, "@%p") != null);
1348     try testing.expect(std.mem.indexOf(u8, text, "max.f32") != null);
1349     try testing.expect(std.mem.indexOf(u8, text, "mov.b64") != null);
1350     try testing.expect(std.mem.indexOf(u8, text, "add.f64") != null);
1351     try testing.expect(std.mem.indexOf(u8, text, "xor.b32") != null);
1352     try testing.expect(std.mem.indexOf(u8, text, "0f00000000") != null);
1353 }
1354 
1355 test "cuda prefix sum scan uses safe shared base index" {
1356     const testing = std.testing;
1357     const scan = accy.kernel.library.scan;
1358     const instance = scan.PrefixSum{ .extent = 64, .threads = 64 };
1359     const entry_name = try scan.prefixSumFamilyEntryName(testing.allocator, instance);
1360     defer testing.allocator.free(entry_name);
1361 
1362     var graph = try scan.PrefixSumRuntimeFamilyF32.buildNamed(
1363         testing.allocator,
1364         scan.PrefixSumRuntimeFamilyF32.Limits.standard,
1365         entry_name,
1366         instance,
1367     );
1368     defer graph.deinit();
1369     const text = try emitLoweredPtx(testing.allocator, entry_name, &graph);
1370     defer testing.allocator.free(text);
1371 
1372     try testing.expect(std.mem.indexOf(u8, text, "ld.shared.f32") != null);
1373     try testing.expect(std.mem.indexOf(u8, text, "selp.u32") != null);
1374     try testing.expect(std.mem.indexOf(u8, text, "max.u32") == null);
1375 }
1376 
1377 test "cuda segment sum clamps signed offsets before index cast" {
1378     const testing = std.testing;
1379     const segmented = accy.kernel.library.segmented;
1380     const instance = segmented.SegmentSum{ .segments = 8, .total = 80, .threads = 64 };
1381     const entry_name = try segmented.segmentSumFamilyEntryName(testing.allocator, instance);
1382     defer testing.allocator.free(entry_name);
1383 
1384     var graph = try segmented.SegmentSumRuntimeFamilyF32.buildNamed(
1385         testing.allocator,
1386         segmented.SegmentSumRuntimeFamilyF32.Limits.standard,
1387         entry_name,
1388         instance,
1389     );
1390     defer graph.deinit();
1391     const text = try emitLoweredPtx(testing.allocator, entry_name, &graph);
1392     defer testing.allocator.free(text);
1393 
1394     try testing.expect(std.mem.indexOf(u8, text, "max.s32") != null);
1395     try testing.expect(std.mem.indexOf(u8, text, "min.s32") != null);
1396     try testing.expect(std.mem.indexOf(u8, text, "max.u32") == null);
1397 }
1398 
1399 test "cuda choir ptx emitter lowers shared memory" {
1400     const testing = std.testing;
1401     const builder_mod = accy.kernel;
1402     const params = [_]builder_mod.Param{
1403         builder_mod.dynamicBuffer(.f32),
1404         builder_mod.dynamicBuffer(.f32),
1405     };
1406 
1407     var builder = try builder_mod.Builder.init(testing.allocator, builder_mod.Builder.Limits.testing, "accy_choir_test_shared_memory", &params);
1408     errdefer builder.deinit();
1409     const out = builder.argument(0);
1410     const input = builder.argument(1);
1411     const index = try builder.globalId(.x);
1412     const lane = try builder.laneId();
1413     const shared = try builder.sharedBuffer(.f32, 32);
1414 
1415     const x = try builder.load(input, index);
1416     try builder.store(x, shared, lane);
1417     try builder.barrier(.block);
1418     const y = try builder.load(shared, lane);
1419     try builder.store(y, out, index);
1420     try builder.return_();
1421 
1422     var program = try builder.finish();
1423     defer program.deinit();
1424     const text = try emitLoweredPtx(testing.allocator, "accy_choir_test_shared_memory", &program);
1425     defer testing.allocator.free(text);
1426 
1427     try testing.expect(std.mem.indexOf(u8, text, ".shared .align 16 .b8 __choir_shared") != null);
1428     try testing.expect(std.mem.indexOf(u8, text, "st.shared.f32") != null);
1429     try testing.expect(std.mem.indexOf(u8, text, "ld.shared.f32") != null);
1430     try testing.expect(std.mem.indexOf(u8, text, "bar.sync") != null);
1431 }
1432 
1433 test "cuda choir ptx emitter lowers dynamic shared memory offsets" {
1434     const testing = std.testing;
1435     const builder_mod = accy.kernel;
1436     const params = [_]builder_mod.Param{
1437         builder_mod.dynamicBuffer(.f32),
1438         builder_mod.dynamicBuffer(.f32),
1439     };
1440 
1441     var builder = try builder_mod.Builder.init(testing.allocator, builder_mod.Builder.Limits.testing, "accy_choir_test_dynamic_shared_memory", &params);
1442     errdefer builder.deinit();
1443     const out = builder.argument(0);
1444     const input = builder.argument(1);
1445     const index = try builder.globalId(.x);
1446     const lane = try builder.laneId();
1447     const first = try builder.dynamicSharedBuffer(.f32, 4, 0);
1448     const second = try builder.dynamicSharedBuffer(.f32, 4, 64);
1449 
1450     const x = try builder.load(input, index);
1451     try builder.store(x, first, lane);
1452     try builder.store(x, second, lane);
1453     try builder.barrier(.block);
1454     const y = try builder.load(second, lane);
1455     try builder.store(y, out, index);
1456     try builder.return_();
1457 
1458     var program = try builder.finish();
1459     defer program.deinit();
1460     const text = try emitLoweredPtx(testing.allocator, "accy_choir_test_dynamic_shared_memory", &program);
1461     defer testing.allocator.free(text);
1462 
1463     try testing.expect(std.mem.indexOf(u8, text, ".extern .shared .align 16 .b8 __choir_dynamic_shared[]") != null);
1464     try testing.expect(std.mem.indexOf(u8, text, ".shared .align 16 .b8 __choir_shared") == null);
1465     try testing.expect(std.mem.indexOf(u8, text, "mov.u32            %r") != null);
1466     try testing.expect(std.mem.indexOf(u8, text, "__choir_dynamic_shared") != null);
1467     try testing.expect(std.mem.indexOf(u8, text, "+64]") != null);
1468     try testing.expect(std.mem.indexOf(u8, text, "st.shared.f32") != null);
1469     try testing.expect(std.mem.indexOf(u8, text, "ld.shared.f32") != null);
1470 }
1471 
1472 test "cuda choir ptx emitter lowers global atomics" {
1473     const testing = std.testing;
1474     const builder_mod = accy.kernel;
1475     const params = [_]builder_mod.Param{
1476         builder_mod.dynamicBuffer(.f32),
1477         builder_mod.dynamicBuffer(.i32),
1478     };
1479 
1480     var builder = try builder_mod.Builder.init(testing.allocator, builder_mod.Builder.Limits.testing, "accy_choir_test_atomics", &params);
1481     errdefer builder.deinit();
1482     const f_bins = builder.argument(0);
1483     const i_bins = builder.argument(1);
1484     const index = try builder.globalId(.x);
1485     const f_delta = try builder.constantFloat(.f32, 1.5);
1486     const i_zero = try builder.constantInt(.i32, 0);
1487     const i_one = try builder.constantInt(.i32, 1);
1488     const i_three = try builder.constantInt(.i32, 3);
1489 
1490     _ = try builder.atomicRmw(.add, f_delta, f_bins, index);
1491     const old_add = try builder.atomicRmw(.add, i_one, i_bins, index);
1492     const old_min = try builder.atomicRmw(.min, old_add, i_bins, index);
1493     const old_max = try builder.atomicRmw(.max, old_min, i_bins, index);
1494     const old_and = try builder.atomicRmw(.bit_and, old_max, i_bins, index);
1495     const old_or = try builder.atomicRmw(.bit_or, old_and, i_bins, index);
1496     const old_xor = try builder.atomicRmw(.bit_xor, old_or, i_bins, index);
1497     _ = try builder.atomicCas(i_zero, i_three, i_bins, index);
1498     _ = try builder.atomicRmw(.exchange, old_xor, i_bins, index);
1499     _ = try builder.atomicRmw(.bit_or, i_one, i_bins, index);
1500     try builder.return_();
1501 
1502     var program = try builder.finish();
1503     defer program.deinit();
1504     const text = try emitLoweredPtx(testing.allocator, "accy_choir_test_atomics", &program);
1505     defer testing.allocator.free(text);
1506 
1507     try testing.expect(std.mem.indexOf(u8, text, "red.global.add.f32") != null);
1508     try testing.expect(std.mem.indexOf(u8, text, "atom.global.add.f32") == null);
1509     try testing.expect(std.mem.indexOf(u8, text, "atom.global.add.u32") != null);
1510     try testing.expect(std.mem.indexOf(u8, text, "atom.global.min.s32") != null);
1511     try testing.expect(std.mem.indexOf(u8, text, "atom.global.max.s32") != null);
1512     try testing.expect(std.mem.indexOf(u8, text, "atom.global.and.b32") != null);
1513     try testing.expect(std.mem.indexOf(u8, text, "atom.global.or.b32") != null);
1514     try testing.expect(std.mem.indexOf(u8, text, "atom.global.xor.b32") != null);
1515     try testing.expect(std.mem.indexOf(u8, text, "atom.global.cas.b32") != null);
1516     try testing.expect(std.mem.indexOf(u8, text, "atom.global.exch.b32") != null);
1517     try testing.expect(std.mem.indexOf(u8, text, "red.global.or.b32") != null);
1518     try testing.expect(std.mem.indexOf(u8, text, "red.global.exch") == null);
1519 }
1520 
1521 test "cuda choir ptx emitter lowers shared atomics" {
1522     const testing = std.testing;
1523     const builder_mod = accy.kernel;
1524     const params = [_]builder_mod.Param{
1525         builder_mod.dynamicBuffer(.i32),
1526         builder_mod.dynamicBuffer(.f32),
1527     };
1528 
1529     var builder = try builder_mod.Builder.init(testing.allocator, builder_mod.Builder.Limits.testing, "accy_choir_test_shared_atomics", &params);
1530     errdefer builder.deinit();
1531     const i_out = builder.argument(0);
1532     const f_out = builder.argument(1);
1533     const index = try builder.globalId(.x);
1534     const i_bins = try builder.sharedBuffer(.i32, 16);
1535     const f_bins = try builder.sharedBuffer(.f32, 16);
1536     const slot = try builder.constantIndex(0);
1537     const i_one = try builder.constantInt(.i32, 1);
1538     const i_three = try builder.constantInt(.i32, 3);
1539     const f_delta = try builder.constantFloat(.f32, 0.5);
1540 
1541     const old_add = try builder.atomicRmw(.add, i_one, i_bins, slot);
1542     _ = try builder.atomicRmw(.max, old_add, i_bins, slot);
1543     _ = try builder.atomicCas(i_one, i_three, i_bins, slot);
1544     const f_old = try builder.atomicRmw(.add, f_delta, f_bins, slot);
1545     try builder.barrier(.block);
1546     const merged = try builder.atomicRmw(.add, old_add, i_out, index);
1547     _ = merged;
1548     try builder.store(f_old, f_out, index);
1549     try builder.return_();
1550 
1551     var program = try builder.finish();
1552     defer program.deinit();
1553     const text = try emitLoweredPtx(testing.allocator, "accy_choir_test_shared_atomics", &program);
1554     defer testing.allocator.free(text);
1555 
1556     try testing.expect(std.mem.indexOf(u8, text, "atom.shared.add.u32") != null);
1557     try testing.expect(std.mem.indexOf(u8, text, "red.shared.max.s32") != null);
1558     try testing.expect(std.mem.indexOf(u8, text, "atom.shared.cas.b32") != null);
1559     try testing.expect(std.mem.indexOf(u8, text, "atom.shared.add.f32") != null);
1560     try testing.expect(std.mem.indexOf(u8, text, "red.global.add.u32") != null);
1561     try testing.expect(std.mem.indexOf(u8, text, "atom.global.add.u32") == null);
1562     try testing.expect(std.mem.indexOf(u8, text, "bar.sync") != null);
1563 }