nikola/llama_cpp
1
0
Fork 0
Code Issues Pull Requests Actions 21 Packages Projects Releases Wiki Activity

revert to using global_invocation_id for cpy shader (#23955)

Browse Source
This commit is contained in:
Masashi Yoshimura
2026-06-01 16:59:06 -07:00
committed by GitHub
parent 5dcb711666
commit b8275a8acc
1 changed files with 4 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-webgpu/wgsl-shaders/cpy.wgsl
+4 -5
View File
@@ -50,13 +50,13 @@ var<uniform> params: Params;
@compute @workgroup_size(WG_SIZE) @compute @workgroup_size(WG_SIZE)
fn main( fn main(
@builtin(global_invocation_index) gindex: u32, @builtin(global_invocation_id) gid: vec3<u32>,
) { ) {
if (gindex >= params.ne) { if (gid.x >= params.ne) {
return; return;
} }
var i = gindex; var i = gid.x;
let i3 = i / (params.src_ne2 * params.src_ne1 * params.src_ne0); let i3 = i / (params.src_ne2 * params.src_ne1 * params.src_ne0);
i = i % (params.src_ne2 * params.src_ne1 * params.src_ne0); i = i % (params.src_ne2 * params.src_ne1 * params.src_ne0);
let i2 = i / (params.src_ne1 * params.src_ne0); let i2 = i / (params.src_ne1 * params.src_ne0);
@@ -64,7 +64,7 @@ fn main(
let i1 = i / params.src_ne0; let i1 = i / params.src_ne0;
let i0 = i % params.src_ne0; let i0 = i % params.src_ne0;
var j = gindex; var j = gid.x;
let j3 = j / (params.dst_ne2 * params.dst_ne1 * params.dst_ne0); let j3 = j / (params.dst_ne2 * params.dst_ne1 * params.dst_ne0);
j = j % (params.dst_ne2 * params.dst_ne1 * params.dst_ne0); j = j % (params.dst_ne2 * params.dst_ne1 * params.dst_ne0);
let j2 = j / (params.dst_ne1 * params.dst_ne0); let j2 = j / (params.dst_ne1 * params.dst_ne0);
@@ -80,4 +80,3 @@ fn main(
dst[params.offset_dst + dst_idx] = DST_TYPE((src[params.offset_src + src_idx])); dst[params.offset_dst + dst_idx] = DST_TYPE((src[params.offset_src + src_idx]));
} }