This is an automated email from the ASF dual-hosted git repository.
junrushao pushed a commit to branch unity
in repository https://gitbox.apache.org/repos/asf/tvm.git
The following commit(s) were added to refs/heads/unity by this push:
new ee5a2c204d [Unity][WEB] Support async pipeline creation (#14362)
ee5a2c204d is described below
commit ee5a2c204d9072cef76e5045d45a98c3c25e7ad1
Author: Tianqi Chen <[email protected]>
AuthorDate: Tue Mar 21 23:41:11 2023 -0400
[Unity][WEB] Support async pipeline creation (#14362)
This PR enables async pipeline creation in webgpu module loading.
This will enable us to report progress in shader compilation
and leverage multi-threading on the host to compile shaders.
---
src/runtime/relax_vm/vm.cc | 32 ++++--
web/apps/browser/rpc_server.html | 6 +-
web/emcc/webgpu_runtime.cc | 32 +++++-
web/src/rpc_server.ts | 12 +--
web/src/runtime.ts | 109 ++++++++++++++++----
web/src/webgpu.ts | 210 ++++++++++++++++++++++++---------------
6 files changed, 284 insertions(+), 117 deletions(-)
diff --git a/src/runtime/relax_vm/vm.cc b/src/runtime/relax_vm/vm.cc
index 6088833dc5..33a19c0482 100644
--- a/src/runtime/relax_vm/vm.cc
+++ b/src/runtime/relax_vm/vm.cc
@@ -188,7 +188,9 @@ class VirtualMachineImpl : public VirtualMachine {
PackedFunc GetFunction(const std::string& name, const ObjectPtr<Object>&
sptr_to_self) override;
- VMClosure GetClosure(const String& func_name) final;
+ VMClosure GetClosure(const String& func_name) final {
+ return this->GetClosureInternal(func_name, false).value();
+ }
void InvokeClosurePacked(const ObjectRef& closure_or_packedfunc, TVMArgs
args,
TVMRetValue* rv) final;
@@ -198,6 +200,14 @@ class VirtualMachineImpl : public VirtualMachine {
//--------------------------------------------------
// Additional support arguments functions for VM
//--------------------------------------------------
+ /*!
+ * \brief Internal implementation of GetClosure which also allow none.
+ * \param func_name The name of the function.
+ * \param allow_missing Whether none is allowed.
+ * \return The result
+ */
+ Optional<VMClosure> GetClosureInternal(const String& func_name, bool
allow_missing);
+
/*!
* \brief Set inputs to a function.
* \param func_name The function name.
@@ -550,10 +560,14 @@ PackedFunc VirtualMachineImpl::GetFunction(const
std::string& name,
});
} else {
// default case, look up closure in VM.
- VMClosure clo = this->GetClosure(name);
- return PackedFunc([sptr_to_self, this, clo](TVMArgs args, TVMRetValue* rv)
{
- this->InvokeClosurePacked(clo, args, rv);
- });
+ if (Optional<VMClosure> opt = this->GetClosureInternal(name, true)) {
+ auto clo = opt.value();
+ return PackedFunc([sptr_to_self, this, clo](TVMArgs args, TVMRetValue*
rv) {
+ this->InvokeClosurePacked(clo, args, rv);
+ });
+ } else {
+ return PackedFunc(nullptr);
+ }
}
}
@@ -653,14 +667,18 @@ void VirtualMachineImpl::SaveClosure(const String&
func_name, const String& save
saved_closures_[save_name] = VMClosure(save_name, impl);
}
-VMClosure VirtualMachineImpl::GetClosure(const String& func_name) {
+Optional<VMClosure> VirtualMachineImpl::GetClosureInternal(const String&
func_name,
+ bool allow_missing)
{
// look up saved closures.
auto saved_it = saved_closures_.find(func_name);
if (saved_it != saved_closures_.end()) {
return saved_it->second;
}
auto it = exec_->func_map.find(func_name);
- CHECK(it != exec_->func_map.end()) << "ValueError: Unknown function: " <<
func_name;
+ if (it == exec_->func_map.end()) {
+ if (allow_missing) return NullOpt;
+ LOG(FATAL) << "ValueError: Unknown function: " << func_name;
+ }
Index gf_idx = it->second;
const VMFuncInfo& finfo = exec_->func_table[gf_idx];
diff --git a/web/apps/browser/rpc_server.html b/web/apps/browser/rpc_server.html
index 4a7b0864e1..07e6fe87fc 100644
--- a/web/apps/browser/rpc_server.html
+++ b/web/apps/browser/rpc_server.html
@@ -43,9 +43,9 @@
}
}
- function fetchProgressCallback(report) {
+ function initProgressCallback(report) {
document.getElementById("rpc-progress-tracker-label").innerHTML =
report.text;
- document.getElementById("rpc-progress-tracker-progress").value =
(report.fetchedBytes / report.totalBytes) * 100;
+ document.getElementById("rpc-progress-tracker-progress").value =
report.progress * 100;
}
function connectRPC() {
@@ -66,7 +66,7 @@
new tvmjs.RPCServer(
proxyUrl, key, getImports, customLog,
- ndarrayCacheUrl, ndarrayCacheDevice, fetchProgressCallback,
+ ndarrayCacheUrl, ndarrayCacheDevice, initProgressCallback,
tvmjsGlobalEnv.asyncOnRPCServerLoad);
}
diff --git a/web/emcc/webgpu_runtime.cc b/web/emcc/webgpu_runtime.cc
index 17efcc8c70..b9a13c9645 100644
--- a/web/emcc/webgpu_runtime.cc
+++ b/web/emcc/webgpu_runtime.cc
@@ -160,6 +160,34 @@ class WebGPUModuleNode final : public runtime::ModuleNode {
const char* type_key() const final { return "webgpu"; }
PackedFunc GetFunction(const std::string& name, const ObjectPtr<Object>&
sptr_to_self) final {
+ // special function
+ if (name == "webgpu.get_fmap") {
+ return PackedFunc([this](TVMArgs args, TVMRetValue* rv) {
+ std::ostringstream os;
+ dmlc::JSONWriter writer(&os);
+ writer.Write(fmap_);
+ *rv = os.str();
+ });
+ } else if (name == "webgpu.get_shader") {
+ return PackedFunc([this](TVMArgs args, TVMRetValue* rv) {
+ std::string name = args[0];
+ auto it = smap_.find(name);
+ ICHECK(it != smap_.end()) << "Cannot find code " << name;
+ *rv = it->second;
+ });
+ } else if (name == "webgpu.update_prebuild") {
+ return PackedFunc([this](TVMArgs args, TVMRetValue* rv) {
+ std::string name = args[0];
+ PackedFunc func = args[1];
+ prebuild_[name] = func;
+ });
+ }
+ // check prebuild cache
+ auto prebuild_it = prebuild_.find(name);
+ if (prebuild_it != prebuild_.end()) {
+ return prebuild_it->second;
+ }
+
auto it = smap_.find(name);
if (it != smap_.end()) {
FunctionInfo info = fmap_.at(name);
@@ -185,12 +213,14 @@ class WebGPUModuleNode final : public runtime::ModuleNode
{
}
private:
- // function information table.
+ // code table
std::unordered_map<std::string, std::string> smap_;
// function information table.
std::unordered_map<std::string, FunctionInfo> fmap_;
// The source
std::string source_;
+ // prebuild_ functions
+ std::unordered_map<std::string, PackedFunc> prebuild_;
// Callback to get the GPU function.
TypedPackedFunc<PackedFunc(std::string finfo, std::string shader)>
create_shader_;
};
diff --git a/web/src/rpc_server.ts b/web/src/rpc_server.ts
index 24acb7ece4..22c4b4617e 100644
--- a/web/src/rpc_server.ts
+++ b/web/src/rpc_server.ts
@@ -83,7 +83,7 @@ export class RPCServer {
getImports: () => Record<string, unknown>;
private ndarrayCacheUrl: string;
private ndarrayCacheDevice: string;
- private fetchProgressCallback?: runtime.FetchProgressCallback;
+ private initProgressCallback?: runtime.InitProgressCallback;
private asyncOnServerLoad?: (inst: runtime.Instance) => Promise<void>;
private pendingSend: Promise<void> = Promise.resolve();
private name: string;
@@ -104,7 +104,7 @@ export class RPCServer {
logger: (msg: string) => void = console.log,
ndarrayCacheUrl: string = "",
ndarrayCacheDevice: string = "cpu",
- fetchProgressCallback: runtime.FetchProgressCallback | undefined =
undefined,
+ initProgressCallback: runtime.InitProgressCallback | undefined = undefined,
asyncOnServerLoad: ((inst: runtime.Instance) => Promise<void>) | undefined
= undefined,
) {
this.url = url;
@@ -114,7 +114,7 @@ export class RPCServer {
this.logger = logger;
this.ndarrayCacheUrl = ndarrayCacheUrl;
this.ndarrayCacheDevice = ndarrayCacheDevice;
- this.fetchProgressCallback = fetchProgressCallback;
+ this.initProgressCallback = initProgressCallback;
this.asyncOnServerLoad = asyncOnServerLoad;
this.checkLittleEndian();
this.socket = compact.createWebSocket(url);
@@ -146,7 +146,7 @@ export class RPCServer {
new RPCServer(
this.url, this.key, this.getImports, this.logger,
this.ndarrayCacheUrl, this.ndarrayCacheDevice,
- this.fetchProgressCallback, this.asyncOnServerLoad);
+ this.initProgressCallback, this.asyncOnServerLoad);
} else {
this.log("Closing the server, final state=" + this.state);
}
@@ -288,8 +288,8 @@ export class RPCServer {
this.inst = inst;
// begin scope to allow handling of objects
this.inst.beginScope();
- if (this.fetchProgressCallback !== undefined) {
- this.inst.registerFetchProgressCallback(this.fetchProgressCallback);
+ if (this.initProgressCallback !== undefined) {
+ this.inst.registerInitProgressCallback(this.initProgressCallback);
}
if (this.ndarrayCacheUrl.length != 0) {
diff --git a/web/src/runtime.ts b/web/src/runtime.ts
index 4550ba9837..f3a6029bbe 100644
--- a/web/src/runtime.ts
+++ b/web/src/runtime.ts
@@ -25,7 +25,7 @@ import { Disposable } from "./types";
import { Memory, CachedCallStack } from "./memory";
import { assert, StringToUint8Array } from "./support";
import { Environment } from "./environment";
-import { WebGPUContext } from "./webgpu";
+import { FunctionInfo, WebGPUContext } from "./webgpu";
import * as compact from "./compact";
import * as ctypes from "./ctypes";
@@ -686,9 +686,10 @@ export class Module implements Disposable {
/**
* Get a function in the module.
* @param name The name of the function.
+ * @param queryImports Whether to also query imports
* @returns The result function.
*/
- getFunction(name: string): PackedFunc {
+ getFunction(name: string, queryImports: boolean = true): PackedFunc {
if (this.handle == 0) {
throw Error("Module has already been disposed");
}
@@ -705,7 +706,7 @@ export class Module implements Disposable {
(this.lib.exports.TVMModGetFunction as ctypes.FTVMModGetFunction)(
this.getHandle(),
stack.ptrFromOffset(nameOffset),
- 1,
+ queryImports? 1 : 0,
outPtr
)
);
@@ -883,6 +884,13 @@ export class VirtualMachine implements Disposable {
getFunction(name: string): PackedFunc {
return this.mod.getFunction(name);
}
+
+ /**
+ * Get the internal module.
+ */
+ getInternalModule(): Module {
+ return this.mod;
+ }
}
/** Code used as the first argument of the async callback. */
@@ -906,14 +914,13 @@ export interface NDArrayShardEntry {
records: Array<NDArrayCacheEntry>;
}
-export interface FetchProgressReport {
- fetchedBytes: number;
- totalBytes: number;
+export interface InitProgressReport {
+ progress: number;
timeElapsed: number;
text: string;
}
-export type FetchProgressCallback = (report: FetchProgressReport) => void;
+export type InitProgressCallback = (report: InitProgressReport) => void;
/**
* TVM runtime instance.
@@ -940,7 +947,7 @@ export class Instance implements Disposable {
private env: Environment;
private objFactory: Map<number, FObjectConstructor>;
private ctx: RuntimeContext;
- private fetchProgressCallback: Array<FetchProgressCallback> = [];
+ private initProgressCallback: Array<InitProgressCallback> = [];
/**
* Internal function(registered by the runtime)
@@ -1281,8 +1288,8 @@ export class Instance implements Disposable {
*
* @param cb the fetch progress callback.
*/
- registerFetchProgressCallback(cb: FetchProgressCallback) {
- this.fetchProgressCallback.push(cb);
+ registerInitProgressCallback(cb: InitProgressCallback) {
+ this.initProgressCallback.push(cb);
}
/**
@@ -1374,26 +1381,24 @@ export class Instance implements Disposable {
const reportCallback = (iter: number)=> {
// report
- for (let j = 0; j < this.fetchProgressCallback.length; ++j) {
+ for (let j = 0; j < this.initProgressCallback.length; ++j) {
let text = "Fetching param cache[" + iter + "/" + list.length+ "]: ";
text += Math.ceil(fetchedBytes / (1024 * 1024)).toString() + "MB
fetched. "
text += Math.floor(fetchedBytes * 100 / totalBytes).toString() + "%
completed, "
text += timeElapsed + " secs elapsed.";
text += " It can take a while when we first visit this page to
populate the cache."
text += " Later refreshes will become faster.";
- this.fetchProgressCallback[j]({
- fetchedBytes: fetchedBytes,
- totalBytes: totalBytes,
+ this.initProgressCallback[j]({
+ progress: fetchedBytes / totalBytes,
timeElapsed: timeElapsed,
text: text
});
}
};
- for (let j = 0; j < this.fetchProgressCallback.length; ++j) {
- this.fetchProgressCallback[j]({
- fetchedBytes: 0,
- totalBytes: totalBytes,
+ for (let j = 0; j < this.initProgressCallback.length; ++j) {
+ this.initProgressCallback[j]({
+ progress: fetchedBytes / totalBytes,
timeElapsed: 0,
text: "Start to fetch params",
});
@@ -1736,6 +1741,71 @@ export class Instance implements Disposable {
this.registerFunc("__async." + name, asyncVariant, override);
}
+ /**
+ * Asynchrously load webgpu pipelines when possible.
+ * @param mod The input module.
+ */
+ async asyncLoadWebGPUPiplines(mod: Module): Promise<void> {
+ 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)() as string;
+ let fmap: Record<string, FunctionInfo> = JSON.parse(fmap_str);
+ const totalFuncs = 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);
+ assert(key == finfo.name);
+ const event = webgpuContext.createShaderAsync(finfo, code).then((func:
Function) => {
+ this.beginScope();
+ fUpdatePrebuild(key, func);
+ this.endScope();
+
+ }).then(() => {
+ finishCounter += 1;
+ const tend = perf.now();
+ const timeReportGap = 1000;
+ // skip report if gap is smaller than 1000
+ if ((tend - tlastReport) < 1000 && finishCounter !=
fmapEntries.length) {
+ return;
+ }
+ tlastReport = tend;
+ const timeElapsed = Math.ceil((perf.now() - tstart) / 1000);
+ // report
+ 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(()=>{});
+ }
+ await allEvents;
+ assert(finishCounter == fmapEntries.length);
+ }
+
/**
* Initialize webgpu in the runtime.
* @param device The given GPU device.
@@ -1748,7 +1818,8 @@ export class Instance implements Disposable {
return webGPUContext.getDeviceAPI(name);
});
this.registerFunc("wasm.WebGPUCreateShader", (info: string, code: string)
=> {
- return webGPUContext.createShader(info, code);
+ const finfo = JSON.parse(info) as FunctionInfo;
+ return webGPUContext.createShader(finfo, code);
});
this.registerAsyncServerFunc("wasm.WebGPUWaitForTasks", async () => {
await webGPUContext.sync();
diff --git a/web/src/webgpu.ts b/web/src/webgpu.ts
index 8137646921..ac39595c76 100644
--- a/web/src/webgpu.ts
+++ b/web/src/webgpu.ts
@@ -58,12 +58,6 @@ export async function detectGPUDevice():
Promise<GPUDeviceDetectOutput | undefin
}
}
-interface FunctionInfo {
- name: string;
- arg_types: Array<string>;
- launch_param_tags: Array<string>;
-}
-
const canvasRenderWGSL =`
@group(0) @binding(0) var my_sampler : sampler;
@group(0) @binding(1) var my_texture : texture_2d<f32>;
@@ -272,6 +266,14 @@ class CanvaRenderManager implements Disposable {
}
}
+/**
+ * Function info from the API
+ */
+export interface FunctionInfo {
+ name: string;
+ arg_types: Array<string>;
+ launch_param_tags: Array<string>;
+}
/**
* WebGPU context
@@ -382,12 +384,41 @@ export class WebGPUContext {
/**
* Create a PackedFunc that runs the given shader
+ * via createComputePipeline
*
- * @param info The function information in json.
+ * @param info The function information already parsed as a record.
* @param code The shader data(in WGSL)
+ * @returns The shader
*/
- createShader(info: string, code: string): Function {
- const finfo = JSON.parse(info);
+ createShader(finfo: FunctionInfo, code: string) : Function {
+ return this.createShadeInternl(finfo, code, false) as Function;
+ }
+
+ /**
+ * Create a PackedFunc that runs the given shader asynchrously
+ * via createComputePipelineAsync
+ *
+ * @param info The function information already parsed as a record.
+ * @param code The shader data(in WGSL)
+ * @returns The shader
+ */
+ async createShaderAsync(finfo: FunctionInfo, code: string) :
Promise<Function> {
+ return await (this.createShadeInternl(finfo, code, true) as
Promise<Function>);
+ }
+
+ /**
+ * Internal impl of createShader for both async and sync mode.
+ *
+ * @param info The function information already parsed as a record.
+ * @param code The shader data(in WGSL)
+ * @param asyncMode Whether use async mode.
+ * @returns The shader function or promise of shader func.
+ */
+ private createShadeInternl(
+ finfo: FunctionInfo,
+ code: string,
+ asyncMode: boolean
+ ): Function | Promise<Function> {
const dispatchToDim: Array<number> = [];
let paramWriteAccess: Array<number> = [];
@@ -432,86 +463,103 @@ export class WebGPUContext {
bindGroupLayouts: [ bindGroupLayout ]
});
- const pipeline = this.device.createComputePipeline({
- layout: pipelineLayout,
- compute: {
- module: this.device.createShaderModule({
- code: code,
- hints: {
- main: {
- layout: pipelineLayout
- }
- }
- }),
- entryPoint: finfo.name
- }
- });
+ // Function to create the pipeline.
+ const createShaderFunc = (pipeline: GPUComputePipeline): Function => {
+ const submitShader = (...args: Array<GPUPointer | number>): void => {
+ if (this.debugShaderSubmitLimit != -1 &&
+ this.shaderSubmitCounter >= this.debugShaderSubmitLimit) {
+ this.shaderSubmitCounter += 1;
+ return;
+ }
- const submitShader = (...args: Array<GPUPointer | number>): void => {
- 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: Array<GPUBindGroupEntry> = [];
+ assert(args.length == layoutEntries.length + dispatchToDim.length);
+
+ for (let i = 0; i < layoutEntries.length; ++i) {
+ bindGroupEntries.push({
+ binding: i,
+ resource: {
+ buffer: this.gpuBufferFromPtr(args[i])
+ }
+ });
+ }
- const commandEncoder = this.device.createCommandEncoder();
- const compute = commandEncoder.beginComputePass();
- compute.setPipeline(pipeline);
- const bindGroupEntries: Array<GPUBindGroupEntry> = [];
- assert(args.length == layoutEntries.length + dispatchToDim.length);
+ compute.setBindGroup(0, this.device.createBindGroup({
+ layout: bindGroupLayout,
+ entries: bindGroupEntries
+ }));
+ const wl: Array<number> = [1, 1, 1, 1, 1, 1];
+ for (let i = 0; i < dispatchToDim.length; ++i) {
+ wl[dispatchToDim[i]] = args[layoutEntries.length + i];
+ }
- for (let i = 0; i < layoutEntries.length; ++i) {
- bindGroupEntries.push({
- binding: i,
- resource: {
- buffer: this.gpuBufferFromPtr(args[i])
+ // get around 65535 restriction of blockIdx.x
+ if (wl[2] != 1) {
+ throw Error("WebGPU: blockIdx.z is reserved for internal use");
+ }
+ // spread thinsg out into blockIdx.z
+ if (wl[0] >= (1 << 16)) {
+ let wl_x = wl[0];
+ let wl_z = wl[2];
+
+ while (wl_x >= (1 << 16)) {
+ if (wl_x % 2 != 0) {
+ throw Error("WebGPU: cannot factorize big gridDim.x=" +
wl[0].toString());
+ }
+ wl_x /= 2;
+ wl_z *= 2;
}
- });
- }
-
- compute.setBindGroup(0, this.device.createBindGroup({
- layout: bindGroupLayout,
- entries: bindGroupEntries
- }));
- const wl: Array<number> = [1, 1, 1, 1, 1, 1];
- for (let i = 0; i < dispatchToDim.length; ++i) {
- wl[dispatchToDim[i]] = args[layoutEntries.length + i];
- }
+ wl[0] = wl_x;
+ wl[2] = wl_z;
+ }
+ compute.dispatchWorkgroups(wl[0], wl[1], wl[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;
+ };
- // get around 65535 restriction of blockIdx.x
- if (wl[2] != 1) {
- throw Error("WebGPU: blockIdx.z is reserved for internal use");
- }
- // spread thinsg out into blockIdx.z
- if (wl[0] >= (1 << 16)) {
- let wl_x = wl[0];
- let wl_z = wl[2];
-
- while (wl_x >= (1 << 16)) {
- if (wl_x % 2 != 0) {
- throw Error("WebGPU: cannot factorize big gridDim.x=" +
wl[0].toString());
- }
- wl_x /= 2;
- wl_z *= 2;
+ const shaderModule = this.device.createShaderModule({
+ code: code,
+ hints: {
+ main: {
+ layout: pipelineLayout
}
- wl[0] = wl_x;
- wl[2] = wl_z;
}
- compute.dispatchWorkgroups(wl[0], wl[1], wl[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;
+ if (asyncMode) {
+ return this.device.createComputePipelineAsync({
+ layout: pipelineLayout,
+ compute: {
+ module: shaderModule,
+ entryPoint: finfo.name
+ }
+ }).then((pipeline: GPUComputePipeline) => {
+ return createShaderFunc(pipeline);
+ });
+ } else {
+ const pipeline = this.device.createComputePipeline({
+ layout: pipelineLayout,
+ compute: {
+ module: shaderModule,
+ entryPoint: finfo.name
+ }
+ });
+ return createShaderFunc(pipeline);
+ }
}
/**
