mirror of
https://github.com/ziglang/zig.git
synced 2025-12-06 06:13:07 +00:00
4b16d81318
error: thread 165232 panic: TODO: rewrite the NvPtx.flushModule function
/home/alexrp/Source/ziglang/zig/src/link/NvPtx.zig:123:5: 0x1ed99ce in flushModule (zig)
@panic("TODO: rewrite the NvPtx.flushModule function");
^
/home/alexrp/Source/ziglang/zig/src/link/NvPtx.zig:110:28: 0x1c2e7e6 in flush (zig)
return self.flushModule(arena, tid, prog_node);
^
/home/alexrp/Source/ziglang/zig/src/link.zig:874:77: 0x1a55bd3 in flush (zig)
return @as(*tag.Type(), @fieldParentPtr("base", base)).flush(arena, tid, prog_node);
^
/home/alexrp/Source/ziglang/zig/src/Compilation.zig:2411:17: 0x1a553f9 in flush (zig)
lf.flush(arena, tid, prog_node) catch |err| switch (err) {
^
/home/alexrp/Source/ziglang/zig/src/Compilation.zig:2348:22: 0x1a595ba in update (zig)
try flush(comp, arena, .{
^
/home/alexrp/Source/ziglang/zig/src/main.zig:4114:32: 0x1ae392b in serve (zig)
try comp.update(main_progress_node);
^
/home/alexrp/Source/ziglang/zig/src/main.zig:3555:22: 0x1b05322 in buildOutputType (zig)
try serve(
^
/home/alexrp/Source/ziglang/zig/src/main.zig:265:31: 0x195faca in mainArgs (zig)
return buildOutputType(gpa, arena, args, .{ .build = .Obj });
^
/home/alexrp/Source/ziglang/zig/src/main.zig:200:20: 0x195c995 in main (zig)
return mainArgs(gpa, arena, args);
^
/home/alexrp/Source/ziglang/zig/lib/std/start.zig:617:37: 0x195c49e in main (zig)
const result = root.main() catch |err| {
^
105 lines
3.2 KiB
Zig
105 lines
3.2 KiB
Zig
const std = @import("std");
|
|
const Cases = @import("src/Cases.zig");
|
|
|
|
pub fn addCases(ctx: *Cases, b: *std.Build) !void {
|
|
const target = b.resolveTargetQuery(.{
|
|
.cpu_arch = .nvptx64,
|
|
.os_tag = .cuda,
|
|
});
|
|
|
|
{
|
|
var case = addPtx(ctx, target, "simple addition and subtraction");
|
|
|
|
case.addCompile(
|
|
\\fn add(a: i32, b: i32) i32 {
|
|
\\ return a + b;
|
|
\\}
|
|
\\
|
|
\\pub export fn add_and_substract(a: i32, out: *i32) callconv(.Kernel) void {
|
|
\\ const x = add(a, 7);
|
|
\\ var y = add(2, 0);
|
|
\\ y -= x;
|
|
\\ out.* = y;
|
|
\\}
|
|
);
|
|
}
|
|
|
|
{
|
|
var case = addPtx(ctx, target, "read special registers");
|
|
|
|
case.addCompile(
|
|
\\fn threadIdX() u32 {
|
|
\\ return asm ("mov.u32 \t%[r], %tid.x;"
|
|
\\ : [r] "=r" (-> u32),
|
|
\\ );
|
|
\\}
|
|
\\
|
|
\\pub export fn special_reg(a: []const i32, out: []i32) callconv(.Kernel) void {
|
|
\\ const i = threadIdX();
|
|
\\ out[i] = a[i] + 7;
|
|
\\}
|
|
);
|
|
}
|
|
|
|
{
|
|
var case = addPtx(ctx, target, "address spaces");
|
|
|
|
case.addCompile(
|
|
\\var x: i32 addrspace(.global) = 0;
|
|
\\
|
|
\\pub export fn increment(out: *i32) callconv(.Kernel) void {
|
|
\\ x += 1;
|
|
\\ out.* = x;
|
|
\\}
|
|
);
|
|
}
|
|
|
|
{
|
|
var case = addPtx(ctx, target, "reduce in shared mem");
|
|
case.addCompile(
|
|
\\fn threadIdX() u32 {
|
|
\\ return asm ("mov.u32 \t%[r], %tid.x;"
|
|
\\ : [r] "=r" (-> u32),
|
|
\\ );
|
|
\\}
|
|
\\
|
|
\\ var _sdata: [1024]f32 addrspace(.shared) = undefined;
|
|
\\ pub export fn reduceSum(d_x: []const f32, out: *f32) callconv(.Kernel) void {
|
|
\\ var sdata: *addrspace(.generic) [1024]f32 = @addrSpaceCast(&_sdata);
|
|
\\ const tid: u32 = threadIdX();
|
|
\\ var sum = d_x[tid];
|
|
\\ sdata[tid] = sum;
|
|
\\ asm volatile ("bar.sync \t0;");
|
|
\\ var s: u32 = 512;
|
|
\\ while (s > 0) : (s = s >> 1) {
|
|
\\ if (tid < s) {
|
|
\\ sum += sdata[tid + s];
|
|
\\ sdata[tid] = sum;
|
|
\\ }
|
|
\\ asm volatile ("bar.sync \t0;");
|
|
\\ }
|
|
\\
|
|
\\ if (tid == 0) {
|
|
\\ out.* = sum;
|
|
\\ }
|
|
\\ }
|
|
);
|
|
}
|
|
}
|
|
|
|
fn addPtx(ctx: *Cases, target: std.Build.ResolvedTarget, name: []const u8) *Cases.Case {
|
|
ctx.cases.append(.{
|
|
.name = name,
|
|
.target = target,
|
|
.updates = std.ArrayList(Cases.Update).init(ctx.cases.allocator),
|
|
.output_mode = .Obj,
|
|
.deps = std.ArrayList(Cases.DepModule).init(ctx.cases.allocator),
|
|
.link_libc = false,
|
|
.emit_bin = false,
|
|
.backend = .llvm,
|
|
// Bug in Debug mode
|
|
.optimize_mode = .ReleaseSafe,
|
|
}) catch @panic("out of memory");
|
|
return &ctx.cases.items[ctx.cases.items.len - 1];
|
|
}
|