|
(function (global, factory) { |
|
typeof exports === 'object' && typeof module !== 'undefined' ? factory(exports, require('perf_hooks'), require('ws')) : |
|
typeof define === 'function' && define.amd ? define(['exports', 'perf_hooks', 'ws'], factory) : |
|
(global = typeof globalThis !== 'undefined' ? globalThis : global || self, factory(global.tvmjs = {}, global.perf_hooks, global.ws)); |
|
})(this, (function (exports, perf_hooks, ws) { 'use strict'; |
|
|
|
function _interopDefaultLegacy (e) { return e && typeof e === 'object' && 'default' in e ? e : { 'default': e }; } |
|
|
|
var perf_hooks__default = _interopDefaultLegacy(perf_hooks); |
|
var ws__default = _interopDefaultLegacy(ws); |
|
|
|
var commonjsGlobal = typeof globalThis !== 'undefined' ? globalThis : typeof window !== 'undefined' ? window : typeof global !== 'undefined' ? global : typeof self !== 'undefined' ? self : {}; |
|
|
|
function unwrapExports (x) { |
|
return x && x.__esModule && Object.prototype.hasOwnProperty.call(x, 'default') ? x['default'] : x; |
|
} |
|
|
|
function createCommonjsModule(fn, module) { |
|
return module = { exports: {} }, fn(module, module.exports), module.exports; |
|
} |
|
|
|
var support = createCommonjsModule(function (module, exports) { |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
Object.defineProperty(exports, "__esModule", { value: true }); |
|
exports.wasmPath = exports.assert = exports.Uint8ArrayToString = exports.StringToUint8Array = void 0; |
|
|
|
|
|
|
|
|
|
|
|
function StringToUint8Array(str) { |
|
const arr = new Uint8Array(str.length + 1); |
|
for (let i = 0; i < str.length; ++i) { |
|
arr[i] = str.charCodeAt(i); |
|
} |
|
arr[str.length] = 0; |
|
return arr; |
|
} |
|
exports.StringToUint8Array = StringToUint8Array; |
|
|
|
|
|
|
|
|
|
|
|
function Uint8ArrayToString(arr) { |
|
const ret = []; |
|
for (const ch of arr) { |
|
ret.push(String.fromCharCode(ch)); |
|
} |
|
return ret.join(""); |
|
} |
|
exports.Uint8ArrayToString = Uint8ArrayToString; |
|
|
|
|
|
|
|
|
|
|
|
function assert(condition, msg) { |
|
if (!condition) { |
|
throw new Error("AssertError:" + (msg || "")); |
|
} |
|
} |
|
exports.assert = assert; |
|
|
|
|
|
|
|
|
|
function wasmPath() { |
|
return __dirname + "/wasm"; |
|
} |
|
exports.wasmPath = wasmPath; |
|
|
|
}); |
|
|
|
unwrapExports(support); |
|
support.wasmPath; |
|
support.assert; |
|
support.Uint8ArrayToString; |
|
support.StringToUint8Array; |
|
|
|
var memory = createCommonjsModule(function (module, exports) { |
|
Object.defineProperty(exports, "__esModule", { value: true }); |
|
exports.CachedCallStack = exports.Memory = void 0; |
|
|
|
|
|
|
|
|
|
class Memory { |
|
constructor(memory) { |
|
this.wasm32 = true; |
|
this.memory = memory; |
|
this.buffer = this.memory.buffer; |
|
this.viewU8 = new Uint8Array(this.buffer); |
|
this.viewU16 = new Uint16Array(this.buffer); |
|
this.viewI32 = new Int32Array(this.buffer); |
|
this.viewU32 = new Uint32Array(this.buffer); |
|
this.viewF32 = new Float32Array(this.buffer); |
|
this.viewF64 = new Float64Array(this.buffer); |
|
} |
|
loadU8(ptr) { |
|
if (this.buffer != this.memory.buffer) { |
|
this.updateViews(); |
|
} |
|
return this.viewU8[ptr >> 0]; |
|
} |
|
loadU16(ptr) { |
|
if (this.buffer != this.memory.buffer) { |
|
this.updateViews(); |
|
} |
|
return this.viewU16[ptr >> 1]; |
|
} |
|
loadU32(ptr) { |
|
if (this.buffer != this.memory.buffer) { |
|
this.updateViews(); |
|
} |
|
return this.viewU32[ptr >> 2]; |
|
} |
|
loadI32(ptr) { |
|
if (this.buffer != this.memory.buffer) { |
|
this.updateViews(); |
|
} |
|
return this.viewI32[ptr >> 2]; |
|
} |
|
loadI64(ptr) { |
|
if (this.buffer != this.memory.buffer) { |
|
this.updateViews(); |
|
} |
|
const base = ptr >> 2; |
|
|
|
return this.viewI32[base]; |
|
} |
|
loadF32(ptr) { |
|
if (this.buffer != this.memory.buffer) { |
|
this.updateViews(); |
|
} |
|
return this.viewF32[ptr >> 2]; |
|
} |
|
loadF64(ptr) { |
|
if (this.buffer != this.memory.buffer) { |
|
this.updateViews(); |
|
} |
|
return this.viewF64[ptr >> 3]; |
|
} |
|
loadPointer(ptr) { |
|
if (this.buffer != this.memory.buffer) { |
|
this.updateViews(); |
|
} |
|
if (this.wasm32) { |
|
return this.loadU32(ptr); |
|
} |
|
else { |
|
return this.loadI64(ptr); |
|
} |
|
} |
|
loadUSize(ptr) { |
|
if (this.buffer != this.memory.buffer) { |
|
this.updateViews(); |
|
} |
|
if (this.wasm32) { |
|
return this.loadU32(ptr); |
|
} |
|
else { |
|
return this.loadI64(ptr); |
|
} |
|
} |
|
sizeofPtr() { |
|
return this.wasm32 ? 4 : 8 ; |
|
} |
|
|
|
|
|
|
|
|
|
|
|
loadRawBytes(ptr, numBytes) { |
|
if (this.buffer != this.memory.buffer) { |
|
this.updateViews(); |
|
} |
|
const result = new Uint8Array(numBytes); |
|
result.set(this.viewU8.slice(ptr, ptr + numBytes)); |
|
return result; |
|
} |
|
|
|
|
|
|
|
|
|
|
|
loadTVMBytes(ptr) { |
|
const data = this.loadPointer(ptr); |
|
const length = this.loadUSize(ptr + this.sizeofPtr()); |
|
return this.loadRawBytes(data, length); |
|
} |
|
|
|
|
|
|
|
|
|
loadCString(ptr) { |
|
if (this.buffer != this.memory.buffer) { |
|
this.updateViews(); |
|
} |
|
|
|
const ret = []; |
|
let ch = 1; |
|
while (ch != 0) { |
|
ch = this.viewU8[ptr]; |
|
if (ch != 0) { |
|
ret.push(String.fromCharCode(ch)); |
|
} |
|
++ptr; |
|
} |
|
return ret.join(""); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
storeRawBytes(ptr, bytes) { |
|
if (this.buffer != this.memory.buffer) { |
|
this.updateViews(); |
|
} |
|
this.viewU8.set(bytes, ptr); |
|
} |
|
|
|
|
|
|
|
updateViews() { |
|
this.buffer = this.memory.buffer; |
|
this.viewU8 = new Uint8Array(this.buffer); |
|
this.viewU16 = new Uint16Array(this.buffer); |
|
this.viewI32 = new Int32Array(this.buffer); |
|
this.viewU32 = new Uint32Array(this.buffer); |
|
this.viewF32 = new Float32Array(this.buffer); |
|
this.viewF64 = new Float64Array(this.buffer); |
|
} |
|
} |
|
exports.Memory = Memory; |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
class CachedCallStack { |
|
constructor(memory, allocSpace, freeSpace) { |
|
|
|
this.tempArgs = []; |
|
this.stackTop = 0; |
|
this.basePtr = 0; |
|
this.addressToSetTargetValue = []; |
|
const initCallStackSize = 128; |
|
this.memory = memory; |
|
this.cAllocSpace = allocSpace; |
|
this.cFreeSpace = freeSpace; |
|
this.buffer = new ArrayBuffer(initCallStackSize); |
|
this.basePtr = this.cAllocSpace(initCallStackSize); |
|
this.viewU8 = new Uint8Array(this.buffer); |
|
this.viewI32 = new Int32Array(this.buffer); |
|
this.viewU32 = new Uint32Array(this.buffer); |
|
this.viewF64 = new Float64Array(this.buffer); |
|
this.updateViews(); |
|
} |
|
dispose() { |
|
if (this.basePtr != 0) { |
|
this.cFreeSpace(this.basePtr); |
|
this.basePtr = 0; |
|
} |
|
} |
|
|
|
|
|
|
|
reset() { |
|
this.stackTop = 0; |
|
support.assert(this.addressToSetTargetValue.length == 0); |
|
while (this.tempArgs.length != 0) { |
|
this.tempArgs.pop().dispose(); |
|
} |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
commitToWasmMemory(nbytes = this.stackTop) { |
|
|
|
while (this.addressToSetTargetValue.length != 0) { |
|
const [targetOffset, valueOffset] = this.addressToSetTargetValue.pop(); |
|
this.storePtr(targetOffset, this.ptrFromOffset(valueOffset)); |
|
} |
|
this.memory.storeRawBytes(this.basePtr, this.viewU8.slice(0, nbytes)); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
allocRawBytes(nbytes) { |
|
|
|
nbytes = ((nbytes + 7) >> 3) << 3; |
|
if (this.stackTop + nbytes > this.buffer.byteLength) { |
|
const newSize = Math.max(this.buffer.byteLength * 2, this.stackTop + nbytes); |
|
const oldU8 = this.viewU8; |
|
this.buffer = new ArrayBuffer(newSize); |
|
this.updateViews(); |
|
this.viewU8.set(oldU8); |
|
if (this.basePtr != 0) { |
|
this.cFreeSpace(this.basePtr); |
|
} |
|
this.basePtr = this.cAllocSpace(newSize); |
|
} |
|
const retOffset = this.stackTop; |
|
this.stackTop += nbytes; |
|
return retOffset; |
|
} |
|
|
|
|
|
|
|
|
|
|
|
allocPtrArray(count) { |
|
return this.allocRawBytes(this.memory.sizeofPtr() * count); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
ptrFromOffset(offset) { |
|
return this.basePtr + offset; |
|
} |
|
|
|
storePtr(offset, value) { |
|
if (this.memory.wasm32) { |
|
this.storeU32(offset, value); |
|
} |
|
else { |
|
this.storeI64(offset, value); |
|
} |
|
} |
|
storeUSize(offset, value) { |
|
if (this.memory.wasm32) { |
|
this.storeU32(offset, value); |
|
} |
|
else { |
|
this.storeI64(offset, value); |
|
} |
|
} |
|
storeI32(offset, value) { |
|
this.viewI32[offset >> 2] = value; |
|
} |
|
storeU32(offset, value) { |
|
this.viewU32[offset >> 2] = value; |
|
} |
|
storeI64(offset, value) { |
|
|
|
|
|
const low = value & 0xffffffff; |
|
const base = offset >> 2; |
|
this.viewI32[base] = low; |
|
this.viewI32[base + 1] = 0; |
|
} |
|
storeF64(offset, value) { |
|
this.viewF64[offset >> 3] = value; |
|
} |
|
storeRawBytes(offset, bytes) { |
|
this.viewU8.set(bytes, offset); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
allocThenSetArgString(offset, data) { |
|
const strOffset = this.allocRawBytes(data.length + 1); |
|
this.storeRawBytes(strOffset, support.StringToUint8Array(data)); |
|
this.addressToSetTargetValue.push([offset, strOffset]); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
allocThenSetArgBytes(offset, data) { |
|
|
|
const headerOffset = this.allocRawBytes(this.memory.sizeofPtr() * 2); |
|
const dataOffset = this.allocRawBytes(data.length); |
|
this.storeRawBytes(dataOffset, data); |
|
this.storeUSize(headerOffset + this.memory.sizeofPtr(), data.length); |
|
this.addressToSetTargetValue.push([offset, headerOffset]); |
|
this.addressToSetTargetValue.push([headerOffset, dataOffset]); |
|
} |
|
|
|
|
|
|
|
updateViews() { |
|
this.viewU8 = new Uint8Array(this.buffer); |
|
this.viewI32 = new Int32Array(this.buffer); |
|
this.viewU32 = new Uint32Array(this.buffer); |
|
this.viewF64 = new Float64Array(this.buffer); |
|
} |
|
} |
|
exports.CachedCallStack = CachedCallStack; |
|
|
|
}); |
|
|
|
unwrapExports(memory); |
|
memory.CachedCallStack; |
|
memory.Memory; |
|
|
|
var environment = createCommonjsModule(function (module, exports) { |
|
Object.defineProperty(exports, "__esModule", { value: true }); |
|
exports.Environment = void 0; |
|
|
|
|
|
|
|
|
|
|
|
|
|
function detectLibraryProvider(importObject) { |
|
if (importObject["wasmLibraryProvider"] && |
|
importObject["wasmLibraryProvider"]["start"] && |
|
importObject["wasmLibraryProvider"]["imports"] !== undefined) { |
|
const item = importObject; |
|
|
|
return { |
|
imports: item.wasmLibraryProvider.imports, |
|
start: (inst) => { |
|
item.wasmLibraryProvider.start(inst); |
|
}, |
|
}; |
|
} |
|
else if (importObject["imports"] && importObject["start"] !== undefined) { |
|
return importObject; |
|
} |
|
else if (importObject["wasiImport"] && importObject["start"] !== undefined) { |
|
|
|
return { |
|
imports: { |
|
"wasi_snapshot_preview1": importObject["wasiImport"], |
|
}, |
|
start: (inst) => { |
|
importObject["start"](inst); |
|
} |
|
}; |
|
} |
|
else { |
|
return undefined; |
|
} |
|
} |
|
|
|
|
|
|
|
class Environment { |
|
constructor(importObject = {}, logger = console.log) { |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
this.packedCFuncTable = [ |
|
undefined, |
|
]; |
|
|
|
|
|
|
|
this.packedCFuncTableFreeId = []; |
|
this.logger = logger; |
|
this.libProvider = detectLibraryProvider(importObject); |
|
|
|
if (this.libProvider !== undefined) { |
|
this.imports = this.libProvider.imports; |
|
} |
|
else { |
|
this.imports = importObject; |
|
} |
|
|
|
this.imports.env = this.environment(this.imports.env); |
|
} |
|
|
|
start(inst) { |
|
if (this.libProvider !== undefined) { |
|
this.libProvider.start(inst); |
|
} |
|
} |
|
environment(initEnv) { |
|
|
|
const defaultEnv = { |
|
"__cxa_thread_atexit": () => { }, |
|
|
|
"emscripten_notify_memory_growth": (index) => { } |
|
}; |
|
const wasmPackedCFunc = (args, typeCodes, nargs, ret, resourceHandle) => { |
|
const cfunc = this.packedCFuncTable[resourceHandle]; |
|
support.assert(cfunc !== undefined); |
|
return cfunc(args, typeCodes, nargs, ret, resourceHandle); |
|
}; |
|
const wasmPackedCFuncFinalizer = (resourceHandle) => { |
|
this.packedCFuncTable[resourceHandle] = undefined; |
|
this.packedCFuncTableFreeId.push(resourceHandle); |
|
}; |
|
const newEnv = { |
|
TVMWasmPackedCFunc: wasmPackedCFunc, |
|
TVMWasmPackedCFuncFinalizer: wasmPackedCFuncFinalizer, |
|
"__console_log": (msg) => { |
|
this.logger(msg); |
|
} |
|
}; |
|
return Object.assign(defaultEnv, initEnv, newEnv); |
|
} |
|
} |
|
exports.Environment = Environment; |
|
|
|
}); |
|
|
|
unwrapExports(environment); |
|
environment.Environment; |
|
|
|
var webgpu = createCommonjsModule(function (module, exports) { |
|
var __awaiter = (commonjsGlobal && commonjsGlobal.__awaiter) || function (thisArg, _arguments, P, generator) { |
|
function adopt(value) { return value instanceof P ? value : new P(function (resolve) { resolve(value); }); } |
|
return new (P || (P = Promise))(function (resolve, reject) { |
|
function fulfilled(value) { try { step(generator.next(value)); } catch (e) { reject(e); } } |
|
function rejected(value) { try { step(generator["throw"](value)); } catch (e) { reject(e); } } |
|
function step(result) { result.done ? resolve(result.value) : adopt(result.value).then(fulfilled, rejected); } |
|
step((generator = generator.apply(thisArg, _arguments || [])).next()); |
|
}); |
|
}; |
|
Object.defineProperty(exports, "__esModule", { value: true }); |
|
exports.WebGPUContext = exports.detectGPUDevice = void 0; |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
function detectGPUDevice() { |
|
return __awaiter(this, void 0, void 0, function* () { |
|
if (typeof navigator !== "undefined" && navigator.gpu !== undefined) { |
|
const adapter = yield navigator.gpu.requestAdapter({ "powerPreference": "high-performance" }); |
|
if (adapter == null) { |
|
throw Error("Cannot find adapter that matches the request"); |
|
} |
|
const computeMB = (value) => { |
|
return Math.ceil(value / (1 << 20)) + "MB"; |
|
}; |
|
|
|
const requiedMaxBufferSize = 1 << 30; |
|
if (requiedMaxBufferSize > adapter.limits.maxBufferSize) { |
|
throw Error(`Cannot initialize runtime because of requested maxBufferSize ` + |
|
`exceeds limit. requested=${computeMB(requiedMaxBufferSize)}, ` + |
|
`limit=${computeMB(adapter.limits.maxBufferSize)}. ` + |
|
`This error may be caused by an older version of the browser (e.g. Chrome 112). ` + |
|
`You can try to upgrade your browser to Chrome 113 or later.`); |
|
} |
|
const requiredMaxStorageBufferBindingSize = 1 << 30; |
|
if (requiredMaxStorageBufferBindingSize > adapter.limits.maxStorageBufferBindingSize) { |
|
throw Error(`Cannot initialize runtime because of requested maxStorageBufferBindingSize ` + |
|
`exceeds limit. requested=${computeMB(requiredMaxStorageBufferBindingSize)}, ` + |
|
`limit=${computeMB(adapter.limits.maxStorageBufferBindingSize)}. `); |
|
} |
|
const requiredMaxComputeWorkgroupStorageSize = 32 << 10; |
|
if (requiredMaxComputeWorkgroupStorageSize > adapter.limits.maxComputeWorkgroupStorageSize) { |
|
throw Error(`Cannot initialize runtime because of requested maxComputeWorkgroupStorageSize ` + |
|
`exceeds limit. requested=${requiredMaxComputeWorkgroupStorageSize}, ` + |
|
`limit=${adapter.limits.maxComputeWorkgroupStorageSize}. `); |
|
} |
|
const adapterInfo = yield adapter.requestAdapterInfo(); |
|
const device = yield adapter.requestDevice({ |
|
requiredLimits: { |
|
maxBufferSize: requiedMaxBufferSize, |
|
maxStorageBufferBindingSize: requiredMaxStorageBufferBindingSize, |
|
maxComputeWorkgroupStorageSize: requiredMaxComputeWorkgroupStorageSize, |
|
} |
|
}); |
|
return { |
|
adapter: adapter, |
|
adapterInfo: adapterInfo, |
|
device: device |
|
}; |
|
} |
|
else { |
|
return undefined; |
|
} |
|
}); |
|
} |
|
exports.detectGPUDevice = detectGPUDevice; |
|
const canvasRenderWGSL = ` |
|
@group(0) @binding(0) var my_sampler : sampler; |
|
@group(0) @binding(1) var my_texture : texture_2d<f32>; |
|
|
|
struct VertexOutput { |
|
@builtin(position) position : vec4<f32>, |
|
@location(0) uv : vec2<f32>, |
|
} |
|
|
|
@vertex |
|
fn vertex_main(@builtin(vertex_index) vidx : u32) -> VertexOutput { |
|
const pos = array( |
|
vec2( 1.0, 1.0), |
|
vec2( 1.0, -1.0), |
|
vec2(-1.0, -1.0), |
|
vec2( 1.0, 1.0), |
|
vec2(-1.0, -1.0), |
|
vec2(-1.0, 1.0), |
|
); |
|
|
|
const uv = array( |
|
vec2(1.0, 0.0), |
|
vec2(1.0, 1.0), |
|
vec2(0.0, 1.0), |
|
vec2(1.0, 0.0), |
|
vec2(0.0, 1.0), |
|
vec2(0.0, 0.0), |
|
); |
|
|
|
var output : VertexOutput; |
|
output.position = vec4(pos[vidx], 0.0, 1.0); |
|
output.uv = uv[vidx]; |
|
return output; |
|
} |
|
|
|
@fragment |
|
fn fragment_main(@location(0) uv : vec2<f32>) -> @location(0) vec4<f32> { |
|
return textureSample(my_texture, my_sampler, uv); |
|
} |
|
|
|
@fragment |
|
fn fragment_clear(@location(0) uv : vec2<f32>) -> @location(0) vec4<f32> { |
|
return vec4(1.0, 1.0, 1.0, 1.0); |
|
} |
|
`; |
|
class CanvaRenderManager { |
|
constructor(device, canvas) { |
|
this.device = device; |
|
const ctx = canvas.getContext("webgpu", { |
|
alpha: false, |
|
antialias: false, |
|
powerPreference: "high-performance", |
|
}); |
|
if (ctx == null) { |
|
throw Error("Cannot bind WebGPU context"); |
|
} |
|
this.canvasContext = ctx; |
|
this.canvasTextureFormat = navigator.gpu.getPreferredCanvasFormat(); |
|
this.canvasContext.configure({ |
|
device: this.device, |
|
format: this.canvasTextureFormat, |
|
alphaMode: "opaque", |
|
}); |
|
this.renderPipeline = device.createRenderPipeline({ |
|
layout: "auto", |
|
vertex: { |
|
module: device.createShaderModule({ |
|
code: canvasRenderWGSL, |
|
}), |
|
entryPoint: "vertex_main", |
|
}, |
|
fragment: { |
|
module: device.createShaderModule({ |
|
code: canvasRenderWGSL, |
|
}), |
|
entryPoint: "fragment_main", |
|
targets: [{ |
|
format: this.canvasTextureFormat, |
|
}], |
|
}, |
|
primitive: { |
|
topology: "triangle-list", |
|
}, |
|
}); |
|
this.clearPipeline = device.createRenderPipeline({ |
|
layout: "auto", |
|
vertex: { |
|
module: device.createShaderModule({ |
|
code: canvasRenderWGSL, |
|
}), |
|
entryPoint: "vertex_main", |
|
}, |
|
fragment: { |
|
module: device.createShaderModule({ |
|
code: canvasRenderWGSL, |
|
}), |
|
entryPoint: "fragment_clear", |
|
targets: [{ |
|
format: this.canvasTextureFormat, |
|
}], |
|
}, |
|
primitive: { |
|
topology: "triangle-list", |
|
}, |
|
}); |
|
this.renderSampler = device.createSampler({ |
|
magFilter: "linear", |
|
minFilter: "linear", |
|
}); |
|
|
|
this.stagingTexture = device.createTexture({ |
|
size: [canvas.height, canvas.width, 1], |
|
format: "rgba8unorm", |
|
usage: GPUTextureUsage.TEXTURE_BINDING | |
|
GPUTextureUsage.COPY_DST | |
|
GPUTextureUsage.RENDER_ATTACHMENT, |
|
}); |
|
} |
|
clear() { |
|
const commandEncoder = this.device.createCommandEncoder(); |
|
const passEncoder = commandEncoder.beginRenderPass({ |
|
colorAttachments: [ |
|
{ |
|
view: this.canvasContext.getCurrentTexture().createView(), |
|
clearValue: { r: 0.0, g: 0.0, b: 0.0, a: 1.0 }, |
|
loadOp: "clear", |
|
storeOp: "store", |
|
}, |
|
], |
|
}); |
|
passEncoder.setPipeline(this.clearPipeline); |
|
const renderBindingGroup = this.device.createBindGroup({ |
|
layout: this.renderPipeline.getBindGroupLayout(0), |
|
entries: [ |
|
{ binding: 0, resource: this.renderSampler }, |
|
{ binding: 1, resource: this.stagingTexture.createView() }, |
|
], |
|
}); |
|
passEncoder.setBindGroup(0, renderBindingGroup); |
|
passEncoder.draw(6, 1, 0, 0); |
|
passEncoder.end(); |
|
this.device.queue.submit([commandEncoder.finish()]); |
|
} |
|
draw(buffer, height, width) { |
|
|
|
if (height != this.stagingTexture.height || width != this.stagingTexture.width) { |
|
this.stagingTexture.destroy(); |
|
this.stagingTexture = this.device.createTexture({ |
|
size: [height, width, 1], |
|
format: "rgba8unorm", |
|
usage: GPUTextureUsage.TEXTURE_BINDING | |
|
GPUTextureUsage.COPY_DST | |
|
GPUTextureUsage.RENDER_ATTACHMENT, |
|
}); |
|
} |
|
const commandEncoder = this.device.createCommandEncoder(); |
|
commandEncoder.copyBufferToTexture({ |
|
buffer: buffer, |
|
offset: 0, |
|
bytesPerRow: this.stagingTexture.width * 4 |
|
}, { |
|
texture: this.stagingTexture |
|
}, { |
|
width: this.stagingTexture.width, |
|
height: this.stagingTexture.height |
|
}); |
|
const passEncoder = commandEncoder.beginRenderPass({ |
|
colorAttachments: [ |
|
{ |
|
view: this.canvasContext.getCurrentTexture().createView(), |
|
clearValue: { r: 0.0, g: 0.0, b: 0.0, a: 1.0 }, |
|
loadOp: "clear", |
|
storeOp: "store", |
|
}, |
|
], |
|
}); |
|
passEncoder.setPipeline(this.renderPipeline); |
|
const renderBindingGroup = this.device.createBindGroup({ |
|
layout: this.renderPipeline.getBindGroupLayout(0), |
|
entries: [ |
|
{ binding: 0, resource: this.renderSampler }, |
|
{ binding: 1, resource: this.stagingTexture.createView() }, |
|
], |
|
}); |
|
passEncoder.setBindGroup(0, renderBindingGroup); |
|
passEncoder.draw(6, 1, 0, 0); |
|
passEncoder.end(); |
|
this.device.queue.submit([commandEncoder.finish()]); |
|
} |
|
dispose() { |
|
this.stagingTexture.destroy(); |
|
} |
|
} |
|
|
|
|
|
|
|
|
|
class WebGPUContext { |
|
constructor(memory, device) { |
|
|
|
this.bufferTable = [undefined]; |
|
this.bufferTableFreeId = []; |
|
this.podArgStagingBuffers = []; |
|
this.canvasRenderManager = undefined; |
|
|
|
this.maxNumPodArgsStagingBuffers = 2; |
|
|
|
|
|
|
|
this.peakAllocatedBytes = 0; |
|
|
|
this.currAllocatedBytes = 0; |
|
|
|
this.allAllocatedBytes = 0; |
|
|
|
this.shaderSubmitCounter = 0; |
|
|
|
this.debugShaderSubmitLimit = -1; |
|
|
|
this.debugLogFinish = false; |
|
this.memory = memory; |
|
this.device = device; |
|
} |
|
|
|
|
|
|
|
dispose() { |
|
var _a, _b, _c; |
|
(_a = this.canvasRenderManager) === null || _a === void 0 ? void 0 : _a.dispose(); |
|
this.bufferTableFreeId = []; |
|
while (this.bufferTable.length != 0) { |
|
(_b = this.bufferTable.pop()) === null || _b === void 0 ? void 0 : _b.destroy(); |
|
} |
|
while (this.podArgStagingBuffers.length != 0) { |
|
(_c = this.podArgStagingBuffers.pop()) === null || _c === void 0 ? void 0 : _c.destroy(); |
|
} |
|
this.device.destroy(); |
|
} |
|
|
|
|
|
|
|
sync() { |
|
return __awaiter(this, void 0, void 0, function* () { |
|
yield this.device.queue.onSubmittedWorkDone(); |
|
}); |
|
} |
|
|
|
|
|
|
|
runtimeStatsText() { |
|
let info = "peak-memory=" + Math.ceil(this.peakAllocatedBytes / (1 << 20)) + " MB"; |
|
info += ", all-memory=" + Math.ceil(this.allAllocatedBytes / (1 << 20)) + " MB"; |
|
info += ", shader-submissions=" + this.shaderSubmitCounter; |
|
return info; |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
drawImageFromBuffer(ptr, height, width) { |
|
if (this.canvasRenderManager == undefined) { |
|
throw Error("Do not have a canvas context, call bindCanvas first"); |
|
} |
|
this.canvasRenderManager.draw(this.gpuBufferFromPtr(ptr), height, width); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
copyRawBytesToBuffer(rawBytes, toPtr, toOffset, nbytes) { |
|
|
|
this.device.queue.writeBuffer(this.gpuBufferFromPtr(toPtr), toOffset, rawBytes, 0, nbytes); |
|
} |
|
|
|
|
|
|
|
clearCanvas() { |
|
var _a; |
|
(_a = this.canvasRenderManager) === null || _a === void 0 ? void 0 : _a.clear(); |
|
} |
|
|
|
|
|
|
|
|
|
bindCanvas(canvas) { |
|
this.canvasRenderManager = new CanvaRenderManager(this.device, canvas); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
createShader(finfo, code) { |
|
return this.createShadeInternal(finfo, code, false); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
createShaderAsync(finfo, code) { |
|
return __awaiter(this, void 0, void 0, function* () { |
|
return yield this.createShadeInternal(finfo, code, true); |
|
}); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
getPodArgsBuffer(nbytes) { |
|
let buffer = undefined; |
|
if (this.podArgStagingBuffers.length >= this.maxNumPodArgsStagingBuffers) { |
|
buffer = this.podArgStagingBuffers.shift(); |
|
} |
|
|
|
let allocSize = 16; |
|
if (buffer !== undefined) { |
|
allocSize = buffer.size; |
|
if (buffer.size < nbytes) { |
|
buffer.destroy(); |
|
buffer = undefined; |
|
} |
|
} |
|
while (allocSize < nbytes) { |
|
allocSize *= 2; |
|
} |
|
if (buffer == undefined) { |
|
|
|
buffer = this.device.createBuffer({ |
|
size: allocSize, |
|
usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, |
|
}); |
|
} |
|
support.assert(nbytes <= buffer.size); |
|
return buffer; |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
createShadeInternal(finfo, code, asyncMode) { |
|
const dispatchToDim = []; |
|
let paramWriteAccess = []; |
|
for (let i = 0; i < finfo.launch_param_tags.length; ++i) { |
|
const tag = finfo.launch_param_tags[i]; |
|
if (tag.startsWith("blockIdx.")) { |
|
const target = tag.charCodeAt(tag.length - 1) - ("x".charCodeAt(0)); |
|
support.assert(target >= 0 && target < 3); |
|
dispatchToDim.push(target); |
|
} |
|
else if (tag.startsWith("threadIdx.")) { |
|
const target = tag.charCodeAt(tag.length - 1) - ("x".charCodeAt(0)); |
|
support.assert(target >= 0 && target < 3); |
|
dispatchToDim.push(target + 3); |
|
} |
|
else if (tag.startsWith("paramWriteAccess:")) { |
|
paramWriteAccess = JSON.parse(tag.substring(17)); |
|
} |
|
else { |
|
throw new Error("Cannot handle thread_axis " + tag); |
|
} |
|
} |
|
const layoutEntries = []; |
|
const bufferArgIndices = []; |
|
const podArgIndices = []; |
|
for (let i = 0; i < finfo.arg_types.length; ++i) { |
|
const dtype = finfo.arg_types[i]; |
|
if (dtype == "handle") { |
|
layoutEntries.push({ |
|
binding: bufferArgIndices.length, |
|
visibility: GPUShaderStage.COMPUTE, |
|
buffer: { |
|
type: paramWriteAccess[bufferArgIndices.length] ? "storage" : "read-only-storage" |
|
} |
|
}); |
|
bufferArgIndices.push(i); |
|
} |
|
else if (dtype.startsWith("int") || dtype.startsWith("uint") || dtype.startsWith("float")) { |
|
podArgIndices.push(i); |
|
} |
|
else { |
|
throw new Error("Cannot handle argument type " + dtype + " in WebGPU shader"); |
|
} |
|
} |
|
support.assert(paramWriteAccess.length == bufferArgIndices.length); |
|
|
|
layoutEntries.push({ |
|
binding: bufferArgIndices.length, |
|
visibility: GPUShaderStage.COMPUTE, |
|
buffer: { |
|
type: "uniform" |
|
} |
|
}); |
|
const bindGroupLayout = this.device.createBindGroupLayout({ |
|
entries: layoutEntries |
|
}); |
|
const pipelineLayout = this.device.createPipelineLayout({ |
|
bindGroupLayouts: [bindGroupLayout] |
|
}); |
|
|
|
const createShaderFunc = (pipeline) => { |
|
const submitShader = (...args) => { |
|
if (this.debugShaderSubmitLimit != -1 && |
|
this.shaderSubmitCounter >= this.debugShaderSubmitLimit) { |
|
this.shaderSubmitCounter += 1; |
|
return; |
|
} |
|
const commandEncoder = this.device.createCommandEncoder(); |
|
const compute = commandEncoder.beginComputePass(); |
|
compute.setPipeline(pipeline); |
|
const bindGroupEntries = []; |
|
const numBufferOrPodArgs = bufferArgIndices.length + podArgIndices.length; |
|
support.assert(args.length == numBufferOrPodArgs + dispatchToDim.length); |
|
const workDim = [1, 1, 1, 1, 1, 1]; |
|
for (let i = 0; i < dispatchToDim.length; ++i) { |
|
workDim[dispatchToDim[i]] = args[numBufferOrPodArgs + i]; |
|
} |
|
|
|
if (workDim[2] != 1) { |
|
throw Error("WebGPU: blockIdx.z is reserved for internal use"); |
|
} |
|
const packDimX = workDim[0]; |
|
|
|
if (workDim[0] >= (1 << 16)) { |
|
let wl_x = workDim[0]; |
|
let wl_z = workDim[2]; |
|
while (wl_x >= (1 << 16)) { |
|
if (wl_x % 2 == 0) { |
|
wl_x = wl_x / 2; |
|
} |
|
else { |
|
|
|
wl_x = (wl_x + 1) / 2; |
|
} |
|
wl_z *= 2; |
|
} |
|
workDim[0] = wl_x; |
|
workDim[2] = wl_z; |
|
support.assert(wl_x * wl_z >= packDimX); |
|
} |
|
for (let i = 0; i < bufferArgIndices.length; ++i) { |
|
bindGroupEntries.push({ |
|
binding: i, |
|
resource: { |
|
buffer: this.gpuBufferFromPtr(args[bufferArgIndices[i]]) |
|
} |
|
}); |
|
} |
|
|
|
const sizeOfI32 = 4; |
|
const podArgBuffer = this.getPodArgsBuffer((podArgIndices.length + 1) * sizeOfI32); |
|
const i32View = new Int32Array(podArgIndices.length + 1); |
|
const u32View = new Uint32Array(i32View.buffer); |
|
const f32View = new Float32Array(i32View.buffer); |
|
for (let i = 0; i < podArgIndices.length; ++i) { |
|
const value = args[podArgIndices[i]]; |
|
const dtype = finfo.arg_types[podArgIndices[i]]; |
|
if (dtype.startsWith("int")) { |
|
i32View[i] = value; |
|
} |
|
else if (dtype.startsWith("uint")) { |
|
u32View[i] = value; |
|
} |
|
else if (dtype.startsWith("float")) { |
|
f32View[i] = value; |
|
} |
|
else { |
|
throw Error("Unknown pod dtype " + dtype); |
|
} |
|
} |
|
|
|
u32View[podArgIndices.length] = packDimX; |
|
this.device.queue.writeBuffer(podArgBuffer, 0, i32View.buffer); |
|
bindGroupEntries.push({ |
|
binding: bufferArgIndices.length, |
|
resource: { |
|
buffer: podArgBuffer, |
|
size: i32View.buffer.byteLength |
|
} |
|
}); |
|
compute.setBindGroup(0, this.device.createBindGroup({ |
|
layout: bindGroupLayout, |
|
entries: bindGroupEntries |
|
})); |
|
compute.dispatchWorkgroups(workDim[0], workDim[1], workDim[2]); |
|
compute.end(); |
|
const command = commandEncoder.finish(); |
|
this.device.queue.submit([command]); |
|
if (this.debugLogFinish) { |
|
const currCounter = this.shaderSubmitCounter; |
|
this.device.queue.onSubmittedWorkDone().then(() => { |
|
console.log("[" + currCounter + "][Debug] finish shader" + finfo.name); |
|
}); |
|
} |
|
this.shaderSubmitCounter += 1; |
|
}; |
|
return submitShader; |
|
}; |
|
const shaderModule = this.device.createShaderModule({ |
|
code: code, |
|
hints: { |
|
main: { |
|
layout: pipelineLayout |
|
} |
|
} |
|
}); |
|
if (asyncMode) { |
|
return this.device.createComputePipelineAsync({ |
|
layout: pipelineLayout, |
|
compute: { |
|
module: shaderModule, |
|
entryPoint: finfo.name |
|
} |
|
}).then((pipeline) => { |
|
return createShaderFunc(pipeline); |
|
}); |
|
} |
|
else { |
|
const pipeline = this.device.createComputePipeline({ |
|
layout: pipelineLayout, |
|
compute: { |
|
module: shaderModule, |
|
entryPoint: finfo.name |
|
} |
|
}); |
|
return createShaderFunc(pipeline); |
|
} |
|
} |
|
|
|
|
|
|
|
|
|
|
|
getDeviceAPI(name) { |
|
if (name == "deviceAllocDataSpace") { |
|
return (nbytes) => { |
|
return this.deviceAllocDataSpace(nbytes); |
|
}; |
|
} |
|
else if (name == "deviceFreeDataSpace") { |
|
return (ptr) => { |
|
return this.deviceFreeDataSpace(ptr); |
|
}; |
|
} |
|
else if (name == "deviceCopyToGPU") { |
|
return (from, to, toOffset, nbytes) => { |
|
this.deviceCopyToGPU(from, to, toOffset, nbytes); |
|
}; |
|
} |
|
else if (name == "deviceCopyFromGPU") { |
|
return (from, fromOffset, to, nbytes) => { |
|
this.deviceCopyFromGPU(from, fromOffset, to, nbytes); |
|
}; |
|
} |
|
else if (name == "deviceCopyWithinGPU") { |
|
return (from, fromOffset, to, toOffset, nbytes) => { |
|
this.deviceCopyWithinGPU(from, fromOffset, to, toOffset, nbytes); |
|
}; |
|
} |
|
else { |
|
throw new Error("Unknown DeviceAPI function " + name); |
|
} |
|
} |
|
|
|
deviceAllocDataSpace(nbytes) { |
|
|
|
if (nbytes == 0) { |
|
nbytes = 1; |
|
} |
|
const buffer = this.device.createBuffer({ |
|
size: nbytes, |
|
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, |
|
}); |
|
this.currAllocatedBytes += nbytes; |
|
this.allAllocatedBytes += nbytes; |
|
if (this.currAllocatedBytes > this.peakAllocatedBytes) { |
|
this.peakAllocatedBytes = this.currAllocatedBytes; |
|
} |
|
const ptr = this.attachToBufferTable(buffer); |
|
return ptr; |
|
} |
|
deviceFreeDataSpace(ptr) { |
|
const idx = ptr; |
|
const buffer = this.bufferTable[idx]; |
|
this.bufferTable[idx] = undefined; |
|
support.assert(buffer !== undefined); |
|
this.bufferTableFreeId.push(idx); |
|
this.currAllocatedBytes -= buffer.size; |
|
buffer.destroy(); |
|
} |
|
deviceCopyToGPU(from, to, toOffset, nbytes) { |
|
|
|
const rawBytes = this.memory.loadRawBytes(from, nbytes); |
|
this.device.queue.writeBuffer(this.gpuBufferFromPtr(to), toOffset, rawBytes, 0, nbytes); |
|
} |
|
deviceCopyFromGPU(from, fromOffset, to, nbytes) { |
|
|
|
const gpuTemp = this.device.createBuffer({ |
|
size: nbytes, |
|
usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST, |
|
}); |
|
const copyEncoder = this.device.createCommandEncoder(); |
|
copyEncoder.copyBufferToBuffer(this.gpuBufferFromPtr(from), fromOffset, gpuTemp, 0, nbytes); |
|
const copyCommands = copyEncoder.finish(); |
|
this.device.queue.submit([copyCommands]); |
|
gpuTemp.mapAsync(GPUMapMode.READ).then(() => { |
|
const data = gpuTemp.getMappedRange(); |
|
this.memory.storeRawBytes(to, new Uint8Array(data)); |
|
gpuTemp.destroy(); |
|
}); |
|
} |
|
deviceCopyWithinGPU(from, fromOffset, to, toOffset, nbytes) { |
|
const copyEncoder = this.device.createCommandEncoder(); |
|
copyEncoder.copyBufferToBuffer(this.gpuBufferFromPtr(from), fromOffset, this.gpuBufferFromPtr(to), toOffset, nbytes); |
|
const copyCommands = copyEncoder.finish(); |
|
this.device.queue.submit([copyCommands]); |
|
} |
|
gpuBufferFromPtr(ptr) { |
|
const buffer = this.bufferTable[ptr]; |
|
support.assert(buffer !== undefined); |
|
return buffer; |
|
} |
|
attachToBufferTable(buffer) { |
|
if (this.bufferTableFreeId.length != 0) { |
|
const idx = this.bufferTableFreeId.pop(); |
|
this.bufferTable[idx] = buffer; |
|
return idx; |
|
} |
|
else { |
|
const idx = this.bufferTable.length; |
|
this.bufferTable.push(buffer); |
|
return idx; |
|
} |
|
} |
|
} |
|
exports.WebGPUContext = WebGPUContext; |
|
|
|
}); |
|
|
|
unwrapExports(webgpu); |
|
webgpu.WebGPUContext; |
|
webgpu.detectGPUDevice; |
|
|
|
var compact = createCommonjsModule(function (module, exports) { |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
Object.defineProperty(exports, "__esModule", { value: true }); |
|
exports.createWebSocket = exports.getPerformance = void 0; |
|
|
|
|
|
|
|
function getPerformance() { |
|
if (typeof performance == "undefined") { |
|
|
|
const performanceNode = perf_hooks__default["default"]; |
|
return performanceNode.performance; |
|
} |
|
else { |
|
return performance; |
|
} |
|
} |
|
exports.getPerformance = getPerformance; |
|
|
|
|
|
|
|
|
|
function createWebSocket(url) { |
|
if (typeof WebSocket == "undefined") { |
|
|
|
const WebSocket = ws__default["default"]; |
|
return new WebSocket(url); |
|
} |
|
else { |
|
return new WebSocket(url); |
|
} |
|
} |
|
exports.createWebSocket = createWebSocket; |
|
|
|
}); |
|
|
|
unwrapExports(compact); |
|
compact.createWebSocket; |
|
compact.getPerformance; |
|
|
|
var runtime = createCommonjsModule(function (module, exports) { |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
var __awaiter = (commonjsGlobal && commonjsGlobal.__awaiter) || function (thisArg, _arguments, P, generator) { |
|
function adopt(value) { return value instanceof P ? value : new P(function (resolve) { resolve(value); }); } |
|
return new (P || (P = Promise))(function (resolve, reject) { |
|
function fulfilled(value) { try { step(generator.next(value)); } catch (e) { reject(e); } } |
|
function rejected(value) { try { step(generator["throw"](value)); } catch (e) { reject(e); } } |
|
function step(result) { result.done ? resolve(result.value) : adopt(result.value).then(fulfilled, rejected); } |
|
step((generator = generator.apply(thisArg, _arguments || [])).next()); |
|
}); |
|
}; |
|
Object.defineProperty(exports, "__esModule", { value: true }); |
|
exports.instantiate = exports.Instance = exports.VirtualMachine = exports.TVMArray = exports.TVMObject = exports.Module = exports.NDArray = exports.DLDataType = exports.DLDevice = exports.Scalar = void 0; |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
class FFILibrary { |
|
constructor(wasmInstance, imports) { |
|
this.recycledCallStacks = []; |
|
this.wasmInstance = wasmInstance; |
|
this.memory = new memory.Memory(this.detectWasmMemory(this.wasmInstance, imports)); |
|
support.assert(this.wasmInstance.exports !== undefined, "Expect the library module contains exports"); |
|
this.exports = this.wasmInstance.exports; |
|
this.wasm32 = this.memory.wasm32; |
|
this.validateInstance(); |
|
} |
|
dispose() { |
|
var _a; |
|
while (this.recycledCallStacks.length != 0) { |
|
this.recycledCallStacks.pop().dispose(); |
|
} |
|
(_a = this.webGPUContext) === null || _a === void 0 ? void 0 : _a.dispose(); |
|
} |
|
sizeofPtr() { |
|
return this.memory.sizeofPtr(); |
|
} |
|
checkCall(code) { |
|
if (code != 0) { |
|
const msgPtr = this.exports |
|
.TVMGetLastError(); |
|
throw new Error("TVMError: " + this.memory.loadCString(msgPtr)); |
|
} |
|
} |
|
getOrAllocCallStack() { |
|
if (this.recycledCallStacks.length != 0) { |
|
return this.recycledCallStacks.pop(); |
|
} |
|
return new memory.CachedCallStack(this.memory, this.exports.TVMWasmAllocSpace, this.exports.TVMWasmFreeSpace); |
|
} |
|
recycleCallStack(callstack) { |
|
callstack.reset(); |
|
this.recycledCallStacks.push(callstack); |
|
} |
|
validateInstance() { |
|
this.checkExports(["TVMWasmAllocSpace", "TVMWasmFreeSpace", "TVMFuncFree"]); |
|
} |
|
checkExports(funcNames) { |
|
const missList = []; |
|
for (const name of funcNames) { |
|
const f = this.exports[name]; |
|
if (!(f instanceof Function)) { |
|
missList.push(name); |
|
} |
|
} |
|
if (missList.length != 0) { |
|
throw new Error("Cannot find " + missList + " in exports"); |
|
} |
|
} |
|
detectWasmMemory(instance, imports) { |
|
if (instance.exports.memory instanceof WebAssembly.Memory) { |
|
return instance.exports.memory; |
|
} |
|
if (imports.env && imports.env.memory instanceof WebAssembly.Memory) { |
|
return imports.env.memory; |
|
} |
|
throw new Error("Cannt detect wasm memory from imports " + |
|
imports + |
|
" or exports" + |
|
instance.exports); |
|
} |
|
} |
|
|
|
|
|
|
|
|
|
class RuntimeContext { |
|
constructor(getGlobalFunc) { |
|
this.autoDisposeScope = []; |
|
this.arrayGetItem = getGlobalFunc("runtime.ArrayGetItem"); |
|
this.arrayGetSize = getGlobalFunc("runtime.ArraySize"); |
|
this.arrayMake = getGlobalFunc("runtime.Array"); |
|
this.getSysLib = getGlobalFunc("runtime.SystemLib"); |
|
this.arrayCacheGet = getGlobalFunc("vm.builtin.ndarray_cache.get"); |
|
this.arrayCacheRemove = getGlobalFunc("vm.builtin.ndarray_cache.remove"); |
|
this.arrayCacheUpdate = getGlobalFunc("vm.builtin.ndarray_cache.update"); |
|
this.arrayCacheClear = getGlobalFunc("vm.builtin.ndarray_cache.clear"); |
|
this.arrayDecodeStorage = getGlobalFunc("tvmjs.array.decode_storage"); |
|
this.paramModuleFromCache = getGlobalFunc("vm.builtin.param_module_from_cache"); |
|
this.makeShapeTuple = getGlobalFunc("runtime.ShapeTuple"); |
|
this.ndarrayCreateView = getGlobalFunc("runtime.TVMArrayCreateView"); |
|
this.sampleTopPFromLogits = getGlobalFunc("vm.builtin.sample_top_p_from_logits"); |
|
} |
|
dispose() { |
|
|
|
this.arrayCacheClear.dispose(); |
|
this.arrayGetItem.dispose(); |
|
this.arrayGetSize.dispose(); |
|
this.arrayMake.dispose(); |
|
this.arrayCacheGet.dispose(); |
|
this.arrayCacheRemove.dispose(); |
|
this.arrayCacheUpdate.dispose(); |
|
this.arrayCacheClear.dispose(); |
|
this.arrayDecodeStorage.dispose(); |
|
this.paramModuleFromCache.dispose(); |
|
this.makeShapeTuple.dispose(); |
|
this.ndarrayCreateView.dispose(); |
|
this.sampleTopPFromLogits.dispose(); |
|
} |
|
beginScope() { |
|
this.autoDisposeScope.push([]); |
|
} |
|
endScope() { |
|
if (this.autoDisposeScope.length == 0) { |
|
throw Error("tvm.endScope called when the stack is empty."); |
|
} |
|
|
|
const currScope = this.autoDisposeScope.pop(); |
|
for (let i = 0; i < currScope.length; ++i) { |
|
const val = currScope[i]; |
|
if (val !== undefined) { |
|
val.dispose(); |
|
} |
|
} |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
attachToCurrentScope(obj) { |
|
if (this.autoDisposeScope.length == 0) { |
|
throw Error("Must call beginScope to use functions that returns TVM objects"); |
|
} |
|
const currScope = this.autoDisposeScope[this.autoDisposeScope.length - 1]; |
|
currScope.push(obj); |
|
return obj; |
|
} |
|
moveToParentScope(obj) { |
|
this.detachFromCurrentScope(obj); |
|
if (this.autoDisposeScope.length < 2) { |
|
throw Error("moveToParentScope: Parent scope do not exist"); |
|
} |
|
const parentScope = this.autoDisposeScope[this.autoDisposeScope.length - 2]; |
|
parentScope.push(obj); |
|
return obj; |
|
} |
|
detachFromCurrentScope(obj) { |
|
const currScope = this.autoDisposeScope[this.autoDisposeScope.length - 1]; |
|
let occurance = 0; |
|
for (let i = 0; i < currScope.length; ++i) { |
|
if (currScope[i] === obj) { |
|
occurance += 1; |
|
currScope[i] = undefined; |
|
} |
|
} |
|
if (occurance == 0) { |
|
throw Error("Cannot find obj in the current auto conversion pool"); |
|
} |
|
if (occurance > 1) { |
|
throw Error("Value attached to scope multiple times"); |
|
} |
|
return obj; |
|
} |
|
} |
|
|
|
|
|
|
|
|
|
class Scalar { |
|
constructor(value, dtype) { |
|
this.value = value; |
|
this.dtype = dtype; |
|
} |
|
} |
|
exports.Scalar = Scalar; |
|
|
|
|
|
|
|
class PackedFuncCell { |
|
constructor(handle, lib) { |
|
this.handle = handle; |
|
this.lib = lib; |
|
} |
|
dispose() { |
|
if (this.handle != 0) { |
|
this.lib.checkCall(this.lib.exports.TVMFuncFree(this.handle)); |
|
this.handle = 0; |
|
} |
|
} |
|
getHandle(requireNotNull = true) { |
|
if (requireNotNull && this.handle == 0) { |
|
throw Error("PackedFunc has already been disposed"); |
|
} |
|
return this.handle; |
|
} |
|
} |
|
const DeviceEnumToStr = { |
|
1: "cpu", |
|
2: "cuda", |
|
4: "opencl", |
|
8: "metal", |
|
15: "webgpu" |
|
}; |
|
const DeviceStrToEnum = { |
|
cpu: 1, |
|
cuda: 2, |
|
cl: 4, |
|
opencl: 4, |
|
vulkan: 7, |
|
metal: 8, |
|
webgpu: 15 |
|
}; |
|
|
|
|
|
|
|
class DLDevice { |
|
constructor(deviceType, deviceId, lib) { |
|
const tp = typeof deviceType; |
|
if (tp == "string") { |
|
this.deviceType = DeviceStrToEnum[deviceType]; |
|
if (this.deviceType == undefined) { |
|
throw new Error("Cannot recogonize deviceType " + deviceType); |
|
} |
|
} |
|
else if (tp == "number") { |
|
this.deviceType = deviceType; |
|
} |
|
else { |
|
throw new Error("Cannot take type " + tp + " as deviceType"); |
|
} |
|
this.deviceId = deviceId; |
|
this.lib = lib; |
|
} |
|
|
|
|
|
|
|
sync() { |
|
return __awaiter(this, void 0, void 0, function* () { |
|
if (this.deviceType == DeviceStrToEnum.webgpu) { |
|
support.assert(this.lib.webGPUContext !== undefined); |
|
yield this.lib.webGPUContext.sync(); |
|
} |
|
}); |
|
} |
|
toString() { |
|
return (DeviceEnumToStr[this.deviceType] + "(" + this.deviceId.toString() + ")"); |
|
} |
|
} |
|
exports.DLDevice = DLDevice; |
|
const DLDataTypeCodeToStr = { |
|
0: "int", |
|
1: "uint", |
|
2: "float", |
|
3: "handle", |
|
}; |
|
|
|
|
|
|
|
class DLDataType { |
|
constructor(code, bits, lanes) { |
|
this.code = code; |
|
this.bits = bits; |
|
this.lanes = lanes; |
|
} |
|
toString() { |
|
const ret = DLDataTypeCodeToStr[this.code] + this.bits.toString(); |
|
if (this.lanes != 1) { |
|
return ret + "x" + this.lanes.toString(); |
|
} |
|
else { |
|
return ret; |
|
} |
|
} |
|
numStorageBytes() { |
|
return (this.bits * this.lanes + 7) >> 3; |
|
} |
|
} |
|
exports.DLDataType = DLDataType; |
|
|
|
|
|
|
|
class NDArray { |
|
constructor(handle, isView, lib, ctx) { |
|
this.handle = handle; |
|
this.isView = isView; |
|
this.lib = lib; |
|
this.ctx = ctx; |
|
if (this.isView) { |
|
this.dltensor = handle; |
|
} |
|
else { |
|
this.dltensor = this.getDLTensorFromArrayHandle(this.handle); |
|
} |
|
|
|
const arrayOffsetData = 0; |
|
const arrayOffsetContext = arrayOffsetData + this.lib.sizeofPtr(); |
|
const arrayOffsetDevType = arrayOffsetContext; |
|
const arrayOffsetDevId = arrayOffsetContext + 4 ; |
|
const arrayOffsetNdim = arrayOffsetContext + 8 ; |
|
const arrayOffsetDtype = arrayOffsetNdim + 4 ; |
|
const arrayOffsetDtypeCode = arrayOffsetDtype; |
|
const arrayOffsetDtypeBits = arrayOffsetDtype + 1 ; |
|
const arrayOffsetDtypeLanes = arrayOffsetDtypeBits + 1 ; |
|
const arrayOffsetShape = arrayOffsetDtype + 4 ; |
|
const arrayOffsetStrides = arrayOffsetShape + this.lib.sizeofPtr(); |
|
const arrayOffsetByteOffset = arrayOffsetStrides + this.lib.sizeofPtr(); |
|
|
|
this.dataPtr = lib.memory.loadPointer(this.dltensor); |
|
|
|
this.ndim = lib.memory.loadI32(this.dltensor + arrayOffsetNdim); |
|
|
|
const cshapePtr = lib.memory.loadPointer(this.dltensor + arrayOffsetShape); |
|
this.shape = []; |
|
for (let i = 0; i < this.ndim; ++i) { |
|
this.shape.push(lib.memory.loadI64(cshapePtr + i * 8 )); |
|
} |
|
|
|
const code = lib.memory.loadU8(this.dltensor + arrayOffsetDtypeCode); |
|
const bits = lib.memory.loadU8(this.dltensor + arrayOffsetDtypeBits); |
|
const lanes = lib.memory.loadU16(this.dltensor + arrayOffsetDtypeLanes); |
|
this.dlDataType = new DLDataType(code, bits, lanes); |
|
this.dtype = this.dlDataType.toString(); |
|
|
|
const deviceType = lib.memory.loadI32(this.dltensor + arrayOffsetDevType); |
|
const deviceId = lib.memory.loadI32(this.dltensor + arrayOffsetDevId); |
|
this.device = new DLDevice(deviceType, deviceId, lib); |
|
|
|
this.byteOffset = lib.memory.loadI64(this.dltensor + arrayOffsetByteOffset); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
view(shape) { |
|
const shapeArray = shape.map((value) => new Scalar(value, "int")); |
|
return this.ctx.ndarrayCreateView(this, this.ctx.makeShapeTuple(...shapeArray)); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
getHandle(requireNotNull = true) { |
|
if (requireNotNull && this.handle == 0) { |
|
throw Error("NDArray has already been disposed"); |
|
} |
|
return this.handle; |
|
} |
|
|
|
|
|
|
|
|
|
|
|
getDataPtr() { |
|
if (this.handle == 0) { |
|
throw Error("NDArray has already been disposed"); |
|
} |
|
return this.dataPtr; |
|
} |
|
dispose() { |
|
if (this.handle != 0 && !this.isView) { |
|
this.lib.checkCall(this.lib.exports.TVMArrayFree(this.handle)); |
|
this.handle = 0; |
|
} |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
copyFrom(data) { |
|
if (data instanceof NDArray) { |
|
this.lib.checkCall(this.lib.exports.TVMArrayCopyFromTo(data.getHandle(), this.getHandle(), 0)); |
|
return this; |
|
} |
|
else { |
|
const size = this.shape.reduce((a, b) => { |
|
return a * b; |
|
}, 1); |
|
if (data.length != size) { |
|
throw new Error("data size and shape mismatch data.length" + |
|
data.length + |
|
" vs " + |
|
size); |
|
} |
|
let buffer; |
|
if (this.dtype == "float32") { |
|
buffer = Float32Array.from(data).buffer; |
|
} |
|
else if (this.dtype == "float64") { |
|
buffer = Float64Array.from(data).buffer; |
|
} |
|
else if (this.dtype == "int32") { |
|
buffer = Int32Array.from(data).buffer; |
|
} |
|
else if (this.dtype == "int8") { |
|
buffer = Int8Array.from(data).buffer; |
|
} |
|
else if (this.dtype == "uint8") { |
|
buffer = Uint8Array.from(data).buffer; |
|
} |
|
else { |
|
throw new Error("Unsupported data type " + this.dtype); |
|
} |
|
return this.copyFromRawBytes(new Uint8Array(buffer)); |
|
} |
|
} |
|
|
|
|
|
|
|
|
|
|
|
copyFromRawBytes(data) { |
|
var _a; |
|
|
|
if (this.device.deviceType == DeviceStrToEnum.webgpu) { |
|
(_a = this.lib.webGPUContext) === null || _a === void 0 ? void 0 : _a.copyRawBytesToBuffer(data, this.getDataPtr(), 0, data.length); |
|
return this; |
|
} |
|
|
|
const size = this.shape.reduce((a, b) => { |
|
return a * b; |
|
}, 1); |
|
const nbytes = this.dlDataType.numStorageBytes() * size; |
|
if (nbytes != data.length) { |
|
throw new Error("Expect the data's length equals nbytes=" + nbytes); |
|
} |
|
const stack = this.lib.getOrAllocCallStack(); |
|
const tempOffset = stack.allocRawBytes(nbytes); |
|
const tempPtr = stack.ptrFromOffset(tempOffset); |
|
this.lib.memory.storeRawBytes(tempPtr, data); |
|
this.lib.checkCall(this.lib.exports.TVMArrayCopyFromBytes(this.getHandle(), tempPtr, nbytes)); |
|
this.lib.recycleCallStack(stack); |
|
return this; |
|
} |
|
|
|
|
|
|
|
|
|
toRawBytes() { |
|
if (this.device.deviceType != DeviceStrToEnum.cpu) { |
|
throw new Error("Can only sync copy CPU array, use cpu_arr.copyfrom(gpu_arr) then sync instead."); |
|
} |
|
const size = this.shape.reduce((a, b) => { |
|
return a * b; |
|
}, 1); |
|
const nbytes = this.dlDataType.numStorageBytes() * size; |
|
const stack = this.lib.getOrAllocCallStack(); |
|
const tempOffset = stack.allocRawBytes(nbytes); |
|
const tempPtr = stack.ptrFromOffset(tempOffset); |
|
this.lib.checkCall(this.lib.exports.TVMArrayCopyToBytes(this.getHandle(), tempPtr, nbytes)); |
|
const ret = this.lib.memory.loadRawBytes(tempPtr, nbytes); |
|
this.lib.recycleCallStack(stack); |
|
return ret; |
|
} |
|
|
|
|
|
|
|
|
|
|
|
toArray() { |
|
const stype = this.dtype; |
|
if (stype == "float32") { |
|
return new Float32Array(this.toRawBytes().buffer); |
|
} |
|
else if (stype == "float64") { |
|
return new Float64Array(this.toRawBytes().buffer); |
|
} |
|
else if (stype == "int32") { |
|
return new Int32Array(this.toRawBytes().buffer); |
|
} |
|
else if (stype == "int8") { |
|
return new Int8Array(this.toRawBytes().buffer); |
|
} |
|
else if (stype == "uint8") { |
|
return new Uint8Array(this.toRawBytes().buffer); |
|
} |
|
else { |
|
throw new Error("Unsupported data type " + this.dtype); |
|
} |
|
} |
|
getDLTensorFromArrayHandle(handle) { |
|
|
|
|
|
return handle; |
|
} |
|
} |
|
exports.NDArray = NDArray; |
|
|
|
|
|
|
|
class Module { |
|
constructor(handle, lib, makePackedFunc) { |
|
this.handle = handle; |
|
this.lib = lib; |
|
this.makePackedFunc = makePackedFunc; |
|
} |
|
dispose() { |
|
if (this.handle != 0) { |
|
this.lib.checkCall(this.lib.exports.TVMModFree(this.handle)); |
|
this.handle = 0; |
|
} |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
getHandle(requireNotNull = true) { |
|
if (requireNotNull && this.handle == 0) { |
|
throw Error("Module has already been disposed"); |
|
} |
|
return this.handle; |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
getFunction(name, queryImports = true) { |
|
if (this.handle == 0) { |
|
throw Error("Module has already been disposed"); |
|
} |
|
const stack = this.lib.getOrAllocCallStack(); |
|
const nameOffset = stack.allocRawBytes(name.length + 1); |
|
stack.storeRawBytes(nameOffset, support.StringToUint8Array(name)); |
|
const outOffset = stack.allocPtrArray(1); |
|
const outPtr = stack.ptrFromOffset(outOffset); |
|
stack.commitToWasmMemory(outOffset); |
|
this.lib.checkCall(this.lib.exports.TVMModGetFunction(this.getHandle(), stack.ptrFromOffset(nameOffset), queryImports ? 1 : 0, outPtr)); |
|
const handle = this.lib.memory.loadPointer(outPtr); |
|
this.lib.recycleCallStack(stack); |
|
if (handle == 0) { |
|
throw Error("Cannot find function " + name); |
|
} |
|
const ret = this.makePackedFunc(handle); |
|
return ret; |
|
} |
|
|
|
|
|
|
|
|
|
importModule(mod) { |
|
this.lib.checkCall(this.lib.exports.TVMModImport(this.getHandle(), mod.getHandle())); |
|
} |
|
} |
|
exports.Module = Module; |
|
|
|
|
|
|
|
class TVMObject { |
|
constructor(handle, lib, ctx) { |
|
this.handle = handle; |
|
this.lib = lib; |
|
this.ctx = ctx; |
|
} |
|
dispose() { |
|
if (this.handle != 0) { |
|
this.lib.checkCall(this.lib.exports.TVMObjectFree(this.handle)); |
|
this.handle = 0; |
|
} |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
getHandle(requireNotNull = true) { |
|
if (requireNotNull && this.handle == 0) { |
|
throw Error("Module has already been disposed"); |
|
} |
|
return this.handle; |
|
} |
|
|
|
typeIndex() { |
|
if (this.handle == 0) { |
|
throw Error("The current Object has already been disposed"); |
|
} |
|
const stack = this.lib.getOrAllocCallStack(); |
|
const outOffset = stack.allocPtrArray(1); |
|
const outPtr = stack.ptrFromOffset(outOffset); |
|
this.lib.checkCall(this.lib.exports.TVMObjectGetTypeIndex(this.getHandle(), outPtr)); |
|
const result = this.lib.memory.loadU32(outPtr); |
|
this.lib.recycleCallStack(stack); |
|
return result; |
|
} |
|
|
|
typeKey() { |
|
const type_index = this.typeIndex(); |
|
const stack = this.lib.getOrAllocCallStack(); |
|
const outOffset = stack.allocPtrArray(1); |
|
const outPtr = stack.ptrFromOffset(outOffset); |
|
this.lib.checkCall(this.lib.exports.TVMObjectTypeIndex2Key(type_index, outPtr)); |
|
const result = this.lib.memory.loadCString(this.lib.memory.loadPointer(outPtr)); |
|
this.lib.recycleCallStack(stack); |
|
return result; |
|
} |
|
} |
|
exports.TVMObject = TVMObject; |
|
|
|
class TVMArray extends TVMObject { |
|
constructor(handle, lib, ctx) { |
|
super(handle, lib, ctx); |
|
} |
|
|
|
|
|
|
|
size() { |
|
return this.ctx.arrayGetSize(this); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
get(index) { |
|
return this.ctx.arrayGetItem(this, new Scalar(index, "int32")); |
|
} |
|
} |
|
exports.TVMArray = TVMArray; |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
class VirtualMachine { |
|
|
|
|
|
|
|
|
|
|
|
constructor(mod, device) { |
|
this.mod = mod; |
|
this.mod.getFunction("vm_initialization")(new Scalar(device.deviceType, "int"), new Scalar(device.deviceId, "int"), new Scalar(2 , "int"), |
|
|
|
new Scalar(DeviceStrToEnum.cpu, "int"), new Scalar(0, "int"), new Scalar(2 , "int")); |
|
} |
|
dispose() { |
|
this.mod.dispose(); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
getFunction(name) { |
|
return this.mod.getFunction(name); |
|
} |
|
|
|
|
|
|
|
getInternalModule() { |
|
return this.mod; |
|
} |
|
} |
|
exports.VirtualMachine = VirtualMachine; |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
class Instance { |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
constructor(wasmModule, importObject = {}, wasmInstance, env) { |
|
this.cacheMetadata = {}; |
|
this.initProgressCallback = []; |
|
if (wasmInstance instanceof WebAssembly.Instance) { |
|
support.assert(env instanceof environment.Environment, "env must be provided when passing in instance"); |
|
} |
|
else { |
|
support.assert(env === undefined); |
|
env = new environment.Environment(importObject); |
|
wasmInstance = new WebAssembly.Instance(wasmModule, env.imports); |
|
} |
|
env.start(wasmInstance); |
|
this.env = env; |
|
this.lib = new FFILibrary(wasmInstance, env.imports); |
|
this.memory = this.lib.memory; |
|
this.exports = this.lib.exports; |
|
this.objFactory = new Map(); |
|
this.ctx = new RuntimeContext((name) => { |
|
const autoAttachToScope = false; |
|
|
|
return this.getGlobalFuncInternal(name, autoAttachToScope); |
|
}); |
|
this.registerEnvGlobalPackedFuncs(); |
|
this.registerObjectFactoryFuncs(); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
benchmark(run, dev, number = 10, repeat = 1) { |
|
return __awaiter(this, void 0, void 0, function* () { |
|
|
|
const perf = compact.getPerformance(); |
|
const results = []; |
|
|
|
this.withNewScope(run); |
|
yield dev.sync(); |
|
for (let k = 0; k < repeat; ++k) { |
|
const tstart = perf.now(); |
|
for (let i = 0; i < number; ++i) { |
|
this.withNewScope(run); |
|
} |
|
yield dev.sync(); |
|
const tend = perf.now(); |
|
results.push((tend - tstart) / number); |
|
} |
|
return results; |
|
}); |
|
} |
|
dispose() { |
|
|
|
|
|
this.ctx.dispose(); |
|
this.lib.dispose(); |
|
} |
|
|
|
|
|
|
|
runtimeStatsText() { |
|
if (this.lib.webGPUContext !== undefined) { |
|
return this.lib.webGPUContext.runtimeStatsText(); |
|
} |
|
else { |
|
return ""; |
|
} |
|
} |
|
|
|
|
|
|
|
beginScope() { |
|
this.ctx.beginScope(); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
endScope() { |
|
this.ctx.endScope(); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
withNewScope(action) { |
|
this.beginScope(); |
|
const val = action(); |
|
this.endScope(); |
|
return val; |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
attachToCurrentScope(obj) { |
|
return this.ctx.attachToCurrentScope(obj); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
moveToParentScope(obj) { |
|
return this.ctx.moveToParentScope(obj); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
detachFromCurrentScope(obj) { |
|
return this.ctx.detachFromCurrentScope(obj); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
systemLib() { |
|
return this.ctx.getSysLib(); |
|
} |
|
|
|
|
|
|
|
|
|
listGlobalFuncNames() { |
|
const stack = this.lib.getOrAllocCallStack(); |
|
const outSizeOffset = stack.allocPtrArray(2); |
|
const outSizePtr = stack.ptrFromOffset(outSizeOffset); |
|
const outArrayPtr = stack.ptrFromOffset(outSizeOffset + this.lib.sizeofPtr()); |
|
this.lib.checkCall(this.exports.TVMFuncListGlobalNames(outSizePtr, outArrayPtr)); |
|
const size = this.memory.loadI32(outSizePtr); |
|
const array = this.memory.loadPointer(outArrayPtr); |
|
const names = []; |
|
for (let i = 0; i < size; ++i) { |
|
names.push(this.memory.loadCString(this.memory.loadPointer(array + this.lib.sizeofPtr() * i))); |
|
} |
|
this.lib.recycleCallStack(stack); |
|
return names; |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
registerFunc(name, func, override = false) { |
|
this.withNewScope(() => { |
|
const autoAttachToScope = true; |
|
|
|
const packedFunc = this.toPackedFuncInternal(func, autoAttachToScope); |
|
const ioverride = override ? 1 : 0; |
|
const stack = this.lib.getOrAllocCallStack(); |
|
const nameOffset = stack.allocRawBytes(name.length + 1); |
|
stack.storeRawBytes(nameOffset, support.StringToUint8Array(name)); |
|
stack.commitToWasmMemory(); |
|
this.lib.checkCall(this.lib.exports.TVMFuncRegisterGlobal(stack.ptrFromOffset(nameOffset), packedFunc._tvmPackedCell.getHandle(), ioverride)); |
|
this.lib.recycleCallStack(stack); |
|
}); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
getGlobalFunc(name) { |
|
return this.getGlobalFuncInternal(name, true); |
|
} |
|
getGlobalFuncInternal(name, autoAttachToScope = true) { |
|
const stack = this.lib.getOrAllocCallStack(); |
|
const nameOffset = stack.allocRawBytes(name.length + 1); |
|
stack.storeRawBytes(nameOffset, support.StringToUint8Array(name)); |
|
const outOffset = stack.allocPtrArray(1); |
|
const outPtr = stack.ptrFromOffset(outOffset); |
|
stack.commitToWasmMemory(outOffset); |
|
this.lib.checkCall(this.exports.TVMFuncGetGlobal(stack.ptrFromOffset(nameOffset), outPtr)); |
|
const handle = this.memory.loadPointer(outPtr); |
|
this.lib.recycleCallStack(stack); |
|
if (handle == 0) { |
|
throw Error("Cannot find global function " + name); |
|
} |
|
const ret = this.makePackedFunc(handle); |
|
if (autoAttachToScope) |
|
this.ctx.attachToCurrentScope(ret); |
|
return ret; |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
isPackedFunc(func) { |
|
|
|
return typeof func == "function" && func.hasOwnProperty("_tvmPackedCell"); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
toPackedFunc(func) { |
|
return this.toPackedFuncInternal(func, true); |
|
} |
|
toPackedFuncInternal(func, autoAttachToScope) { |
|
if (this.isPackedFunc(func)) |
|
return func; |
|
const ret = this.createPackedFuncFromCFunc(this.wrapJSFuncAsPackedCFunc(func)); |
|
if (autoAttachToScope) |
|
return this.ctx.attachToCurrentScope(ret); |
|
return ret; |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
createVirtualMachine(dev) { |
|
const mod = this.ctx.detachFromCurrentScope(this.systemLib().getFunction("vm_load_executable")()); |
|
return this.ctx.attachToCurrentScope(new VirtualMachine(mod, dev)); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
registerInitProgressCallback(cb) { |
|
this.initProgressCallback.push(cb); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
getParamsFromCache(prefix, numParams) { |
|
return this.ctx.paramModuleFromCache(prefix, new Scalar(numParams, "int32")).getFunction("get_params")(); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
ndarrayCacheGet(name) { |
|
return this.ctx.arrayCacheGet(name); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
ndarrayCacheRemove(name) { |
|
return this.ctx.arrayCacheRemove(name); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
ndarrayCacheUpdate(name, arr, override = false) { |
|
this.ctx.arrayCacheUpdate(name, arr, this.scalar(override ? 1 : 0, "int32")); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
ndarrayCacheClear() { |
|
this.ctx.arrayCacheClear(); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
fetchNDArrayCache(ndarrayCacheUrl, device) { |
|
return __awaiter(this, void 0, void 0, function* () { |
|
const cacheExists = yield caches.has("tvmjs"); |
|
const jsonUrl = cacheExists |
|
? "/lib/WebLLM/vicuna-7b/ndarray-cache.json" |
|
: new URL("ndarray-cache.json", ndarrayCacheUrl).href; |
|
var list; |
|
try { |
|
list = yield (yield fetch(jsonUrl)).json(); |
|
} |
|
catch (err) { |
|
this.env.logger("Cannot fetch " + jsonUrl); |
|
} |
|
yield this.fetchNDArrayCacheInternal(ndarrayCacheUrl, list["records"], device); |
|
this.cacheMetadata = Object.assign(Object.assign({}, this.cacheMetadata), list["metadata"]); |
|
}); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
fetchNDArrayCacheInternal(ndarrayCacheUrl, list, device) { |
|
return __awaiter(this, void 0, void 0, function* () { |
|
const perf = compact.getPerformance(); |
|
let tstart = perf.now(); |
|
let totalBytes = 0; |
|
for (let i = 0; i < list.length; ++i) { |
|
totalBytes += list[i].nbytes; |
|
} |
|
let fetchedBytes = 0; |
|
let timeElapsed = 0; |
|
const reportCallback = (iter) => { |
|
|
|
for (let j = 0; j < this.initProgressCallback.length; ++j) { |
|
let text = "[" + iter + "/" + list.length + "]: "; |
|
text += Math.ceil(fetchedBytes / (1024 * 1024)).toString() + " MB, "; |
|
text += Math.floor(fetchedBytes * 100 / totalBytes).toString() + "% complete, "; |
|
text += timeElapsed + " secs"; |
|
this.initProgressCallback[j]({ |
|
progress: fetchedBytes / totalBytes, |
|
timeElapsed: timeElapsed, |
|
text: text |
|
}); |
|
} |
|
}; |
|
for (let j = 0; j < this.initProgressCallback.length; ++j) { |
|
this.initProgressCallback[j]({ |
|
progress: fetchedBytes / totalBytes, |
|
timeElapsed: 0, |
|
text: "Start to fetch params", |
|
}); |
|
} |
|
const cache = yield caches.open("tvmjs"); |
|
for (let i = 0; i < list.length; ++i) { |
|
reportCallback(i); |
|
fetchedBytes += list[i].nbytes; |
|
const dataUrl = new URL(list[i].dataPath, ndarrayCacheUrl).href; |
|
const request = new Request(dataUrl); |
|
let buffer; |
|
try { |
|
|
|
let result = yield cache.match(request); |
|
if (result === undefined) { |
|
yield cache.add(request); |
|
result = yield cache.match(request); |
|
} |
|
if (result == undefined) { |
|
this.env.logger("Error: Cannot cache " + dataUrl + ", reloading will be slow"); |
|
result = yield fetch(request); |
|
} |
|
buffer = yield result.arrayBuffer(); |
|
} |
|
catch (err) { |
|
this.env.logger("Error: Cannot fetch " + dataUrl + " err= " + err); |
|
throw err; |
|
} |
|
const shardRecords = list[i].records; |
|
for (let j = 0; j < shardRecords.length; ++j) { |
|
const rec = shardRecords[j]; |
|
const cpu_arr = this.withNewScope(() => { |
|
return this.detachFromCurrentScope(this.empty(rec.shape, rec.dtype, this.cpu())); |
|
}); |
|
const recSource = buffer.slice(rec.byteOffset, rec.byteOffset + rec.nbytes); |
|
|
|
this.ctx.arrayDecodeStorage(cpu_arr, new Uint8Array(recSource), rec.format); |
|
|
|
if (device.deviceType == DeviceStrToEnum.cpu) { |
|
this.ndarrayCacheUpdate(rec.name, cpu_arr, false); |
|
cpu_arr.dispose(); |
|
} |
|
else { |
|
|
|
const gpu_arr = this.withNewScope(() => { |
|
return this.detachFromCurrentScope(this.empty(rec.shape, rec.dtype, device)); |
|
}); |
|
gpu_arr.copyFrom(cpu_arr); |
|
yield device.sync(); |
|
this.ndarrayCacheUpdate(rec.name, gpu_arr, false); |
|
cpu_arr.dispose(); |
|
gpu_arr.dispose(); |
|
} |
|
} |
|
timeElapsed = Math.ceil((perf.now() - tstart) / 1000); |
|
} |
|
reportCallback(list.length); |
|
}); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
toDLDataType(dtype) { |
|
if (dtype instanceof DLDataType) |
|
return dtype; |
|
if (typeof dtype == "string") { |
|
let pattern = dtype; |
|
let code, bits = 32, lanes = 1; |
|
if (pattern.substring(0, 5) == "float") { |
|
pattern = pattern.substring(5, pattern.length); |
|
code = 2 ; |
|
} |
|
else if (pattern.substring(0, 3) == "int") { |
|
pattern = pattern.substring(3, pattern.length); |
|
code = 0 ; |
|
} |
|
else if (pattern.substring(0, 4) == "uint") { |
|
pattern = pattern.substring(4, pattern.length); |
|
code = 1 ; |
|
} |
|
else if (pattern.substring(0, 6) == "handle") { |
|
pattern = pattern.substring(5, pattern.length); |
|
code = 3 ; |
|
bits = 64; |
|
} |
|
else { |
|
throw new Error("Unknown dtype " + dtype); |
|
} |
|
const arr = pattern.split("x"); |
|
if (arr.length >= 1) { |
|
const parsed = parseInt(arr[0]); |
|
if (parsed + "" == arr[0]) { |
|
bits = parsed; |
|
} |
|
} |
|
if (arr.length >= 2) { |
|
lanes = parseInt(arr[1]); |
|
} |
|
return new DLDataType(code, bits, lanes); |
|
} |
|
else { |
|
throw new Error("Unknown dtype " + dtype); |
|
} |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
scalar(value, dtype) { |
|
return new Scalar(value, dtype); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
device(deviceType, deviceId = 0) { |
|
return new DLDevice(deviceType, deviceId, this.lib); |
|
} |
|
|
|
|
|
|
|
|
|
cpu(deviceId = 0) { |
|
return this.device("cpu", deviceId); |
|
} |
|
|
|
|
|
|
|
|
|
webgpu(deviceId = 0) { |
|
return this.device("webgpu", deviceId); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
empty(shape, dtype = "float32", dev = this.device("cpu", 0)) { |
|
dtype = this.toDLDataType(dtype); |
|
shape = typeof shape == "number" ? [shape] : shape; |
|
const stack = this.lib.getOrAllocCallStack(); |
|
const shapeOffset = stack.allocRawBytes(shape.length * 8 ); |
|
for (let i = 0; i < shape.length; ++i) { |
|
stack.storeI64(shapeOffset + i * 8 , shape[i]); |
|
} |
|
const outOffset = stack.allocPtrArray(1); |
|
const outPtr = stack.ptrFromOffset(outOffset); |
|
stack.commitToWasmMemory(outOffset); |
|
this.lib.checkCall(this.exports.TVMArrayAlloc(stack.ptrFromOffset(shapeOffset), shape.length, dtype.code, dtype.bits, dtype.lanes, dev.deviceType, dev.deviceId, outPtr)); |
|
const ret = this.ctx.attachToCurrentScope(new NDArray(this.memory.loadPointer(outPtr), false, this.lib, this.ctx)); |
|
this.lib.recycleCallStack(stack); |
|
return ret; |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
uniform(shape, low, high, dev) { |
|
const ret = this.empty(shape, "float32", dev); |
|
const size = shape.reduce((a, b) => { |
|
return a * b; |
|
}, 1); |
|
const scale = high - low; |
|
const input = new Float32Array(size); |
|
for (let i = 0; i < input.length; ++i) { |
|
input[i] = low + Math.random() * scale; |
|
} |
|
return ret.copyFrom(input); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
sampleTopPFromLogits(logits, temperature, top_p) { |
|
return this.ctx.sampleTopPFromLogits(logits, temperature, top_p, Math.random()); |
|
} |
|
|
|
|
|
|
|
|
|
bindCanvas(canvas) { |
|
var _a; |
|
(_a = this.lib.webGPUContext) === null || _a === void 0 ? void 0 : _a.bindCanvas(canvas); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
showImage(dataRGBA) { |
|
var _a; |
|
if (dataRGBA.shape.length != 2) { |
|
throw Error("Require a height x width uint32 NDArray in RGBA" + |
|
"get shape=" + dataRGBA.shape.toString() + " instead."); |
|
} |
|
if (dataRGBA.device.deviceType != DeviceStrToEnum.webgpu) { |
|
throw new Error("Can only run showImage on WebGPU array, " + |
|
"get " + DeviceEnumToStr[dataRGBA.device.deviceType] + " instead."); |
|
} |
|
if (dataRGBA.dtype != "uint32") { |
|
throw Error("Require a height x width uint32 NDArray in RGBA, " + |
|
"get " + dataRGBA.dtype + " instead."); |
|
} |
|
(_a = this.lib.webGPUContext) === null || _a === void 0 ? void 0 : _a.drawImageFromBuffer(dataRGBA.getDataPtr(), dataRGBA.shape[0], dataRGBA.shape[1]); |
|
} |
|
|
|
|
|
|
|
clearCanvas() { |
|
var _a; |
|
(_a = this.lib.webGPUContext) === null || _a === void 0 ? void 0 : _a.clearCanvas(); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
makeTVMArray(inputs) { |
|
return this.ctx.arrayMake(...inputs); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
makeShapeTuple(shape) { |
|
const shapeArray = shape.map((value) => new Scalar(value, "int")); |
|
return this.ctx.makeShapeTuple(...shapeArray); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
typeKey2Index(typeKey) { |
|
const stack = this.lib.getOrAllocCallStack(); |
|
const typeKeyOffset = stack.allocRawBytes(typeKey.length + 1); |
|
stack.storeRawBytes(typeKeyOffset, support.StringToUint8Array(typeKey)); |
|
const outOffset = stack.allocPtrArray(1); |
|
const outPtr = stack.ptrFromOffset(outOffset); |
|
stack.commitToWasmMemory(outOffset); |
|
this.lib.checkCall(this.lib.exports.TVMObjectTypeKey2Index(stack.ptrFromOffset(typeKeyOffset), outPtr)); |
|
const typeIndex = this.memory.loadU32(outPtr); |
|
this.lib.recycleCallStack(stack); |
|
return typeIndex; |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
registerObjectConstructor(typeKey, func, override = false) { |
|
const typeIndex = this.typeKey2Index(typeKey); |
|
if (this.objFactory.has(typeIndex)) { |
|
if (!override) { |
|
throw new Error("Type " + typeKey + " already registered"); |
|
} |
|
} |
|
this.objFactory.set(typeIndex, func); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
registerAsyncServerFunc(name, func, override = false) { |
|
const asyncVariant = (...args) => { |
|
const fargs = args.slice(0, args.length - 1); |
|
|
|
const callback = this.detachFromCurrentScope(args[args.length - 1]); |
|
const promise = func(...fargs); |
|
promise.then((rv) => { |
|
callback(this.scalar(4 , "int32"), rv); |
|
callback.dispose(); |
|
}); |
|
}; |
|
this.registerFunc("__async." + name, asyncVariant, override); |
|
} |
|
|
|
|
|
|
|
|
|
asyncLoadWebGPUPiplines(mod) { |
|
return __awaiter(this, void 0, void 0, function* () { |
|
if (this.lib.webGPUContext == undefined) |
|
throw Error("WebGPU not initialied"); |
|
const webgpuContext = this.lib.webGPUContext; |
|
this.beginScope(); |
|
const fmap_str = mod.getFunction("webgpu.get_fmap", true)(); |
|
let fmap = JSON.parse(fmap_str); |
|
fmap.length; |
|
const fGetShader = this.detachFromCurrentScope(mod.getFunction("webgpu.get_shader")); |
|
const fUpdatePrebuild = this.detachFromCurrentScope(mod.getFunction("webgpu.update_prebuild")); |
|
this.endScope(); |
|
const perf = compact.getPerformance(); |
|
const tstart = perf.now(); |
|
let tlastReport = tstart; |
|
let finishCounter = 0; |
|
const fmapEntries = Object.entries(fmap); |
|
let allEvents = Promise.resolve(); |
|
for (const [key, finfo] of fmapEntries) { |
|
const code = fGetShader(key); |
|
support.assert(key == finfo.name); |
|
const event = webgpuContext.createShaderAsync(finfo, code).then((func) => { |
|
this.beginScope(); |
|
fUpdatePrebuild(key, func); |
|
this.endScope(); |
|
}).then(() => { |
|
finishCounter += 1; |
|
const tend = perf.now(); |
|
|
|
if ((tend - tlastReport) < 1000 && finishCounter != fmapEntries.length) { |
|
return; |
|
} |
|
tlastReport = tend; |
|
const timeElapsed = Math.ceil((perf.now() - tstart) / 1000); |
|
|
|
for (let j = 0; j < this.initProgressCallback.length; ++j) { |
|
const progress = finishCounter / fmapEntries.length; |
|
let text = "Loading GPU shader modules[" + finishCounter + "/" + fmapEntries.length + "]: "; |
|
text += Math.floor(progress * 100).toString() + "% completed, "; |
|
text += timeElapsed + " secs elapsed."; |
|
this.initProgressCallback[j]({ |
|
progress: progress, |
|
timeElapsed: timeElapsed, |
|
text: text |
|
}); |
|
} |
|
}); |
|
allEvents = Promise.all([allEvents, event]).then(() => { }); |
|
} |
|
yield allEvents; |
|
support.assert(finishCounter == fmapEntries.length); |
|
}); |
|
} |
|
|
|
|
|
|
|
|
|
initWebGPU(device) { |
|
const webGPUContext = new webgpu.WebGPUContext(this.memory, device); |
|
this.registerFunc("wasm.WebGPUDeviceAPI", (name) => { |
|
return webGPUContext.getDeviceAPI(name); |
|
}); |
|
this.registerFunc("wasm.WebGPUCreateShader", (info, code) => { |
|
const finfo = JSON.parse(info); |
|
return webGPUContext.createShader(finfo, code); |
|
}); |
|
this.registerAsyncServerFunc("wasm.WebGPUWaitForTasks", () => __awaiter(this, void 0, void 0, function* () { |
|
yield webGPUContext.sync(); |
|
})); |
|
this.lib.webGPUContext = webGPUContext; |
|
} |
|
|
|
registerObjectFactoryFuncs() { |
|
this.registerObjectConstructor("Array", (handle, lib, ctx) => { |
|
return new TVMArray(handle, lib, ctx); |
|
}); |
|
} |
|
|
|
registerEnvGlobalPackedFuncs() { |
|
|
|
const perf = compact.getPerformance(); |
|
|
|
const timeExecution = (finvoke, dev, nstep, repeat, minRepeatMs, limitZeroTimeIterations, cooldownIntervalMs, repeatsToCooldown) => __awaiter(this, void 0, void 0, function* () { |
|
|
|
|
|
|
|
this.ctx.detachFromCurrentScope(finvoke); |
|
finvoke(this.scalar(1, "int32")); |
|
yield dev.sync(); |
|
const result = []; |
|
let setupNumber = nstep; |
|
for (let i = 0; i < repeat; ++i) { |
|
let durationMs = 0.0; |
|
let absoluteZeroTimes = 0; |
|
do { |
|
if (durationMs > 0.0) { |
|
let golden_ratio = 1.618; |
|
setupNumber = Math.floor(Math.max(minRepeatMs / (durationMs / setupNumber) + 1, setupNumber * golden_ratio)); |
|
} |
|
const tstart = perf.now(); |
|
finvoke(this.scalar(setupNumber, "int32")); |
|
yield dev.sync(); |
|
const tend = perf.now(); |
|
durationMs = tend - tstart; |
|
if (durationMs == 0) { |
|
absoluteZeroTimes++; |
|
} |
|
} while (durationMs < minRepeatMs && absoluteZeroTimes < limitZeroTimeIterations); |
|
const speed = durationMs / setupNumber / 1000; |
|
result.push(speed); |
|
if (cooldownIntervalMs > 0.0 && (i % repeatsToCooldown) == 0) { |
|
yield new Promise(r => setTimeout(r, cooldownIntervalMs)); |
|
} |
|
} |
|
const ret = new Float64Array(result.length); |
|
ret.set(result); |
|
|
|
finvoke.dispose(); |
|
return new Uint8Array(ret.buffer); |
|
}); |
|
const addOne = (x) => __awaiter(this, void 0, void 0, function* () { |
|
yield new Promise(resolve => setTimeout(resolve, 100)); |
|
return x + 1; |
|
}); |
|
this.registerAsyncServerFunc("wasm.TimeExecution", timeExecution); |
|
this.registerAsyncServerFunc("testing.asyncAddOne", addOne); |
|
} |
|
createPackedFuncFromCFunc(func) { |
|
let findex = this.env.packedCFuncTable.length; |
|
if (this.env.packedCFuncTableFreeId.length != 0) { |
|
findex = this.env.packedCFuncTableFreeId.pop(); |
|
} |
|
else { |
|
this.env.packedCFuncTable.push(undefined); |
|
} |
|
this.env.packedCFuncTable[findex] = func; |
|
const stack = this.lib.getOrAllocCallStack(); |
|
const outOffset = stack.allocPtrArray(1); |
|
const outPtr = stack.ptrFromOffset(outOffset); |
|
this.lib.checkCall(this.exports |
|
.TVMWasmFuncCreateFromCFunc(findex, outPtr)); |
|
const ret = this.makePackedFunc(this.memory.loadPointer(outPtr)); |
|
this.lib.recycleCallStack(stack); |
|
return ret; |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
setPackedArguments(stack, args, argsValue, argsCode) { |
|
for (let i = 0; i < args.length; ++i) { |
|
let val = args[i]; |
|
const tp = typeof val; |
|
const valueOffset = argsValue + i * 8 ; |
|
const codeOffset = argsCode + i * 4 ; |
|
if (val instanceof NDArray) { |
|
if (!val.isView) { |
|
stack.storePtr(valueOffset, val.getHandle()); |
|
stack.storeI32(codeOffset, 13 ); |
|
} |
|
else { |
|
stack.storePtr(valueOffset, val.getHandle()); |
|
stack.storeI32(codeOffset, 7 ); |
|
} |
|
} |
|
else if (val instanceof Scalar) { |
|
if (val.dtype.startsWith("int") || val.dtype.startsWith("uint")) { |
|
stack.storeI64(valueOffset, val.value); |
|
stack.storeI32(codeOffset, 0 ); |
|
} |
|
else if (val.dtype.startsWith("float")) { |
|
stack.storeF64(valueOffset, val.value); |
|
stack.storeI32(codeOffset, 2 ); |
|
} |
|
else { |
|
support.assert(val.dtype == "handle", "Expect handle"); |
|
stack.storePtr(valueOffset, val.value); |
|
stack.storeI32(codeOffset, 3 ); |
|
} |
|
} |
|
else if (val instanceof DLDevice) { |
|
stack.storeI32(valueOffset, val.deviceType); |
|
stack.storeI32(valueOffset + 4 , val.deviceType); |
|
stack.storeI32(codeOffset, 6 ); |
|
} |
|
else if (tp == "number") { |
|
stack.storeF64(valueOffset, val); |
|
stack.storeI32(codeOffset, 2 ); |
|
|
|
} |
|
else if (tp == "function" && val.hasOwnProperty("_tvmPackedCell")) { |
|
stack.storePtr(valueOffset, val._tvmPackedCell.getHandle()); |
|
stack.storeI32(codeOffset, 10 ); |
|
} |
|
else if (val === null || val == undefined) { |
|
stack.storePtr(valueOffset, 0); |
|
stack.storeI32(codeOffset, 4 ); |
|
} |
|
else if (tp == "string") { |
|
stack.allocThenSetArgString(valueOffset, val); |
|
stack.storeI32(codeOffset, 11 ); |
|
} |
|
else if (val instanceof Uint8Array) { |
|
stack.allocThenSetArgBytes(valueOffset, val); |
|
stack.storeI32(codeOffset, 12 ); |
|
} |
|
else if (val instanceof Function) { |
|
val = this.toPackedFuncInternal(val, false); |
|
stack.tempArgs.push(val); |
|
stack.storePtr(valueOffset, val._tvmPackedCell.getHandle()); |
|
stack.storeI32(codeOffset, 10 ); |
|
} |
|
else if (val instanceof Module) { |
|
stack.storePtr(valueOffset, val.getHandle()); |
|
stack.storeI32(codeOffset, 9 ); |
|
} |
|
else if (val instanceof TVMObject) { |
|
stack.storePtr(valueOffset, val.getHandle()); |
|
stack.storeI32(codeOffset, 8 ); |
|
} |
|
else { |
|
throw new Error("Unsupported argument type " + tp); |
|
} |
|
} |
|
} |
|
wrapJSFuncAsPackedCFunc(func) { |
|
const lib = this.lib; |
|
return (argValues, argCodes, nargs, ret, |
|
|
|
_handle) => { |
|
const jsArgs = []; |
|
|
|
this.ctx.beginScope(); |
|
for (let i = 0; i < nargs; ++i) { |
|
const valuePtr = argValues + i * 8 ; |
|
const codePtr = argCodes + i * 4 ; |
|
let tcode = lib.memory.loadI32(codePtr); |
|
if (tcode == 8 || |
|
tcode == 14 || |
|
tcode == 10 || |
|
tcode == 13 || |
|
tcode == 9 ) { |
|
lib.checkCall(lib.exports.TVMCbArgToReturn(valuePtr, codePtr)); |
|
} |
|
tcode = lib.memory.loadI32(codePtr); |
|
jsArgs.push(this.retValueToJS(valuePtr, tcode, true)); |
|
} |
|
const rv = func(...jsArgs); |
|
|
|
this.ctx.endScope(); |
|
if (rv !== undefined && rv !== null) { |
|
const stack = lib.getOrAllocCallStack(); |
|
const valueOffset = stack.allocRawBytes(8 ); |
|
const codeOffset = stack.allocRawBytes(4 ); |
|
this.setPackedArguments(stack, [rv], valueOffset, codeOffset); |
|
const valuePtr = stack.ptrFromOffset(valueOffset); |
|
const codePtr = stack.ptrFromOffset(codeOffset); |
|
stack.commitToWasmMemory(); |
|
lib.checkCall(lib.exports.TVMCFuncSetReturn(ret, valuePtr, codePtr, 1)); |
|
lib.recycleCallStack(stack); |
|
} |
|
return 0; |
|
}; |
|
} |
|
makePackedFunc(handle) { |
|
const cell = new PackedFuncCell(handle, this.lib); |
|
const packedFunc = (...args) => { |
|
const stack = this.lib.getOrAllocCallStack(); |
|
const valueOffset = stack.allocRawBytes(8 * args.length); |
|
const tcodeOffset = stack.allocRawBytes(4 * args.length); |
|
this.setPackedArguments(stack, args, valueOffset, tcodeOffset); |
|
const rvalueOffset = stack.allocRawBytes(8 ); |
|
const rcodeOffset = stack.allocRawBytes(4 ); |
|
const rvaluePtr = stack.ptrFromOffset(rvalueOffset); |
|
const rcodePtr = stack.ptrFromOffset(rcodeOffset); |
|
|
|
stack.commitToWasmMemory(rvalueOffset); |
|
this.lib.checkCall(this.exports.TVMFuncCall(cell.getHandle(), stack.ptrFromOffset(valueOffset), stack.ptrFromOffset(tcodeOffset), args.length, rvaluePtr, rcodePtr)); |
|
const ret = this.retValueToJS(rvaluePtr, this.memory.loadI32(rcodePtr), false); |
|
this.lib.recycleCallStack(stack); |
|
return ret; |
|
}; |
|
|
|
|
|
const ret = packedFunc; |
|
ret.dispose = () => { |
|
cell.dispose(); |
|
}; |
|
ret._tvmPackedCell = cell; |
|
return ret; |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
retValueToJS(rvaluePtr, tcode, callbackArg) { |
|
switch (tcode) { |
|
case 0 : |
|
case 1 : |
|
return this.memory.loadI64(rvaluePtr); |
|
case 2 : |
|
return this.memory.loadF64(rvaluePtr); |
|
case 3 : { |
|
return this.memory.loadPointer(rvaluePtr); |
|
} |
|
case 13 : { |
|
return this.ctx.attachToCurrentScope(new NDArray(this.memory.loadPointer(rvaluePtr), false, this.lib, this.ctx)); |
|
} |
|
case 7 : { |
|
support.assert(callbackArg); |
|
|
|
return new NDArray(this.memory.loadPointer(rvaluePtr), true, this.lib, this.ctx); |
|
} |
|
case 10 : { |
|
return this.ctx.attachToCurrentScope(this.makePackedFunc(this.memory.loadPointer(rvaluePtr))); |
|
} |
|
case 9 : { |
|
return this.ctx.attachToCurrentScope(new Module(this.memory.loadPointer(rvaluePtr), this.lib, (ptr) => { |
|
return this.ctx.attachToCurrentScope(this.makePackedFunc(ptr)); |
|
})); |
|
} |
|
case 8 : { |
|
const obj = new TVMObject(this.memory.loadPointer(rvaluePtr), this.lib, this.ctx); |
|
const func = this.objFactory.get(obj.typeIndex()); |
|
if (func != undefined) { |
|
return this.ctx.attachToCurrentScope(func(obj.getHandle(), this.lib, this.ctx)); |
|
} |
|
else { |
|
return this.ctx.attachToCurrentScope(obj); |
|
} |
|
} |
|
case 4 : return undefined; |
|
case 6 : { |
|
const deviceType = this.memory.loadI32(rvaluePtr); |
|
const deviceId = this.memory.loadI32(rvaluePtr + 4 ); |
|
return this.device(deviceType, deviceId); |
|
} |
|
case 11 : { |
|
const ret = this.memory.loadCString(this.memory.loadPointer(rvaluePtr)); |
|
return ret; |
|
} |
|
case 12 : { |
|
return this.memory.loadTVMBytes(this.memory.loadPointer(rvaluePtr)); |
|
} |
|
default: |
|
throw new Error("Unsupported return type code=" + tcode); |
|
} |
|
} |
|
} |
|
exports.Instance = Instance; |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
function instantiate(bufferSource, importObject = {}, logger = console.log) { |
|
const env = new environment.Environment(importObject, logger); |
|
return WebAssembly.instantiate(bufferSource, env.imports).then((result) => { |
|
return new Instance(result.module, {}, result.instance, env); |
|
}); |
|
} |
|
exports.instantiate = instantiate; |
|
|
|
}); |
|
|
|
unwrapExports(runtime); |
|
runtime.instantiate; |
|
runtime.Instance; |
|
runtime.VirtualMachine; |
|
runtime.TVMArray; |
|
runtime.TVMObject; |
|
runtime.Module; |
|
runtime.NDArray; |
|
runtime.DLDataType; |
|
runtime.DLDevice; |
|
runtime.Scalar; |
|
|
|
var rpc_server = createCommonjsModule(function (module, exports) { |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
var __awaiter = (commonjsGlobal && commonjsGlobal.__awaiter) || function (thisArg, _arguments, P, generator) { |
|
function adopt(value) { return value instanceof P ? value : new P(function (resolve) { resolve(value); }); } |
|
return new (P || (P = Promise))(function (resolve, reject) { |
|
function fulfilled(value) { try { step(generator.next(value)); } catch (e) { reject(e); } } |
|
function rejected(value) { try { step(generator["throw"](value)); } catch (e) { reject(e); } } |
|
function step(result) { result.done ? resolve(result.value) : adopt(result.value).then(fulfilled, rejected); } |
|
step((generator = generator.apply(thisArg, _arguments || [])).next()); |
|
}); |
|
}; |
|
Object.defineProperty(exports, "__esModule", { value: true }); |
|
exports.RPCServer = void 0; |
|
|
|
|
|
|
|
|
|
var RPCServerState; |
|
(function (RPCServerState) { |
|
RPCServerState[RPCServerState["InitHeader"] = 0] = "InitHeader"; |
|
RPCServerState[RPCServerState["InitHeaderKey"] = 1] = "InitHeaderKey"; |
|
RPCServerState[RPCServerState["InitServer"] = 2] = "InitServer"; |
|
RPCServerState[RPCServerState["WaitForCallback"] = 3] = "WaitForCallback"; |
|
RPCServerState[RPCServerState["ReceivePacketHeader"] = 4] = "ReceivePacketHeader"; |
|
RPCServerState[RPCServerState["ReceivePacketBody"] = 5] = "ReceivePacketBody"; |
|
})(RPCServerState || (RPCServerState = {})); |
|
|
|
const RPC_MAGIC = 0xff271; |
|
|
|
|
|
|
|
class ByteStreamReader { |
|
constructor(bytes) { |
|
this.offset = 0; |
|
this.bytes = bytes; |
|
} |
|
readU32() { |
|
const i = this.offset; |
|
const b = this.bytes; |
|
const val = b[i] | (b[i + 1] << 8) | (b[i + 2] << 16) | (b[i + 3] << 24); |
|
this.offset += 4; |
|
return val; |
|
} |
|
readU64() { |
|
const val = this.readU32(); |
|
this.offset += 4; |
|
return val; |
|
} |
|
readByteArray() { |
|
const len = this.readU64(); |
|
support.assert(this.offset + len <= this.bytes.byteLength); |
|
const ret = new Uint8Array(len); |
|
ret.set(this.bytes.slice(this.offset, this.offset + len)); |
|
this.offset += len; |
|
return ret; |
|
} |
|
} |
|
|
|
|
|
|
|
class RPCServer { |
|
constructor(url, key, getImports, logger = console.log, ndarrayCacheUrl = "", ndarrayCacheDevice = "cpu", initProgressCallback = undefined, asyncOnServerLoad = undefined) { |
|
this.state = RPCServerState.InitHeader; |
|
this.pendingSend = Promise.resolve(); |
|
this.inst = undefined; |
|
this.globalObjects = []; |
|
this.currPacketLength = 0; |
|
this.remoteKeyLength = 0; |
|
this.pendingBytes = 0; |
|
this.buffredBytes = 0; |
|
this.messageQueue = []; |
|
this.url = url; |
|
this.key = key; |
|
this.name = "WebSocketRPCServer[" + this.key + "]: "; |
|
this.getImports = getImports; |
|
this.logger = logger; |
|
this.ndarrayCacheUrl = ndarrayCacheUrl; |
|
this.ndarrayCacheDevice = ndarrayCacheDevice; |
|
this.initProgressCallback = initProgressCallback; |
|
this.asyncOnServerLoad = asyncOnServerLoad; |
|
this.checkLittleEndian(); |
|
this.socket = compact.createWebSocket(url); |
|
this.socket.binaryType = "arraybuffer"; |
|
this.socket.addEventListener("open", (event) => { |
|
return this.onOpen(event); |
|
}); |
|
this.socket.addEventListener("message", (event) => { |
|
return this.onMessage(event); |
|
}); |
|
this.socket.addEventListener("close", (event) => { |
|
return this.onClose(event); |
|
}); |
|
} |
|
|
|
onClose(_event) { |
|
if (this.inst !== undefined) { |
|
this.globalObjects.forEach(obj => { |
|
obj.dispose(); |
|
}); |
|
this.log(this.inst.runtimeStatsText()); |
|
this.inst.dispose(); |
|
} |
|
if (this.state == RPCServerState.ReceivePacketHeader) { |
|
this.log("Closing the server in clean state"); |
|
this.log("Automatic reconnecting.."); |
|
new RPCServer(this.url, this.key, this.getImports, this.logger, this.ndarrayCacheUrl, this.ndarrayCacheDevice, this.initProgressCallback, this.asyncOnServerLoad); |
|
} |
|
else { |
|
this.log("Closing the server, final state=" + this.state); |
|
} |
|
} |
|
|
|
onOpen(_event) { |
|
|
|
let bkey = support.StringToUint8Array("server:" + this.key); |
|
bkey = bkey.slice(0, bkey.length - 1); |
|
const intbuf = new Int32Array(1); |
|
intbuf[0] = RPC_MAGIC; |
|
this.socket.send(intbuf); |
|
intbuf[0] = bkey.length; |
|
this.socket.send(intbuf); |
|
this.socket.send(bkey); |
|
this.log("connected..."); |
|
|
|
this.requestBytes(4 + 4 ); |
|
this.state = RPCServerState.InitHeader; |
|
} |
|
|
|
onMessage(event) { |
|
const buffer = event.data; |
|
this.buffredBytes += buffer.byteLength; |
|
this.messageQueue.push(new Uint8Array(buffer)); |
|
this.processEvents(); |
|
} |
|
|
|
processEvents() { |
|
while (this.buffredBytes >= this.pendingBytes && this.pendingBytes != 0) { |
|
this.onDataReady(); |
|
} |
|
} |
|
|
|
onDataReady() { |
|
switch (this.state) { |
|
case RPCServerState.InitHeader: { |
|
this.handleInitHeader(); |
|
break; |
|
} |
|
case RPCServerState.InitHeaderKey: { |
|
this.handleInitHeaderKey(); |
|
break; |
|
} |
|
case RPCServerState.ReceivePacketHeader: { |
|
this.currPacketHeader = this.readFromBuffer(8 ); |
|
const reader = new ByteStreamReader(this.currPacketHeader); |
|
this.currPacketLength = reader.readU64(); |
|
support.assert(this.pendingBytes == 0); |
|
this.requestBytes(this.currPacketLength); |
|
this.state = RPCServerState.ReceivePacketBody; |
|
break; |
|
} |
|
case RPCServerState.ReceivePacketBody: { |
|
const body = this.readFromBuffer(this.currPacketLength); |
|
support.assert(this.pendingBytes == 0); |
|
support.assert(this.currPacketHeader !== undefined); |
|
this.onPacketReady(this.currPacketHeader, body); |
|
break; |
|
} |
|
case RPCServerState.WaitForCallback: { |
|
support.assert(this.pendingBytes == 0); |
|
break; |
|
} |
|
default: { |
|
throw new Error("Cannot handle state " + this.state); |
|
} |
|
} |
|
} |
|
onPacketReady(header, body) { |
|
if (this.inst === undefined) { |
|
|
|
const reader = new ByteStreamReader(body); |
|
|
|
reader.readU32(); |
|
|
|
support.Uint8ArrayToString(reader.readByteArray()); |
|
const nargs = reader.readU32(); |
|
const tcodes = []; |
|
const args = []; |
|
for (let i = 0; i < nargs; ++i) { |
|
tcodes.push(reader.readU32()); |
|
} |
|
for (let i = 0; i < nargs; ++i) { |
|
const tcode = tcodes[i]; |
|
if (tcode == 11 ) { |
|
const str = support.Uint8ArrayToString(reader.readByteArray()); |
|
args.push(str); |
|
} |
|
else if (tcode == 12 ) { |
|
args.push(reader.readByteArray()); |
|
} |
|
else { |
|
throw new Error("cannot support type code " + tcode); |
|
} |
|
} |
|
this.onInitServer(args, header, body); |
|
} |
|
else { |
|
support.assert(this.serverRecvData !== undefined); |
|
this.serverRecvData(header, body); |
|
this.requestBytes(8 ); |
|
this.state = RPCServerState.ReceivePacketHeader; |
|
} |
|
} |
|
|
|
onInitServer(args, header, body) { |
|
|
|
support.assert(args[0] == "rpc.WasmSession"); |
|
support.assert(this.pendingBytes == 0); |
|
const asyncInitServer = () => __awaiter(this, void 0, void 0, function* () { |
|
support.assert(args[1] instanceof Uint8Array); |
|
const inst = yield runtime.instantiate(args[1].buffer, this.getImports(), this.logger); |
|
try { |
|
const output = yield webgpu.detectGPUDevice(); |
|
if (output !== undefined) { |
|
const label = "WebGPU: " + output.adapterInfo.description; |
|
this.log("Initialize GPU device: " + label); |
|
inst.initWebGPU(output.device); |
|
} |
|
else { |
|
this.log("Cannot find WebGPU device in the env"); |
|
} |
|
} |
|
catch (err) { |
|
this.log("Cannnot initialize WebGPU, " + err.toString()); |
|
} |
|
this.inst = inst; |
|
|
|
this.inst.beginScope(); |
|
if (this.initProgressCallback !== undefined) { |
|
this.inst.registerInitProgressCallback(this.initProgressCallback); |
|
} |
|
if (this.ndarrayCacheUrl.length != 0) { |
|
if (this.ndarrayCacheDevice == "cpu") { |
|
yield this.inst.fetchNDArrayCache(this.ndarrayCacheUrl, this.inst.cpu()); |
|
} |
|
else { |
|
support.assert(this.ndarrayCacheDevice == "webgpu"); |
|
yield this.inst.fetchNDArrayCache(this.ndarrayCacheUrl, this.inst.webgpu()); |
|
} |
|
} |
|
support.assert(this.inst !== undefined); |
|
if (this.asyncOnServerLoad !== undefined) { |
|
yield this.asyncOnServerLoad(this.inst); |
|
} |
|
const fcreate = this.inst.getGlobalFunc("rpc.CreateEventDrivenServer"); |
|
const messageHandler = fcreate((cbytes) => { |
|
support.assert(this.inst !== undefined); |
|
if (this.socket.readyState == 1) { |
|
|
|
|
|
|
|
const sendDataWithCongestionControl = () => __awaiter(this, void 0, void 0, function* () { |
|
const packetSize = 4 << 10; |
|
const maxBufferAmount = 4 * packetSize; |
|
const waitTimeMs = 20; |
|
for (let offset = 0; offset < cbytes.length; offset += packetSize) { |
|
const end = Math.min(offset + packetSize, cbytes.length); |
|
while (this.socket.bufferedAmount >= maxBufferAmount) { |
|
yield new Promise((r) => setTimeout(r, waitTimeMs)); |
|
} |
|
this.socket.send(cbytes.slice(offset, end)); |
|
} |
|
}); |
|
|
|
this.pendingSend = this.pendingSend.then(sendDataWithCongestionControl); |
|
|
|
return this.inst.scalar(cbytes.length, "int32"); |
|
} |
|
else { |
|
return this.inst.scalar(0, "int32"); |
|
} |
|
}, this.name, this.key); |
|
|
|
this.globalObjects.push(this.inst.detachFromCurrentScope(messageHandler)); |
|
const writeFlag = this.inst.scalar(3, "int32"); |
|
this.serverRecvData = (header, body) => { |
|
if (messageHandler(header, writeFlag) == 0) { |
|
this.socket.close(); |
|
} |
|
if (messageHandler(body, writeFlag) == 0) { |
|
this.socket.close(); |
|
} |
|
}; |
|
|
|
|
|
|
|
|
|
const flocal = this.inst.getGlobalFunc("wasm.LocalSession"); |
|
const localSession = flocal(); |
|
support.assert(localSession instanceof runtime.Module); |
|
|
|
this.inst.registerFunc("rpc.WasmSession", |
|
|
|
(_args) => { |
|
return localSession; |
|
}); |
|
messageHandler(header, writeFlag); |
|
messageHandler(body, writeFlag); |
|
this.log("Finish initializing the Wasm Server.."); |
|
this.requestBytes(8 ); |
|
this.state = RPCServerState.ReceivePacketHeader; |
|
|
|
this.processEvents(); |
|
|
|
this.inst.endScope(); |
|
}); |
|
this.state = RPCServerState.WaitForCallback; |
|
asyncInitServer(); |
|
} |
|
log(msg) { |
|
this.logger(this.name + msg); |
|
} |
|
handleInitHeader() { |
|
const reader = new ByteStreamReader(this.readFromBuffer(4 * 2)); |
|
const magic = reader.readU32(); |
|
if (magic == RPC_MAGIC + 1) { |
|
throw new Error("key: " + this.key + " has already been used in proxy"); |
|
} |
|
else if (magic == RPC_MAGIC + 2) { |
|
throw new Error("RPCProxy do not have matching client key " + this.key); |
|
} |
|
support.assert(magic == RPC_MAGIC, this.url + " is not an RPC Proxy"); |
|
this.remoteKeyLength = reader.readU32(); |
|
support.assert(this.pendingBytes == 0); |
|
this.requestBytes(this.remoteKeyLength); |
|
this.state = RPCServerState.InitHeaderKey; |
|
} |
|
handleInitHeaderKey() { |
|
|
|
support.Uint8ArrayToString(this.readFromBuffer(this.remoteKeyLength)); |
|
support.assert(this.pendingBytes == 0); |
|
this.requestBytes(8 ); |
|
this.state = RPCServerState.ReceivePacketHeader; |
|
} |
|
checkLittleEndian() { |
|
const a = new ArrayBuffer(4); |
|
const b = new Uint8Array(a); |
|
const c = new Uint32Array(a); |
|
b[0] = 0x11; |
|
b[1] = 0x22; |
|
b[2] = 0x33; |
|
b[3] = 0x44; |
|
support.assert(c[0] === 0x44332211, "RPCServer little endian to work"); |
|
} |
|
requestBytes(nbytes) { |
|
this.pendingBytes += nbytes; |
|
} |
|
readFromBuffer(nbytes) { |
|
const ret = new Uint8Array(nbytes); |
|
let ptr = 0; |
|
while (ptr < nbytes) { |
|
support.assert(this.messageQueue.length != 0); |
|
const nleft = nbytes - ptr; |
|
if (this.messageQueue[0].byteLength <= nleft) { |
|
const buffer = this.messageQueue.shift(); |
|
ret.set(buffer, ptr); |
|
ptr += buffer.byteLength; |
|
} |
|
else { |
|
const buffer = this.messageQueue[0]; |
|
ret.set(buffer.slice(0, nleft), ptr); |
|
this.messageQueue[0] = buffer.slice(nleft, buffer.byteLength); |
|
ptr += nleft; |
|
} |
|
} |
|
this.buffredBytes -= nbytes; |
|
this.pendingBytes -= nbytes; |
|
return ret; |
|
} |
|
} |
|
exports.RPCServer = RPCServer; |
|
|
|
}); |
|
|
|
unwrapExports(rpc_server); |
|
rpc_server.RPCServer; |
|
|
|
var dist = createCommonjsModule(function (module, exports) { |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
Object.defineProperty(exports, "__esModule", { value: true }); |
|
|
|
Object.defineProperty(exports, "Scalar", { enumerable: true, get: function () { return runtime.Scalar; } }); |
|
Object.defineProperty(exports, "DLDevice", { enumerable: true, get: function () { return runtime.DLDevice; } }); |
|
Object.defineProperty(exports, "DLDataType", { enumerable: true, get: function () { return runtime.DLDataType; } }); |
|
Object.defineProperty(exports, "Module", { enumerable: true, get: function () { return runtime.Module; } }); |
|
Object.defineProperty(exports, "NDArray", { enumerable: true, get: function () { return runtime.NDArray; } }); |
|
Object.defineProperty(exports, "TVMArray", { enumerable: true, get: function () { return runtime.TVMArray; } }); |
|
Object.defineProperty(exports, "Instance", { enumerable: true, get: function () { return runtime.Instance; } }); |
|
Object.defineProperty(exports, "instantiate", { enumerable: true, get: function () { return runtime.instantiate; } }); |
|
|
|
Object.defineProperty(exports, "RPCServer", { enumerable: true, get: function () { return rpc_server.RPCServer; } }); |
|
|
|
Object.defineProperty(exports, "wasmPath", { enumerable: true, get: function () { return support.wasmPath; } }); |
|
|
|
Object.defineProperty(exports, "detectGPUDevice", { enumerable: true, get: function () { return webgpu.detectGPUDevice; } }); |
|
var support_2 = support; |
|
Object.defineProperty(exports, "assert", { enumerable: true, get: function () { return support_2.assert; } }); |
|
|
|
}); |
|
|
|
var index = unwrapExports(dist); |
|
var dist_1 = dist.Scalar; |
|
var dist_2 = dist.DLDevice; |
|
var dist_3 = dist.DLDataType; |
|
var dist_4 = dist.Module; |
|
var dist_5 = dist.NDArray; |
|
var dist_6 = dist.TVMArray; |
|
var dist_7 = dist.Instance; |
|
var dist_8 = dist.instantiate; |
|
var dist_9 = dist.RPCServer; |
|
var dist_10 = dist.wasmPath; |
|
var dist_11 = dist.detectGPUDevice; |
|
var dist_12 = dist.assert; |
|
|
|
exports.DLDataType = dist_3; |
|
exports.DLDevice = dist_2; |
|
exports.Instance = dist_7; |
|
exports.Module = dist_4; |
|
exports.NDArray = dist_5; |
|
exports.RPCServer = dist_9; |
|
exports.Scalar = dist_1; |
|
exports.TVMArray = dist_6; |
|
exports.assert = dist_12; |
|
exports["default"] = index; |
|
exports.detectGPUDevice = dist_11; |
|
exports.instantiate = dist_8; |
|
exports.wasmPath = dist_10; |
|
|
|
Object.defineProperty(exports, '__esModule', { value: true }); |
|
|
|
})); |
|
|