/**
* Warranty & Liability
* To the extent permitted by applicable law and unless explicitly
* otherwise agreed upon, XLOG Technologies AG makes no warranties
* regarding the provided information. XLOG Technologies AG assumes
* no liability that any problems might be solved with the information
* provided by XLOG Technologies AG.
*
* Rights & License
* All industrial property rights regarding the information - copyright
* and patent rights in particular - are the sole property of XLOG
* Technologies AG. If the company was not the originator of some
* excerpts, XLOG Technologies AG has at least obtained the right to
* reproduce, change and translate the information.
*
* Reproduction is restricted to the whole unaltered document. Reproduction
* of the information is only allowed for non-commercial uses. Selling,
* giving away or letting of the execution of the library is prohibited.
* The library can be distributed as part of your applications and libraries
* for execution provided this comment remains unchanged.
*
* Restrictions
* Only to be distributed with programs that add significant and primary
* functionality to the library. Not to be distributed with additional
* software intended to replace any components of the library.
*
* Trademarks
* Jekejeke is a registered trademark of XLOG Technologies AG.
*/
import {
fs
} from "../../nova/core.mjs";
export let GROUP_SIZE = 32;
export let device;
export let pipeline;
if (fs !== undefined) {
let webgpu = await import("webgpu");
Object.assign(globalThis, webgpu.globals);
globalThis.navigator.gpu = webgpu.create([]);
}
async function init_device() {
let adapter = await navigator.gpu.requestAdapter();
return await adapter.requestDevice();
}
function init_pipeline() {
let bootstrap = `
fn hack_get(instr: i32, base : i32, pc : i32) -> i32 {
var act : i32 = instr >> 20;
var obj : i32 = ((instr >> 10) & 0x03FF) - 512;
switch ((act >> 4) & 0x0F) {
case 0: { /* get */
return state[base+obj];
}
case 1: { /* cst */
return obj;
}
case 2: { /* tbl */
return code[pc+1+obj];
}
case 3: { /* set */
return 0;
}
case 4: { /* out */
return 0;
}
case 5: { /* in */
/* t.b.d. */
return 0;
}
default: {
// illegal mode - no exception throwing in WGSL
return 0;
}
}
}
fn hack_fun(instr: i32, accu: i32, value: i32) -> i32 {
var act : i32 = instr >> 20;
switch (act >> 8) {
case 0: { /* val */
return value;
}
case 1: { /* acc */
return accu;
}
case 2: { /* add */
return value + accu;
}
case 3: { /* sub */
return value - accu;
}
case 4: { /* mul */
return value * accu;
}
case 5: { /* zdiv */
return value / accu;
}
case 6: { /* rem */
return value % accu;
}
default: {
// illegal fun - no exception throwing in WGSL
return 0;
}
}
}
fn hack_set(instr: i32, accu: i32, base: i32) {
var act : i32 = instr >> 20;
var obj : i32 = ((instr >> 10) & 0x03FF) - 512;
switch ((act >> 4) & 0x0F) {
case 0: { /* get */
return;
}
case 1: { /* cst */
return;
}
case 2: { /* tbl */
return;
}
case 3: { /* set */
state[base+obj] = accu;
return;
}
case 4: { /* out */
// t.b.d.
return;
}
case 5: { /* in */
return;
}
default: {
// illegal mode - no exception throwing in WGSL
return;
}
}
}
fn hack_jump(instr: i32, accu: i32) -> i32 {
var act : i32 = instr >> 20;
var rel : i32 = (instr & 0x03FF) - 512;
switch (act & 0x0F) {
case 0: { /* true */
return rel;
}
case 1: { /* eq */
return select(0, rel, accu == 0);
}
case 2: { /* nq */
return select(0, rel, accu != 0);
}
case 3: { /* ls */
return select(0, rel, accu < 0);
}
case 4: { /* gr */
return select(0, rel, accu > 0);
}
case 5: { /* lq */
return select(0, rel, accu <= 0);
}
case 6: { /* gq */
return select(0, rel, accu >= 0);
}
default: {
return 0; // illegal cond
}
}
}
fn run_id(id : i32) {
if (!(id < i32(arrayLength(&offset)))) {
return;
}
var base : i32 = offset[id];
var pc : i32 = state[base-2];
var accu : i32 = state[base-1];
while (pc < i32(arrayLength(&code))) {
var instr : i32 = code[pc];
var value : i32 = hack_get(instr, base, pc);
accu = hack_fun(instr, accu, value);
hack_set(instr, accu, base);
pc = pc+1+hack_jump(instr, accu);
}
state[base-2] = pc;
state[base-1] = accu;
}
`
let shader = `
@group(0) @binding(0) var<storage, read> code: array<i32>;
@group(0) @binding(1) var<storage, read> offset: array<i32>;
@group(0) @binding(2) var<storage, read_write> state: array<i32>;
`
shader += bootstrap;
shader += `
@compute @workgroup_size(${GROUP_SIZE})
fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
run_id(i32(global_id.x));
}
`
let module = device.createShaderModule({
code: shader,
});
return device.createComputePipeline({
layout: "auto",
compute: {
module: module,
entryPoint: 'main',
},
});
}
device = await init_device();
pipeline = init_pipeline();