Constant STENCIL_EVAL_WGSL 

pub const STENCIL_EVAL_WGSL: &str = "// Uniform stencil evaluation compute kernel.\n//\n// Ported from opensubdiv-petite\'s `shaders/wgsl/stencil_eval.wgsl` (which itself\n// mirrors OpenSubdiv\'s `glslComputeKernel.glsl`), specialised to subdiv-kernels\'\n// CSR `StencilTable`: row `i`\'s entries span `offsets[i]..offsets[i+1]`, so the\n// per-row size is derived from the offsets and no separate `sizes` buffer is\n// needed. Positions/primvars only -- the limit-derivative (b1) path is separate.\n\noverride WORKGROUP_SIZE: u32 = 64u;\n\nstruct Params {\n    // All offsets/strides are in floats.\n    src_offset: u32,\n    dst_offset: u32,\n    src_stride: u32,\n    dst_stride: u32,\n    length: u32,       // components per element (e.g. 3 for xyz)\n    batch_start: u32,  // first output row (inclusive)\n    batch_end: u32,    // last output row (exclusive)\n}\n\n@group(0) @binding(0) var<uniform> params: Params;\n@group(0) @binding(1) var<storage, read> src_buffer: array<f32>;\n@group(0) @binding(2) var<storage, read_write> dst_buffer: array<f32>;\n// CSR row offsets, length = output_count + 1.\n@group(0) @binding(3) var<storage, read> stencil_offsets: array<u32>;\n@group(0) @binding(4) var<storage, read> stencil_indices: array<u32>;\n@group(0) @binding(5) var<storage, read> stencil_weights: array<f32>;\n// Indirection: output-row indices for the sparse (indexed) entry point.\n@group(0) @binding(6) var<storage, read> stencil_indirection: array<u32>;\n\n// Cap per-element component count (matches the host MAX_COMPONENTS).\nconst MAX_LENGTH: u32 = 32u;\n\n// Evaluate output row `current` into dst_buffer.\nfn eval_row(current: u32) {\n    let row_start = stencil_offsets[current];\n    let row_end = stencil_offsets[current + 1u];\n    let dst_base = params.dst_offset + current * params.dst_stride;\n\n    for (var c: u32 = 0u; c < params.length && c < MAX_LENGTH; c = c + 1u) {\n        var sum: f32 = 0.0;\n        for (var si: u32 = row_start; si < row_end; si = si + 1u) {\n            let vi = params.src_offset + stencil_indices[si] * params.src_stride + c;\n            sum = sum + stencil_weights[si] * src_buffer[vi];\n        }\n        dst_buffer[dst_base + c] = sum;\n    }\n}\n\n// Dense: invocation gid.x handles output row gid.x + batch_start, in\n// [batch_start, batch_end).\n@compute @workgroup_size(WORKGROUP_SIZE)\nfn eval_stencils(@builtin(global_invocation_id) gid: vec3<u32>) {\n    let current = gid.x + params.batch_start;\n    if (current >= params.batch_end) {\n        return;\n    }\n    eval_row(current);\n}\n\n// Sparse: invocation gid.x handles the output row named by the indirection\n// buffer; `batch_end` is the number of indirection entries.\n@compute @workgroup_size(WORKGROUP_SIZE)\nfn eval_stencils_indexed(@builtin(global_invocation_id) gid: vec3<u32>) {\n    if (gid.x >= params.batch_end) {\n        return;\n    }\n    eval_row(stencil_indirection[gid.x]);\n}\n";

Available on crate feature wgpu only.

Canonical WGSL source for the stencil-eval compute kernel.