export declare const cgKernelsWGSL = "\n/**\n * Conjugate Gradient Kernels \u2014 Sparse Linear Algebra on WebGPU\n */\n\nstruct SolverArgs {\n  num_rows: u32,\n  vector_width: u32,\n  n: u32,\n  alpha: f32,\n};\n\n@group(0) @binding(0) var<storage, read> csr_val: array<f32>;\n@group(0) @binding(1) var<storage, read> csr_col: array<u32>;\n@group(0) @binding(2) var<storage, read> csr_row: array<u32>;\n\n@group(1) @binding(0) var<storage, read> vec_in: array<f32>;\n@group(1) @binding(1) var<storage, read_write> vec_out: array<f32>;\n\n@group(2) @binding(0) var<uniform> args: SolverArgs;\n@group(2) @binding(1) var<storage, read> dot_vec_b: array<f32>;\n\n@group(3) @binding(0) var<storage, read_write> partial_sums: array<f32>;\n@group(3) @binding(1) var<storage, read_write> scalar_result: array<f32>;\n\n// Scalar workspace for on-device alpha/beta (avoids per-iteration CPU readback).\n// scal_num / scal_den are inputs, scal_out is the result. Referenced only by the\n// _buf kernels and divide_scalar; with layout:'auto' each pipeline binds just what\n// it uses.\n@group(2) @binding(2) var<storage, read> scal_num: array<f32>;\n@group(2) @binding(3) var<storage, read> scal_den: array<f32>;\n@group(2) @binding(4) var<storage, read_write> scal_out: array<f32>;\n\nvar<workgroup> spmv_shared: array<f32, 256>;\n\n@compute @workgroup_size(256)\nfn spmv_vector(\n    @builtin(global_invocation_id) global_id: vec3<u32>,\n    @builtin(local_invocation_id) local_id: vec3<u32>\n) {\n    let tid = local_id[0];\n    let gid = global_id[0];\n    let threads_per_row = args.vector_width;\n    let row = gid / threads_per_row;\n    let lane = gid % threads_per_row;\n\n    var sum: f32 = 0.0;\n    if (row < args.num_rows) {\n        let row_start = csr_row[row];\n        let row_end = csr_row[row + 1];\n\n        for (var i = row_start + lane; i < row_end; i = i + threads_per_row) {\n            sum += csr_val[i] * vec_in[csr_col[i]];\n        }\n    }\n\n    spmv_shared[tid] = sum;\n    workgroupBarrier();\n\n    for (var s = threads_per_row / 2u; s > 0u; s >>= 1u) {\n        if (lane < s) {\n            spmv_shared[tid] += spmv_shared[tid + s];\n        }\n        workgroupBarrier();\n    }\n\n    if (row < args.num_rows && lane == 0u) {\n        vec_out[row] = spmv_shared[tid];\n    }\n}\n\n@compute @workgroup_size(64)\nfn spmv(@builtin(global_invocation_id) global_id: vec3<u32>) {\n    let row = global_id[0];\n    if (row >= args.num_rows) {\n        return;\n    }\n\n    let row_start = csr_row[row];\n    let row_end = csr_row[row + 1];\n\n    var sum: f32 = 0.0;\n    for (var i = row_start; i < row_end; i = i + 1u) {\n        sum += csr_val[i] * vec_in[csr_col[i]];\n    }\n\n    vec_out[row] = sum;\n}\n\n@compute @workgroup_size(256)\nfn saxpy(@builtin(global_invocation_id) global_id: vec3<u32>) {\n    let idx = global_id[0];\n    if (idx >= args.n) {\n        return;\n    }\n    vec_out[idx] = args.alpha * vec_in[idx] + vec_out[idx];\n}\n\n@compute @workgroup_size(256)\nfn p_update(@builtin(global_invocation_id) global_id: vec3<u32>) {\n    let idx = global_id[0];\n    if (idx >= args.n) {\n        return;\n    }\n    vec_out[idx] = vec_in[idx] + args.alpha * vec_out[idx];\n}\n\n@compute @workgroup_size(256)\nfn vec_copy(@builtin(global_invocation_id) global_id: vec3<u32>) {\n    let idx = global_id[0];\n    if (idx >= args.n) {\n        return;\n    }\n    vec_out[idx] = vec_in[idx];\n}\n\n@compute @workgroup_size(256)\nfn vec_zero(@builtin(global_invocation_id) global_id: vec3<u32>) {\n    let idx = global_id[0];\n    if (idx >= args.n) {\n        return;\n    }\n    vec_out[idx] = 0.0;\n}\n\nvar<workgroup> dot_shared: array<f32, 256>;\n\n@compute @workgroup_size(256)\nfn dot_product(\n    @builtin(global_invocation_id) global_id: vec3<u32>,\n    @builtin(local_invocation_id) local_id: vec3<u32>,\n    @builtin(workgroup_id) workgroup_id: vec3<u32>\n) {\n    let idx = global_id[0];\n    let tid = local_id[0];\n\n    if (idx < args.n) {\n        dot_shared[tid] = vec_in[idx] * dot_vec_b[idx];\n    } else {\n        dot_shared[tid] = 0.0;\n    }\n\n    workgroupBarrier();\n\n    for (var s = 128u; s > 0u; s >>= 1u) {\n        if (tid < s) {\n            dot_shared[tid] += dot_shared[tid + s];\n        }\n        workgroupBarrier();\n    }\n\n    if (tid == 0u) {\n        partial_sums[workgroup_id[0]] = dot_shared[0];\n    }\n}\n\nvar<workgroup> reduce_shared: array<f32, 256>;\n\n@compute @workgroup_size(256)\nfn final_reduce(@builtin(local_invocation_id) local_id: vec3<u32>) {\n    let tid = local_id[0];\n    let count = args.n;\n\n    var acc: f32 = 0.0;\n    var i = tid;\n    loop {\n        if (i >= count) {\n            break;\n        }\n        acc += partial_sums[i];\n        i += 256u;\n    }\n    reduce_shared[tid] = acc;\n\n    workgroupBarrier();\n\n    for (var s = 128u; s > 0u; s >>= 1u) {\n        if (tid < s) {\n            reduce_shared[tid] += reduce_shared[tid + s];\n        }\n        workgroupBarrier();\n    }\n\n    if (tid == 0u) {\n        scalar_result[0] = reduce_shared[0];\n    }\n}\n\n// \u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\n// 8. Jacobi preconditioner \u2014 extract inverse diagonal of A (M^-1 = 1/diag)\n//    vec_out[row] = 1/A[row,row] (or 1.0 if diagonal is ~0)\n//    Uses: groups 0 (CSR), 1 (binding 1 = out), 2 (binding 0 = args)\n// \u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\n\n@compute @workgroup_size(256)\nfn extract_inv_diagonal(@builtin(global_invocation_id) global_id: vec3<u32>) {\n    let row = global_id[0];\n    if (row >= args.num_rows) {\n        return;\n    }\n    let row_start = csr_row[row];\n    let row_end = csr_row[row + 1];\n    var diag: f32 = 0.0;\n    for (var i = row_start; i < row_end; i = i + 1u) {\n        if (csr_col[i] == row) {\n            diag = csr_val[i];\n        }\n    }\n    if (abs(diag) > 1e-30) {\n        vec_out[row] = 1.0 / diag;\n    } else {\n        vec_out[row] = 1.0;\n    }\n}\n\n// \u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\n// 9. Apply preconditioner \u2014 z = M^-1 \u2218 r (elementwise)\n//    vec_in = r, dot_vec_b = invDiag, vec_out = z\n//    Uses: groups 1 (vecs), 2 (bindings 0=args, 1=invDiag)\n// \u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\n\n@compute @workgroup_size(256)\nfn apply_precond(@builtin(global_invocation_id) global_id: vec3<u32>) {\n    let idx = global_id[0];\n    if (idx >= args.n) {\n        return;\n    }\n    vec_out[idx] = vec_in[idx] * dot_vec_b[idx];\n}\n\n// \u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\n// 10. Scalar divide \u2014 scal_out[0] = scal_num[0] / (scal_den[0] + eps)\n//     Single-thread. On-device alpha/beta to avoid CPU readback.\n//     Uses: group 2 (bindings 2=num, 3=den, 4=out)\n// \u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\n\n@compute @workgroup_size(1)\nfn divide_scalar() {\n    let den = scal_den[0];\n    var inv: f32 = 0.0;\n    if (abs(den) > 1e-30) {\n        inv = scal_num[0] / den;\n    }\n    scal_out[0] = inv;\n}\n\n// \u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\n// 11. SAXPY (scalar from buffer): vec_out = s * vec_in + vec_out\n//     s = scal_num[0]. Uses: groups 1 (vecs), 2 (bindings 0=args, 2=s)\n// \u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\n\n@compute @workgroup_size(256)\nfn saxpy_buf(@builtin(global_invocation_id) global_id: vec3<u32>) {\n    let idx = global_id[0];\n    if (idx >= args.n) {\n        return;\n    }\n    vec_out[idx] = scal_num[0] * vec_in[idx] + vec_out[idx];\n}\n\n// \u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\n// 12. SAXPY-negate (scalar from buffer): vec_out = -s * vec_in + vec_out\n//     s = scal_num[0]. Uses: groups 1 (vecs), 2 (bindings 0=args, 2=s)\n// \u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\n\n@compute @workgroup_size(256)\nfn saxpy_neg_buf(@builtin(global_invocation_id) global_id: vec3<u32>) {\n    let idx = global_id[0];\n    if (idx >= args.n) {\n        return;\n    }\n    vec_out[idx] = -scal_num[0] * vec_in[idx] + vec_out[idx];\n}\n\n// \u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\n// 13. p-update (scalar from buffer): vec_out = vec_in + s * vec_out\n//     s = scal_num[0]. Uses: groups 1 (vecs), 2 (bindings 0=args, 2=s)\n// \u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\u2550\n\n@compute @workgroup_size(256)\nfn p_update_buf(@builtin(global_invocation_id) global_id: vec3<u32>) {\n    let idx = global_id[0];\n    if (idx >= args.n) {\n        return;\n    }\n    vec_out[idx] = vec_in[idx] + scal_num[0] * vec_out[idx];\n}\n";
//# sourceMappingURL=cg_kernels.d.ts.map