MCPcopy Index your code
hub / github.com/apache/tvm / submitShader

Method submitShader

web/src/webgpu.ts:716–822  ·  view source on GitHub ↗
(...args: Array<GPUPointer | number>)

Source from the content-addressed store, hash-verified

714 // Function to create the pipeline.
715 const createShaderFunc = (pipeline: GPUComputePipeline): Function => {
716 const submitShader = (...args: Array<GPUPointer | number>): void => {
717 if (this.debugShaderSubmitLimit != -1 &&
718 this.shaderSubmitCounter >= this.debugShaderSubmitLimit) {
719 this.shaderSubmitCounter += 1;
720 return;
721 }
722
723 // Reuse a single command encoder across dispatches; only flush on sync/readback.
724 if (!this.pendingEncoder) {
725 this.pendingEncoder = this.device.createCommandEncoder();
726 }
727
728 const compute = this.pendingEncoder.beginComputePass();
729 compute.setPipeline(pipeline);
730 const bindGroupEntries: Array<GPUBindGroupEntry> = [];
731 const numBufferOrPodArgs = bufferArgIndices.length + podArgIndices.length;
732
733 assert(args.length == numBufferOrPodArgs + dispatchToDim.length);
734
735 const workDim: Array<number> = [1, 1, 1, 1, 1, 1];
736 for (let i = 0; i < dispatchToDim.length; ++i) {
737 workDim[dispatchToDim[i]] = args[numBufferOrPodArgs + i];
738 }
739
740 // get around 65535 restriction of blockIdx.x
741 if (workDim[2] != 1) {
742 throw Error("WebGPU: blockIdx.z is reserved for internal use");
743 }
744 const packDimX = workDim[0];
745 // spread thinsg out into blockIdx.z
746 if (workDim[0] >= (1 << 16)) {
747 let wl_x = workDim[0];
748 let wl_z = workDim[2];
749
750 while (wl_x >= (1 << 16)) {
751 if (wl_x % 2 == 0) {
752 wl_x = wl_x / 2;
753 } else {
754 // pad up
755 wl_x = (wl_x + 1) / 2;
756 }
757 wl_z *= 2;
758 }
759 workDim[0] = wl_x;
760 workDim[2] = wl_z;
761 assert(wl_x * wl_z >= packDimX);
762 }
763
764 for (let i = 0; i < bufferArgIndices.length; ++i) {
765 bindGroupEntries.push({
766 binding: i,
767 resource: {
768 buffer: this.gpuBufferFromPtr(args[bufferArgIndices[i]])
769 }
770 });
771 }
772
773 const sizeOfI32 = 4;

Callers

nothing calls this directly

Calls 6

gpuBufferFromPtrMethod · 0.95
getUniformFromPoolMethod · 0.95
flushCommandsMethod · 0.95
assertFunction · 0.90
endMethod · 0.80
logMethod · 0.45

Tested by

no test coverage detected