* 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
)
| 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 |
no test coverage detected