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

Method createShadeInternal

web/src/webgpu.ts:649–856  ·  view source on GitHub ↗

* Internal impl of createShader for both async and sync mode. * * @param finfo 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.

(
    finfo: FunctionInfo,
    code: string,
    asyncMode: boolean
  )

Source from the content-addressed store, hash-verified

647 * @returns The shader function or promise of shader func.
648 */
649 private createShadeInternal(
650 finfo: FunctionInfo,
651 code: string,
652 asyncMode: boolean
653 ): Function | Promise<Function> {
654 const dispatchToDim: Array<number> = [];
655 let paramWriteAccess: Array<number> = [];
656
657 for (let i = 0; i < finfo.launch_param_tags.length; ++i) {
658 const tag: string = finfo.launch_param_tags[i];
659 if (tag.startsWith("blockIdx.")) {
660 const target: number = tag.charCodeAt(tag.length - 1) - ("x".charCodeAt(0));
661 assert(target >= 0 && target < 3);
662 dispatchToDim.push(target);
663 } else if (tag.startsWith("threadIdx.")) {
664 const target: number = tag.charCodeAt(tag.length - 1) - ("x".charCodeAt(0));
665 assert(target >= 0 && target < 3);
666 dispatchToDim.push(target + 3);
667 } else if (tag.startsWith("paramWriteAccess:")) {
668 paramWriteAccess = JSON.parse(tag.substring(17));
669 } else {
670 throw new Error("Cannot handle thread_axis " + tag);
671 }
672 }
673
674
675 const layoutEntries: Array<GPUBindGroupLayoutEntry> = [];
676 const bufferArgIndices: Array<number> = [];
677 const podArgIndices: Array<number> = [];
678
679 for (let i = 0; i < finfo.arg_types.length; ++i) {
680 const dtype = finfo.arg_types[i];
681 if (dtype == "handle") {
682 layoutEntries.push({
683 binding: bufferArgIndices.length,
684 visibility: GPUShaderStage.COMPUTE,
685 buffer: {
686 type: paramWriteAccess[bufferArgIndices.length] ? "storage" : "read-only-storage"
687 }
688 });
689 bufferArgIndices.push(i);
690 } else if (dtype.startsWith("int") || dtype.startsWith("uint") || dtype.startsWith("float")) {
691 podArgIndices.push(i);
692 } else {
693 throw new Error("Cannot handle argument type " + dtype + " in WebGPU shader");
694 }
695 }
696
697 assert(paramWriteAccess.length == bufferArgIndices.length);
698 // POD arguments are pass in the end
699 layoutEntries.push({
700 binding: bufferArgIndices.length,
701 visibility: GPUShaderStage.COMPUTE,
702 buffer: {
703 type: "uniform"
704 }
705 });
706

Callers 2

createShaderMethod · 0.95
createShaderAsyncMethod · 0.95

Calls 2

assertFunction · 0.90
parseMethod · 0.80

Tested by

no test coverage detected