---
title: "Omega Lab: write a kernel, run it in one click | AienOS"
description: "A data-parallel kernel in Omega beside the same kernel in CUDA C. Run it in your browser as WebAssembly, checked against JavaScript, compile your edit on my DGX Spark, and see every back end and every GPU result that has and has not run."
last_updated: "2026-10-26"
canonical_url: "https://www.aienos.com/lab/"
---

01 / RUN IT

## Same kernel. Two languages.

Omega axpy.om · editable [GitHub](https://github.com/aien-dev/aienos.com/blob/main/om/lab-page/axpy.om)
note y[i] = a*x[i] + y[i]: one element, the work of one CUDA thread.
fn axpy a long x long y long to long
ret add mul a x y
end
note The grid: elements lo to hi-1, with x[i] = i and y[i] = 2i+1 generated in place.
note Returns the sum of the results so a host can check every element.
fn run a long lo long hi long to long
s = long 0
i = lo
while lt i hi bound 999999
s = add s axpy a i add mul long 2 i long 1
i = add i long 1
end
ret s
end

CUDA C axpy.cu · for comparison [GitHub](https://github.com/aien-dev/aienos.com/blob/main/om/lab-page/axpy.cu)

// The same kernel in CUDA C: y[i] = a*x[i] + y[i], one element per thread.
#include <cstdio>
#include <cuda_runtime.h>

__global__ void axpy(long long a, const long long *x, long long *y, long long n) {
long long i = (long long)blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) y[i] = a * x[i] + y[i];
}

int main() {
const long long n = 1 << 20, a = 3;
long long *x, *y;
cudaMallocManaged(&x, n * sizeof *x);
cudaMallocManaged(&y, n * sizeof *y);
for (long long i = 0; i < n; i++) { x[i] = i; y[i] = 2 * i + 1; }
axpy<<<(n + 255) / 256, 256>>>(a, x, y, n);
cudaDeviceSynchronize();
long long s = 0;
for (long long i = 0; i < n; i++) s += y[i];
printf("%lld\n", s);
cudaFree(x);
cudaFree(y);
return 0;
}

a
elements 1,02465,5361,048,5764,194,304
Run in my browser
Compile my edit
Break it on purpose
Reset

Loading the prebuilt kernel.

Omega kernel, WebAssembly, this tab not run yet

Same math in plain JavaScript not run yet

Kernel checksum not run yet

JavaScript checksum not run yet

Kernel throughput not run yet

Compile on my Spark press Compile

Round trip from this tab press Compile

Compiler pins press Compile

How the check works: for every element the kernel computes y[i] = a·x[i] + y[i] with x[i] = i and y[i] = 2i + 1, and returns the sum of all results. JavaScript computes the same sum on its own. The page calls run in blocks of at most 999,999 elements, the way a CUDA grid splits work into blocks. Times are the median of three runs in a worker on your device. Integers only for now: the pinned WebAssembly path does not take f32 yet. Omega traps on integer overflow where CUDA C wraps silently. Break it on purpose changes add to sub so you can watch the check catch it.

02 / ONE SOURCE, EVERY BACK END

## What my compiler actually emitted.

Recorded on my DGX Spark by record.sh, except the WebAssembly, which is whatever this tab is running right now.

WebAssembly runs in this tab

Loading.

[axpy.wasm on GitHub](https://github.com/aien-dev/aienos.com/blob/main/public/lab/axpy.wasm)

AArch64 ran on the Spark CPU

compiler omc sha256 734f0a39319fbfa930bcc62cb30dbf06d47d78e126372b74f5baf478e6126b61 (the site pin)
command omc < axpy.om > axpy.bin
compile 1.455 ms, median of 7, Linux aarch64, 2 of 20 CPU cores allowed, DGX Spark (GB10), load1 7.90, 2026-10-27T00:42:59Z
output 276 bytes of raw AArch64 code, sha256 7bdd7e87bb3c4c79897509c7b22f63b2e738195301e9afa665a40387172d2bb7
program axpy.om + harness.om into one Linux program: 1.881 ms, median of 7, load1 7.90
run a=3, n=1048576: checksum 2748777496576, awk reference 2748777496576, MATCH
control add changed to sub in the kernel: 549754241024, MISMATCH as it should be
run time 4.345 ms per process (start, 1048576 elements, print, exit), median of 7, load1 7.90
disasm objdump -D -b binary -m aarch64 -M no-aliases (GNU objdump (GNU Binutils for Ubuntu) 2.42)

0: 1400000d b 0x34
4: 9b417c04 smulh x4, x0, x1
8: 9b017c03 madd x3, x0, x1, xzr
c: eb83fc9f subs xzr, x4, x3, asr #63
10: 540000e1 b.ne 0x2c // b.any
14: ab020060 adds x0, x3, x2
18: 540000a6 b.vs 0x2c
1c: d65f03c0 ret
20: aa0703e0 orr x0, xzr, x7
24: f94008f0 ldr x16, [x7, #16]
28: d61f0200 br x16
2c: d2800021 movz x1, #0x1
30: 17fffffc b 0x20
34: d2800003 movz x3, #0x0
38: aa0103e4 orr x4, xzr, x1
3c: d2800005 movz x5, #0x0
40: eb02009f subs xzr, x4, x2
44: 5400056a b.ge 0xf0 // b.tcont
48: d28847e6 movz x6, #0x423f
4c: f2a001e6 movk x6, #0xf, lsl #16
50: eb0600bf subs xzr, x5, x6
54: 540005c2 b.cs 0x10c // b.hs, b.nlast
58: 910004a5 add x5, x5, #0x1
5c: d2800046 movz x6, #0x2
60: 9b447cc9 smulh x9, x6, x4
64: 9b047cc8 madd x8, x6, x4, xzr
68: eb88fd3f subs xzr, x9, x8, asr #63
6c: 540004c1 b.ne 0x104 // b.any
70: b100050a adds x10, x8, #0x1
74: 54000486 b.vs 0x104
78: d10183ff sub sp, sp, #0x60
7c: f90003e0 str x0, [sp]
80: f90007e1 str x1, [sp, #8]
84: f9000be2 str x2, [sp, #16]
88: f9000fe3 str x3, [sp, #24]
8c: f90013e4 str x4, [sp, #32]
90: f90017e5 str x5, [sp, #40]
94: f9001be6 str x6, [sp, #48]
98: f9001fe8 str x8, [sp, #56]
9c: f90023e9 str x9, [sp, #64]
a0: f90027ea str x10, [sp, #72]
a4: f9002bfe str x30, [sp, #80]
a8: f94003e0 ldr x0, [sp]
ac: f94013e1 ldr x1, [sp, #32]
b0: f94027e2 ldr x2, [sp, #72]
b4: 97ffffd4 bl 0x4
b8: aa0003e6 orr x6, xzr, x0
bc: f94003e0 ldr x0, [sp]
c0: f94007e1 ldr x1, [sp, #8]
c4: f9400be2 ldr x2, [sp, #16]
c8: f9400fe3 ldr x3, [sp, #24]
cc: f94013e4 ldr x4, [sp, #32]
d0: f94017e5 ldr x5, [sp, #40]
d4: f9402bfe ldr x30, [sp, #80]
d8: 910183ff add sp, sp, #0x60
dc: ab060063 adds x3, x3, x6
e0: 54000126 b.vs 0x104
e4: b1000484 adds x4, x4, #0x1
e8: 540000e6 b.vs 0x104
ec: 17ffffd5 b 0x40
f0: aa0303e0 orr x0, xzr, x3
f4: d65f03c0 ret
f8: aa0703e0 orr x0, xzr, x7
fc: f94008f0 ldr x16, [x7, #16]
100: d61f0200 br x16
104: d2800021 movz x1, #0x1
108: 17fffffc b 0xf8
10c: d2800081 movz x1, #0x4
110: 17fffffa b 0xf8

[aarch64.txt on GitHub](https://github.com/aien-dev/aienos.com/blob/main/om/lab-page/captures/aarch64.txt)

SPIR-V validated, ran on PoCL (CPU) GPU: NOT_RUN

compiler omega-mv 369dbc96daa1684db2d7c2c7085bbf16b3fd77ee: omc sha256 2c82480596bbdc50cda9626866209ef662d5834f278f839ff4ce0250c495029f builds irspirv
command { printf 'ir\n'; cat axpy.om; } | omc > axpy.ir && irspirv < axpy.ir > axpy.spv
translate 1.625 ms for irspirv, median of 7, Linux aarch64, 2 of 20 CPU cores allowed, DGX Spark (GB10), load1 7.90
output 4140 bytes, sha256 adeb8af7af76a547329eb9d91d839032a66c1d70f0a9b3fd0f78000bd77d8f48
validate spirv-val --target-env opencl2.1: valid (SPIRV-Tools v2025.1 unknown hash, 2026-05-13T06:03:38+00:00)
run PoCL 7.2 on the CPU (cpu-cortex-a720-0xd87): run 3 0 1000 = 2498500, trap 0; reference 2498500, MATCH
not run no GPU has executed this SPIR-V yet

; SPIR-V
; Version: 1.0
; Generator: Khronos; 0
; Bound: 212
; Schema: 0
OpCapability Addresses
OpCapability Kernel
OpCapability Int64
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %33 "axpy"
OpEntryPoint Kernel %35 "run"
%void = OpTypeVoid
%bool = OpTypeBool
%ulong = OpTypeInt 64 0
%_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
%6 = OpTypeFunction %void %_ptr_CrossWorkgroup_ulong %_ptr_CrossWorkgroup_ulong
%7 = OpTypeFunction %ulong %_ptr_Function_ulong
%8 = OpTypeFunction %ulong %_ptr_Function_ulong %ulong
%9 = OpTypeFunction %ulong %_ptr_Function_ulong %ulong %ulong
%10 = OpTypeFunction %ulong %_ptr_Function_ulong %ulong %ulong %ulong
%11 = OpTypeFunction %ulong %_ptr_Function_ulong %ulong %ulong %ulong %ulong
%12 = OpTypeFunction %ulong %_ptr_Function_ulong %ulong %ulong %ulong %ulong %ulong
%13 = OpTypeFunction %ulong %_ptr_Function_ulong %ulong %ulong %ulong %ulong %ulong %ulong
%ulong_0 = OpConstant %ulong 0
%ulong_9223372036854775808 = OpConstant %ulong 9223372036854775808
%ulong_9223372036854775807 = OpConstant %ulong 9223372036854775807
%ulong_1 = OpConstant %ulong 1
%ulong_101 = OpConstant %ulong 101
%ulong_2 = OpConstant %ulong 2
%ulong_999999 = OpConstant %ulong 999999
%ulong_104 = OpConstant %ulong 104
%32 = OpFunction %ulong None %10
%36 = OpFunctionParameter %_ptr_Function_ulong
%37 = OpFunctionParameter %ulong
%38 = OpFunctionParameter %ulong
%39 = OpFunctionParameter %ulong
%40 = OpLabel
%41 = OpVariable %_ptr_Function_ulong Function
%42 = OpVariable %_ptr_Function_ulong Function
%43 = OpVariable %_ptr_Function_ulong Function
%44 = OpVariable %_ptr_Function_ulong Function
%45 = OpVariable %_ptr_Function_ulong Function
OpStore %41 %37
OpStore %42 %38
OpStore %43 %39
OpBranch %46
%46 = OpLabel
%47 = OpLoad %ulong %41
%48 = OpLoad %ulong %42
%50 = OpSLessThan %bool %47 %ulong_0
%51 = OpSLessThan %bool %48 %ulong_0
%52 = OpISub %ulong %ulong_0 %47
%53 = OpSelect %ulong %50 %52 %47
%54 = OpISub %ulong %ulong_0 %48
%55 = OpSelect %ulong %51 %54 %48
%56 = OpLogicalNotEqual %bool %50 %51
%59 = OpSelect %ulong %56 %ulong_9223372036854775808 %ulong_9223372036854775807
%60 = OpIEqual %bool %55 %ulong_0
%62 = OpSelect %ulong %60 %ulong_1 %55
%63 = OpUDiv %ulong %59 %62
%64 = OpUGreaterThan %bool %53 %63
%65 = OpINotEqual %bool %55 %ulong_0
%66 = OpLogicalAnd %bool %64 %65
OpBranchConditional %66 %67 %68
%67 = OpLabel
OpStore %36 %ulong_101
OpReturnValue %ulong_0
%68 = OpLabel
%70 = OpIMul %ulong %47 %48
OpStore %44 %70
%71 = OpLoad %ulong %44
%72 = OpLoad %ulong %43
%73 = OpIAdd %ulong %71 %72
%74 = OpBitwiseXor %ulong %71 %73
%75 = OpBitwiseXor %ulong %72 %73
%76 = OpBitwiseAnd %ulong %74 %75
%77 = OpSLessThan %bool %76 %ulong_0
OpBranchConditional %77 %78 %79
%78 = OpLabel
OpStore %36 %ulong_101
OpReturnValue %ulong_0
%79 = OpLabel
OpStore %45 %73
%80 = OpLoad %ulong %45
OpReturnValue %80
OpFunctionEnd
%33 = OpFunction %void None %6
%81 = OpFunctionParameter %_ptr_CrossWorkgroup_ulong
%82 = OpFunctionParameter %_ptr_CrossWorkgroup_ulong
%83 = OpLabel
%84 = OpVariable %_ptr_Function_ulong Function
OpStore %84 %ulong_0
%85 = OpPtrAccessChain %_ptr_CrossWorkgroup_ulong %81 %ulong_0
%86 = OpLoad %ulong %85
%87 = OpPtrAccessChain %_ptr_CrossWorkgroup_ulong %81 %ulong_1
%88 = OpLoad %ulong %87
%89 = OpPtrAccessChain %_ptr_CrossWorkgroup_ulong %81 %ulong_2
%91 = OpLoad %ulong %89
%92 = OpFunctionCall %ulong %32 %84 %86 %88 %91
OpStore %82 %92
%93 = OpPtrAccessChain %_ptr_CrossWorkgroup_ulong %82 %ulong_1
%94 = OpLoad %ulong %84
OpStore %93 %94
OpReturn
OpFunctionEnd
%34 = OpFunction %ulong None %10
%95 = OpFunctionParameter %_ptr_Function_ulong
%96 = OpFunctionParameter %ulong
%97 = OpFunctionParameter %ulong
%98 = OpFunctionParameter %ulong
%99 = OpLabel
%100 = OpVariable %_ptr_Function_ulong Function
%101 = OpVariable %_ptr_Function_ulong Function
%102 = OpVariable %_ptr_Function_ulong Function
%103 = OpVariable %_ptr_Function_ulong Function
%104 = OpVariable %_ptr_Function_ulong Function
%105 = OpVariable %_ptr_Function_ulong Function
%106 = OpVariable %_ptr_Function_ulong Function
%107 = OpVariable %_ptr_Function_ulong Function
%108 = OpVariable %_ptr_Function_ulong Function
%109 = OpVariable %_ptr_Function_ulong Function
%110 = OpVariable %_ptr_Function_ulong Function
%111 = OpVariable %_ptr_Function_ulong Function
%112 = OpVariable %_ptr_Function_ulong Function
%113 = OpVariable %_ptr_Function_ulong Function
%114 = OpVariable %_ptr_Function_ulong Function
%115 = OpVariable %_ptr_Function_ulong Function
%116 = OpVariable %_ptr_Function_ulong Function
OpStore %100 %96
OpStore %101 %97
OpStore %102 %98
OpBranch %117
%117 = OpLabel
OpStore %103 %ulong_0
OpStore %106 %ulong_0
OpStore %107 %ulong_1
OpStore %110 %ulong_999999
OpStore %112 %ulong_2
OpStore %114 %ulong_1
%124 = OpLoad %ulong %103
OpStore %104 %124
%125 = OpLoad %ulong %101
OpStore %105 %125
%126 = OpLoad %ulong %106
OpStore %108 %126
OpBranch %118
%118 = OpLabel
%127 = OpLoad %ulong %105
%128 = OpLoad %ulong %102
%129 = OpSLessThan %bool %127 %128
%130 = OpSelect %ulong %129 %ulong_1 %ulong_0
OpStore %109 %130
%131 = OpLoad %ulong %109
%132 = OpINotEqual %bool %131 %ulong_0
OpBranchConditional %132 %119 %122
%119 = OpLabel
%133 = OpLoad %ulong %108
%134 = OpLoad %ulong %110
%135 = OpUGreaterThanEqual %bool %133 %134
%136 = OpSelect %ulong %135 %ulong_1 %ulong_0
OpStore %111 %136
%137 = OpLoad %ulong %111
%138 = OpINotEqual %bool %137 %ulong_0
OpBranchConditional %138 %120 %121
%121 = OpLabel
%139 = OpLoad %ulong %108
%140 = OpLoad %ulong %107
%141 = OpIAdd %ulong %139 %140
OpStore %108 %141
%142 = OpLoad %ulong %112
%143 = OpLoad %ulong %105
%144 = OpSLessThan %bool %142 %ulong_0
%145 = OpSLessThan %bool %143 %ulong_0
%146 = OpISub %ulong %ulong_0 %142
%147 = OpSelect %ulong %144 %146 %142
%148 = OpISub %ulong %ulong_0 %143
%149 = OpSelect %ulong %145 %148 %143
%150 = OpLogicalNotEqual %bool %144 %145
%151 = OpSelect %ulong %150 %ulong_9223372036854775808 %ulong_9223372036854775807
%152 = OpIEqual %bool %149 %ulong_0
%153 = OpSelect %ulong %152 %ulong_1 %149
%154 = OpUDiv %ulong %151 %153
%155 = OpUGreaterThan %bool %147 %154
%156 = OpINotEqual %bool %149 %ulong_0
%157 = OpLogicalAnd %bool %155 %156
OpBranchConditional %157 %158 %159
%158 = OpLabel
OpStore %95 %ulong_101
OpReturnValue %ulong_0
%159 = OpLabel
%160 = OpIMul %ulong %142 %143
OpStore %113 %160
%161 = OpLoad %ulong %113
%162 = OpLoad %ulong %114
%163 = OpIAdd %ulong %161 %162
%164 = OpBitwiseXor %ulong %161 %163
%165 = OpBitwiseXor %ulong %162 %163
%166 = OpBitwiseAnd %ulong %164 %165
%167 = OpSLessThan %bool %166 %ulong_0
OpBranchConditional %167 %168 %169
%168 = OpLabel
OpStore %95 %ulong_101
OpReturnValue %ulong_0
%169 = OpLabel
OpStore %115 %163
%170 = OpLoad %ulong %100
%171 = OpLoad %ulong %105
%172 = OpLoad %ulong %115
%173 = OpFunctionCall %ulong %32 %95 %170 %171 %172
%174 = OpLoad %ulong %95
%177 = OpINotEqual %bool %174 %ulong_0
OpBranchConditional %177 %175 %176
%175 = OpLabel
OpReturnValue %ulong_0
%176 = OpLabel
OpStore %116 %173
%178 = OpLoad %ulong %104
%179 = OpLoad %ulong %116
%180 = OpIAdd %ulong %178 %179
%181 = OpBitwiseXor %ulong %178 %180
%182 = OpBitwiseXor %ulong %179 %180
%183 = OpBitwiseAnd %ulong %181 %182
%184 = OpSLessThan %bool %183 %ulong_0
OpBranchConditional %184 %185 %186
%185 = OpLabel
OpStore %95 %ulong_101
OpReturnValue %ulong_0
%186 = OpLabel
OpStore %104 %180
%187 = OpLoad %ulong %105
%188 = OpLoad %ulong %114
%189 = OpIAdd %ulong %187 %188
%190 = OpBitwiseXor %ulong %187 %189
%191 = OpBitwiseXor %ulong %188 %189
%192 = OpBitwiseAnd %ulong %190 %191
%193 = OpSLessThan %bool %192 %ulong_0
OpBranchConditional %193 %194 %195
%194 = OpLabel
OpStore %95 %ulong_101
OpReturnValue %ulong_0
%195 = OpLabel
OpStore %105 %189
OpBranch %118
%120 = OpLabel
OpStore %95 %ulong_104
OpReturnValue %ulong_0
%197 = OpLabel
OpBranch %122
%122 = OpLabel
%198 = OpLoad %ulong %104
OpReturnValue %198
OpFunctionEnd
%35 = OpFunction %void None %6
%199 = OpFunctionParameter %_ptr_CrossWorkgroup_ulong
%200 = OpFunctionParameter %_ptr_CrossWorkgroup_ulong
%201 = OpLabel
%202 = OpVariable %_ptr_Function_ulong Function
OpStore %202 %ulong_0
%203 = OpPtrAccessChain %_ptr_CrossWorkgroup_ulong %199 %ulong_0
%204 = OpLoad %ulong %203
%205 = OpPtrAccessChain %_ptr_CrossWorkgroup_ulong %199 %ulong_1
%206 = OpLoad %ulong %205
%207 = OpPtrAccessChain %_ptr_CrossWorkgroup_ulong %199 %ulong_2
%208 = OpLoad %ulong %207
%209 = OpFunctionCall %ulong %34 %202 %204 %206 %208
OpStore %200 %209
%210 = OpPtrAccessChain %_ptr_CrossWorkgroup_ulong %200 %ulong_1
%211 = OpLoad %ulong %202
OpStore %210 %211
OpReturn
OpFunctionEnd

[spirv.txt on GitHub](https://github.com/aien-dev/aienos.com/blob/main/om/lab-page/captures/spirv.txt)

NVIDIA SASS this kernel: NOT_RUN

omc does not lower this kernel to SASS yet. Below are two real listings for comparison: what nvcc emits for the CUDA C version (compiled here, not launched), then the vector add that did run on the GB10, written in asm2 and launched without any CUDA library.

compiler nvcc 13.0 (V13.0.88)
command nvcc -arch=sm_121 -o axpy axpy.cu (host + device, a runnable program)
compile 593.800 ms, median of 7, Linux aarch64, 2 of 20 CPU cores allowed, DGX Spark (GB10), load1 8.15
device nvcc -arch=sm_121 -cubin: 192.561 ms, median of 7
not run this CUDA program has not been launched; that needs a GPU window
disasm cuobjdump -sass axpy.cubin

/*0000*/ LDC R1, c[0x0][0x37c] ;
/*0010*/ S2R R2, SR_TID.X ;
/*0020*/ S2UR UR4, SR_CTAID.X ;
/*0030*/ LDCU.64 UR6, c[0x0][0x398] ;
/*0040*/ HFMA2 R3, -RZ, RZ, 0, 0 ;
/*0050*/ LDC R5, c[0x0][0x360] ;
/*0060*/ IMAD.WIDE.U32 R2, R5, UR4, R2 ;
/*0070*/ ISETP.GE.S64.AND P0, PT, R2, UR6, PT ;
/*0080*/ @P0 EXIT ;
/*0090*/ LDCU.64 UR6, c[0x0][0x388] ;
/*00a0*/ SHF.L.U64.HI R3, R2.reuse, 0x3, R3 ;
/*00b0*/ IMAD.SHL.U32 R2, R2, 0x8, RZ ;
/*00c0*/ LDCU.64 UR8, c[0x0][0x390] ;
/*00d0*/ LDCU.64 UR4, c[0x0][0x358] ;
/*00e0*/ IADD.64 R4, R2.reuse, UR6 ;
/*00f0*/ LDCU.64 UR6, c[0x0][0x380] ;
/*0100*/ IADD.64 R2, R2, UR8 ;
/*0110*/ LDG.E.64 R4, desc[UR4][R4.64] ;
/*0120*/ LDG.E.64 R6, desc[UR4][R2.64] ;
/*0130*/ IMAD R9, R5, UR6, RZ ;
/*0140*/ IMAD.WIDE.U32 R6, R4, UR6, R6 ;
/*0150*/ IMAD R9, R4, UR7, R9 ;
/*0160*/ IADD3 R7, PT, PT, R7, R9, RZ ;
/*0170*/ STG.E.64 desc[UR4][R2.64], R6 ;
/*0180*/ EXIT ;
/*0190*/ BRA 0x190;
/*01a0*/ NOP;
/*01b0*/ NOP;
/*01c0*/ NOP;
/*01d0*/ NOP;
/*01e0*/ NOP;
/*01f0*/ NOP;
/*0200*/ NOP;
/*0210*/ NOP;
/*0220*/ NOP;
/*0230*/ NOP;
/*0240*/ NOP;
/*0250*/ NOP;
/*0260*/ NOP;
/*0270*/ NOP;

source omega-mv seed/sm121/vecadd.s2, written in asm2 (not compiled from .om)
run GB10 (sm_121), 1048576 sums, 0 wrong, launch_to_done_ns=214257, 2026-10-19
control a mutated copy of the same kernel fails the same check
note launched by omega-mv's bytes tool with no CUDA library (seed/sm121/chip6.sh); the Spark still runs NVIDIA DGX OS and the GPU loads NVIDIA signed firmware; the run record says dirty=3 (the bytes tool, copied in)
disasm nvdisasm, as committed in evidence/SEED-0/sm121/vecadd.sass.txt

/*0000*/ LDC R1, c[0x0][0x37c] ;
/*0010*/ S2R R9, SR_CTAID.X ;
/*0020*/ LDCU UR4, c[0x0][0x360] ;
/*0030*/ S2R R0, SR_TID.X ;
/*0040*/ LDCU UR5, c[0x0][0x398] ;
/*0050*/ IMAD R9, R9, UR4, R0 ;
/*0060*/ ISETP.GE.U32.AND P0, PT, R9, UR5, PT ;
/*0070*/ @P0 EXIT ;
/*0080*/ LDC.64 R2, c[0x0][0x380] ;
/*0090*/ LDCU.64 UR4, c[0x0][0x358] ;
/*00a0*/ LDC.64 R4, c[0x0][0x388] ;
/*00b0*/ LDC.64 R6, c[0x0][0x390] ;
/*00c0*/ IMAD.WIDE.U32 R2, R9, 0x4, R2 ;
/*00d0*/ LDG.E R2, desc[UR4][R2.64] ;
/*00e0*/ IMAD.WIDE.U32 R4, R9, 0x4, R4 ;
/*00f0*/ LDG.E R5, desc[UR4][R4.64] ;
/*0100*/ IMAD.WIDE.U32 R6, R9, 0x4, R6 ;
/*0110*/ IADD3 R9, PT, PT, R2, R5, RZ ;
/*0120*/ STG.E desc[UR4][R6.64], R9 ;
/*0130*/ EXIT ;
/*0140*/ BRA 0x140;

[nvcc.txt](https://github.com/aien-dev/aienos.com/blob/main/om/lab-page/captures/nvcc.txt) · [gb10.txt](https://github.com/aien-dev/aienos.com/blob/main/om/lab-page/captures/gb10.txt)

03 / THE HONEST RACE

## Where it has run, and where it has not.

Where What ran Elements Time Check Evidence

Your device, this tab this kernel, Omega to WebAssembly press Run press Run press Run [kernel](https://github.com/aien-dev/aienos.com/blob/main/om/lab-page/axpy.om)

DGX Spark, Arm CPU this kernel, Omega to AArch64, whole process 1,048,576 see record.tsv MATCH; sub control fails [aarch64.txt](https://github.com/aien-dev/aienos.com/blob/main/om/lab-page/captures/aarch64.txt)

DGX Spark, Arm CPU this kernel, Omega to SPIR-V on PoCL 1,000 not timed MATCH [spirv.txt](https://github.com/aien-dev/aienos.com/blob/main/om/lab-page/captures/spirv.txt)

GB10 GPU vector add in asm2, no CUDA library (not this kernel) 1,048,576 214,257 ns launch to done 0 wrong [omega-mv SEED-0](https://github.com/aien-dev/omega-mv/commit/3e3da942ad630035ff773982b1aba89b36ba76d6)

GB10 GPU hello_gpu float add, kernel from my own builder (not this kernel) 1,048,576 947.5 µs GPU 0 wrong [CUDApp log](https://github.com/aien-dev/CUDApp/blob/2b4b2b14f1c67473b80878d2eff73b32b3522d54/evidence/l424/hello-chip.log)

GB10 GPU this kernel from omc NOT_RUN: omc has no SASS lowering for it yet [omega-mv](https://github.com/aien-dev/omega-mv)

GB10 GPU this kernel from nvcc NOT_RUN: built, not launched; needs a GPU window [nvcc.txt](https://github.com/aien-dev/aienos.com/blob/main/om/lab-page/captures/nvcc.txt)

Different kernels, sizes and machine load, so this is a ledger, not a ranking. The GB10 runs use NVIDIA's driver-loaded signed firmware on DGX OS; what they skip is the CUDA library. The asm2 run shared the machine at load1 61; the hello_gpu run had a quiet window.

04 / ASK FOR A KERNEL

## Let Aien write the next one.

Aien is my assistant. It runs on my GB10 and writes Omega in the chat, with the same compile and run buttons you just used.

[Talk to Aien](https://www.aienos.com/aien/) [source](https://github.com/aien-dev/aienos.com/tree/main/om/allen-page)

05 / TAKE IT HOME

## Three lines, straight from GitHub.

git clone --depth 1 https://github.com/aien-dev/aienos.com
cd aienos.com
sh om/lab-page/run.sh
Copy

Linux on 64-bit Arm runs the whole thing: it checks the compiler pins, builds the kernel, runs it, compares with awk, runs the broken control and writes the same axpy.wasm this page runs. Other machines get a plain NOT_RUN and a pointer back here. [Read run.sh first.](https://github.com/aien-dev/aienos.com/blob/main/om/lab-page/run.sh)

What it printed from a fresh clone

$ git clone --depth 1 https://github.com/aien-dev/aienos.com
(recorded from branch feat/omega-lab before merge)
$ cd aienos.com
$ sh om/lab-page/run.sh
pin omc sha256 734f0a39319fbfa930bcc62cb30dbf06d47d78e126372b74f5baf478e6126b61
pin irwasm sha256 2d235a70c39ac1461a09c41e10e058d962b767ae83893cdcfc693ef6c47182e6
build omc: axpy.om + harness.om -> out/axpy in 1048 us
run out/axpy a=3 n=1048576 -> 2748777496576 in 3779 us (whole process)
check awk reference 2748777496576
control add -> sub gives 549754241024: mismatch, so the check can fail
wasm out/axpy.wasm 519 bytes, sha256 7607811d1d4440195bd39c5bab57385dd3b81f87321dc6e83f3855b2a5fba76e
PASS
(recorded 2026-10-27T00:43:06Z in an empty directory, Linux aarch64, load1 8.21)

[home.txt on GitHub](https://github.com/aien-dev/aienos.com/blob/main/om/lab-page/captures/home.txt)

### Start your own kernel repository.

One command makes a public repository from my template. Edit kernel.om and the Python reference beside it, push, and GitHub's ubuntu-24.04-arm runner builds your kernel for AArch64 and SPIR-V, runs both (SPIR-V on PoCL, a CPU), and checks every result against the reference. No GPU runs there.

gh repo create my-kernel --template aien-dev/omega-kernel-template --public --clone

[omega-kernel-template](https://github.com/aien-dev/omega-kernel-template) · [passing check](https://github.com/aien-dev/omega-kernel-template/actions/runs/40426781699) · [failing control: an overflowing kernel must fail](https://github.com/aien-dev/omega-kernel-template/actions/runs/40426783628)

### Do not trust my compiler? Rebuild it.

The seed challenge rebuilds the Omega seed chain from 408 bytes and must land on the same fixed output. It passes on GitHub's own ubuntu-24.04-arm runner, and its tampered-rung, wrong-value and forged-log controls fail as required. From a fresh full clone of omega-mv main (52eba88) on my Spark it printed CHALLENGE PASS in 1.06 s wall. It is a fixed-output reproducibility match, not general self-hosting and not diverse double compiling. It needs the full history, so no --depth.

git clone https://github.com/aien-dev/omega-mv
cd omega-mv
sh seed/challenge/run.sh

[CHALLENGE.md](https://github.com/aien-dev/omega-mv/blob/main/seed/challenge/CHALLENGE.md) · [passing run on GitHub](https://github.com/aien-dev/omega-seed-challenge/actions/runs/40426178383) · [post your result](https://github.com/aien-dev/omega-seed-challenge/issues/new?template=reproduced.yml)

### Have an NVIDIA GPU? Put it on the board.

The bench kit loads GPU binaries my encoder wrote (sm_80, sm_86, sm_89) and checks every element. It loads them through NVIDIA's own libcuda, as an oracle for my machine code, so this path is not CUDA++ without NVIDIA user space. The board starts with three rented machines (RTX 4090, A100, RTX 3090), not timed. The kit's own run on real hardware: NOT_RUN so far. Yours would be the first.

git clone https://github.com/aien-dev/omega-gpu-bench
cd omega-gpu-bench
./run

[omega-gpu-bench](https://github.com/aien-dev/omega-gpu-bench) · [the board](https://aien-dev.github.io/omega-gpu-bench/) · [receipts.json](https://aien-dev.github.io/omega-gpu-bench/receipts.json)

Page structure and script generated by seeded Omega: [om/lab-page](https://github.com/aien-dev/aienos.com/tree/main/om/lab-page).

## Sitemap

See the full [sitemap](https://www.aienos.com/sitemap.md) for every page on aienos.com.