`nx_ptx_emit.nx` source

1// nx_ptx_emit.nx -- the SEED of the sovereign PTX emitter: a NishiLang program (compiled by nx_cc) that GENERATES
2// dispatchable PTX text for an elementwise integer kernel, parameterized by OP. Not hand-written, not nvcc -- OUR
3// stack emits the GPU kernel. Writes emitted_vadd.ptx; the driver harness then launches it on the RTX 5080.
4// Proves nx_cc -> (program that emits) PTX -> GPU. Grow this into the nx_cc backend proper. license_tier: ORIGINAL
5import "nx_syscalls.nx"
6
7func w(fd: i64, s: *u8) -> i64 { var n: i64=0; while s[n]!=(0 as u8){n=n+1} sys_write(fd,s,n); return 0 }
8
9const OP: i64 = 0   // 0=add.s32  1=mul.lo.s32  2=max.s32
10
11func emit_head(fd: i64) -> i64 {
12    w(fd, ".version 7.0\n.target sm_70\n.address_size 64\n\n" as *u8)
13    w(fd, ".visible .entry vadd(\n    .param .u64 vadd_a,\n    .param .u64 vadd_b,\n    .param .u64 vadd_c,\n    .param .u32 vadd_n\n)\n{\n" as *u8)
14    w(fd, "    .reg .pred  %p<2>;\n    .reg .b32   %r<9>;\n    .reg .b64   %rd<11>;\n" as *u8)
15    return 0
16}
17func emit_index(fd: i64) -> i64 {
18    w(fd, "    ld.param.u64  %rd1, [vadd_a];\n    ld.param.u64  %rd2, [vadd_b];\n    ld.param.u64  %rd3, [vadd_c];\n    ld.param.u32  %r1,  [vadd_n];\n" as *u8)
19    w(fd, "    mov.u32   %r2, %ntid.x;\n    mov.u32   %r3, %ctaid.x;\n    mov.u32   %r4, %tid.x;\n    mad.lo.s32 %r5, %r3, %r2, %r4;\n" as *u8)
20    w(fd, "    setp.ge.u32 %p1, %r5, %r1;\n    @%p1 bra $L_END;\n" as *u8)
21    return 0
22}
23func emit_load(fd: i64) -> i64 {
24    w(fd, "    cvta.to.global.u64 %rd4, %rd1;\n    cvta.to.global.u64 %rd5, %rd2;\n    cvta.to.global.u64 %rd6, %rd3;\n" as *u8)
25    w(fd, "    mul.wide.u32 %rd7, %r5, 4;\n    add.s64  %rd8,  %rd4, %rd7;\n    add.s64  %rd9,  %rd5, %rd7;\n    add.s64  %rd10, %rd6, %rd7;\n" as *u8)
26    w(fd, "    ld.global.u32 %r6, [%rd8];\n    ld.global.u32 %r7, [%rd9];\n" as *u8)
27    return 0
28}
29func emit_op(fd: i64) -> i64 {
30    if OP==0 { w(fd, "    add.s32  %r8, %r6, %r7;\n" as *u8) }
31    else { if OP==1 { w(fd, "    mul.lo.s32  %r8, %r6, %r7;\n" as *u8) }
32    else { w(fd, "    max.s32  %r8, %r6, %r7;\n" as *u8) } }
33    return 0
34}
35func emit_tail(fd: i64) -> i64 {
36    w(fd, "    st.global.u32 [%rd10], %r8;\n$L_END:\n    ret;\n}\n" as *u8)
37    return 0
38}
39
40// --- GEMM emitter: a REDUCTION-LOOP integer-MAC kernel (the matmul core), integer-exact/deterministic ---
41func emit_gemm_head(fd: i64) -> i64 {
42    w(fd, ".version 7.0\n.target sm_70\n.address_size 64\n\n" as *u8)
43    w(fd, ".visible .entry gemm(\n    .param .u64 g_a, .param .u64 g_b, .param .u64 g_c,\n    .param .u32 g_m, .param .u32 g_n, .param .u32 g_k\n)\n{\n" as *u8)
44    w(fd, "    .reg .pred %p<3>;\n    .reg .b32  %r<20>;\n    .reg .b64  %rd<14>;\n" as *u8)
45    return 0
46}
47func emit_gemm_idx(fd: i64) -> i64 {
48    w(fd, "    ld.param.u64 %rd1, [g_a];\n    ld.param.u64 %rd2, [g_b];\n    ld.param.u64 %rd3, [g_c];\n" as *u8)
49    w(fd, "    ld.param.u32 %r1, [g_m];\n    ld.param.u32 %r2, [g_n];\n    ld.param.u32 %r3, [g_k];\n" as *u8)
50    w(fd, "    mov.u32 %r4, %ntid.x; mov.u32 %r5, %ctaid.x; mov.u32 %r6, %tid.x;\n    mad.lo.s32 %r7, %r5, %r4, %r6;\n" as *u8)
51    w(fd, "    mov.u32 %r8, %ntid.y; mov.u32 %r9, %ctaid.y; mov.u32 %r10, %tid.y;\n    mad.lo.s32 %r11, %r9, %r8, %r10;\n" as *u8)
52    return 0
53}
54func emit_gemm_bounds(fd: i64) -> i64 {
55    w(fd, "    setp.ge.u32 %p1, %r7, %r2;\n    setp.ge.u32 %p2, %r11, %r1;\n    or.pred %p1, %p1, %p2;\n    @%p1 bra $DONE;\n" as *u8)
56    w(fd, "    cvta.to.global.u64 %rd4, %rd1;\n    cvta.to.global.u64 %rd5, %rd2;\n    cvta.to.global.u64 %rd6, %rd3;\n" as *u8)
57    w(fd, "    mov.s32 %r12, 0;\n    mov.s32 %r13, 0;\n$LOOP:\n    setp.ge.u32 %p1, %r13, %r3;\n    @%p1 bra $ENDLOOP;\n" as *u8)
58    return 0
59}
60func emit_gemm_body(fd: i64) -> i64 {
61    w(fd, "    mad.lo.s32 %r14, %r11, %r3, %r13;\n    mul.wide.s32 %rd7, %r14, 4;\n    add.s64 %rd8, %rd4, %rd7;\n    ld.global.u32 %r15, [%rd8];\n" as *u8)
62    w(fd, "    mad.lo.s32 %r16, %r13, %r2, %r7;\n    mul.wide.s32 %rd9, %r16, 4;\n    add.s64 %rd10, %rd5, %rd9;\n    ld.global.u32 %r17, [%rd10];\n" as *u8)
63    w(fd, "    mad.lo.s32 %r12, %r15, %r17, %r12;\n    add.s32 %r13, %r13, 1;\n    bra $LOOP;\n" as *u8)
64    w(fd, "$ENDLOOP:\n    mad.lo.s32 %r18, %r11, %r2, %r7;\n    mul.wide.s32 %rd11, %r18, 4;\n    add.s64 %rd12, %rd6, %rd11;\n    st.global.u32 [%rd12], %r12;\n$DONE:\n    ret;\n}\n" as *u8)
65    return 0
66}
67
68func main() -> i64 {
69    let path: *u8 = "/mnt/c/Users/elder/AppData/Local/Temp/claude/C--Users-elder/10730ed9-0e7d-4843-ad8a-e2d4a8122cbd/scratchpad/emitted_gemm.ptx" as *u8
70    let fd: i64 = sys_openat_wr(path)
71    if fd < 0 { w(1, "NX-PTX-EMIT: FAILED to open output file\n" as *u8); return 1 }
72    emit_gemm_head(fd); emit_gemm_idx(fd); emit_gemm_bounds(fd); emit_gemm_body(fd)
73    sys_close(fd)
74    w(1, "NX-PTX-EMIT: nx_cc-compiled emitter WROTE emitted_gemm.ptx -- a REDUCTION-LOOP integer-MAC matmul kernel.\n" as *u8)
75    w(1, "  The sovereign stack now generates a MATMUL GPU kernel (loops+reduction+integer MAC), not just elementwise. Driver launches it next.\n" as *u8)
76    return 0
77}
nx_ptx_emit.nx source

`nx_ptx_emit.nx` source
