| ... | @@ -1,105 +0,0 @@ |
| 1 | const std = @import("std"); |
| 2 | const Cases = @import("src/Cases.zig"); |
| 3 | |
| 4 | pub fn addCases(ctx: *Cases, b: *std.Build) !void { |
| 5 | const target = b.resolveTargetQuery(.{ |
| 6 | .cpu_arch = .nvptx64, |
| 7 | .os_tag = .cuda, |
| 8 | }); |
| 9 | |
| 10 | { |
| 11 | var case = addPtx(ctx, target, "simple addition and subtraction"); |
| 12 | |
| 13 | case.addCompile( |
| 14 | \\fn add(a: i32, b: i32) i32 { |
| 15 | \\ return a + b; |
| 16 | \\} |
| 17 | \\ |
| 18 | \\pub export fn add_and_substract(a: i32, out: *i32) callconv(.kernel) void { |
| 19 | \\ const x = add(a, 7); |
| 20 | \\ var y = add(2, 0); |
| 21 | \\ y -= x; |
| 22 | \\ out.* = y; |
| 23 | \\} |
| 24 | ); |
| 25 | } |
| 26 | |
| 27 | { |
| 28 | var case = addPtx(ctx, target, "read special registers"); |
| 29 | |
| 30 | case.addCompile( |
| 31 | \\fn threadIdX() u32 { |
| 32 | \\ return asm ("mov.u32 \t%[r], %tid.x;" |
| 33 | \\ : [r] "=r" (-> u32), |
| 34 | \\ ); |
| 35 | \\} |
| 36 | \\ |
| 37 | \\pub export fn special_reg(a: []const i32, out: []i32) callconv(.kernel) void { |
| 38 | \\ const i = threadIdX(); |
| 39 | \\ out[i] = a[i] + 7; |
| 40 | \\} |
| 41 | ); |
| 42 | } |
| 43 | |
| 44 | { |
| 45 | var case = addPtx(ctx, target, "address spaces"); |
| 46 | |
| 47 | case.addCompile( |
| 48 | \\var x: i32 addrspace(.global) = 0; |
| 49 | \\ |
| 50 | \\pub export fn increment(out: *i32) callconv(.kernel) void { |
| 51 | \\ x += 1; |
| 52 | \\ out.* = x; |
| 53 | \\} |
| 54 | ); |
| 55 | } |
| 56 | |
| 57 | { |
| 58 | var case = addPtx(ctx, target, "reduce in shared mem"); |
| 59 | case.addCompile( |
| 60 | \\fn threadIdX() u32 { |
| 61 | \\ return asm ("mov.u32 \t%[r], %tid.x;" |
| 62 | \\ : [r] "=r" (-> u32), |
| 63 | \\ ); |
| 64 | \\} |
| 65 | \\ |
| 66 | \\ var _sdata: [1024]f32 addrspace(.shared) = undefined; |
| 67 | \\ pub export fn reduceSum(d_x: []const f32, out: *f32) callconv(.kernel) void { |
| 68 | \\ var sdata: *addrspace(.generic) [1024]f32 = @addrSpaceCast(&_sdata); |
| 69 | \\ const tid: u32 = threadIdX(); |
| 70 | \\ var sum = d_x[tid]; |
| 71 | \\ sdata[tid] = sum; |
| 72 | \\ asm volatile ("bar.sync \t0;"); |
| 73 | \\ var s: u32 = 512; |
| 74 | \\ while (s > 0) : (s = s >> 1) { |
| 75 | \\ if (tid < s) { |
| 76 | \\ sum += sdata[tid + s]; |
| 77 | \\ sdata[tid] = sum; |
| 78 | \\ } |
| 79 | \\ asm volatile ("bar.sync \t0;"); |
| 80 | \\ } |
| 81 | \\ |
| 82 | \\ if (tid == 0) { |
| 83 | \\ out.* = sum; |
| 84 | \\ } |
| 85 | \\ } |
| 86 | ); |
| 87 | } |
| 88 | } |
| 89 | |
| 90 | fn addPtx(ctx: *Cases, target: std.Build.ResolvedTarget, name: []const u8) *Cases.Case { |
| 91 | ctx.cases.append(.{ |
| 92 | .name = name, |
| 93 | .target = target, |
| 94 | .files = .init(ctx.arena), |
| 95 | .case = null, |
| 96 | .output_mode = .Obj, |
| 97 | .deps = .init(ctx.arena), |
| 98 | .link_libc = false, |
| 99 | .emit_bin = false, |
| 100 | .backend = .llvm, |
| 101 | // Bug in Debug mode |
| 102 | .optimize_mode = .ReleaseSafe, |
| 103 | }) catch @panic("out of memory"); |
| 104 | return &ctx.cases.items[ctx.cases.items.len - 1]; |
| 105 | } |