Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[JS/WebGPU] fix an error in Clip #18799

Merged
merged 5 commits into from
Dec 19, 2023
Merged

[JS/WebGPU] fix an error in Clip #18799

Show file tree
Hide file tree
Changes from 4 commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
c00d9dc
Make sure that min/max input data type is float before reading float …
satyajandhyala Dec 13, 2023
acf7e9c
Defined both global_idx and local_idx in common.ts
satyajandhyala Dec 13, 2023
78f3fd1
typo
satyajandhyala Dec 13, 2023
5df50d9
Code review changes.
satyajandhyala Dec 13, 2023
238aa47
Replaced global_id.x with global_idx in gemm.ts.
satyajandhyala Dec 15, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 3 additions & 3 deletions js/web/lib/wasm/jsep/webgpu/ops/common.ts
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -772,14 +772,14 @@ class ShaderHelperImpl implements ShaderHelper {
const is1DimensionDispatch = this.normalizedDispatchGroup[1] === 1 && this.normalizedDispatchGroup[2] === 1;
const paramList = is1DimensionDispatch ? `@builtin(global_invocation_id) global_id : vec3<u32>,
@builtin(local_invocation_id) local_id : vec3<u32>` :
`@builtin(local_invocation_index) local_index : u32,
`@builtin(local_invocation_index) local_idx : u32,
@builtin(workgroup_id) workgroup_id : vec3<u32>,
@builtin(num_workgroups) num_workgroups : vec3<u32>`;
satyajandhyala marked this conversation as resolved.
Show resolved Hide resolved
const globalIdxDefinition = is1DimensionDispatch ?
'let global_idx = global_id.x;' :
'let global_idx = global_id.x; let local_idx = local_id.x;' :
satyajandhyala marked this conversation as resolved.
Show resolved Hide resolved
`let global_idx = (workgroup_id.z * num_workgroups[0] * num_workgroups[1] +
workgroup_id.y * num_workgroups[0] + workgroup_id.x) * ${
workgroupSizeX * workgroupSizeY * workgroupSizeZ}u + local_index;`;
workgroupSizeX * workgroupSizeY * workgroupSizeZ}u + local_idx;`;

return `@compute @workgroup_size(${workgroupSizeX}, ${workgroupSizeY}, ${workgroupSizeZ})
fn main(${paramList}) {
Expand Down
1 change: 0 additions & 1 deletion js/web/lib/wasm/jsep/webgpu/ops/reduce-shared.ts
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -141,7 +141,6 @@ export const createReduceSharedProgramInfo =
return ((a - 1u) / b + 1u);
}
${shaderHelper.mainStart(workgroupSize)}
let local_idx = local_id.x;

let outputIndex = global_idx / ${workgroupSize};
let offset = outputIndex * uniforms.reduceSize;
Expand Down
4 changes: 2 additions & 2 deletions js/web/lib/wasm/jsep/webgpu/ops/softmax.ts
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -73,8 +73,8 @@ const createSoftmaxProgramInfo = (input: TensorView, attributes: SoftmaxAttribut
}
${shaderHelper.registerUniform('packedCols', 'i32').declareVariables(x, output)}
${shaderHelper.mainStart()}
let gindex = i32(global_id.x);
satyajandhyala marked this conversation as resolved.
Show resolved Hide resolved
let lindex = i32(local_id.x);
let gindex = i32(global_idx);
let lindex = i32(local_idx);
const wg = ${WG};
let row = gindex / wg;
let cols = uniforms.packedCols;
Expand Down
4 changes: 2 additions & 2 deletions js/web/lib/wasm/jsep/webgpu/ops/unary-op.ts
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -125,8 +125,8 @@ export interface ClipAttributes extends AttributeWithCacheKey {
}

const generateClipAttributesFromInputs = (inputs: readonly TensorView[]): ClipAttributes => {
const min = (inputs.length >= 2) ? inputs[1].getFloat32Array()[0] : MIN_CLIP;
const max = (inputs.length >= 3) ? inputs[2].getFloat32Array()[0] : MAX_CLIP;
const min = (inputs.length >= 2 && inputs[1].data !== 0) ? inputs[1].getFloat32Array()[0] : MIN_CLIP;
const max = (inputs.length >= 3 && inputs[2].data !== 0) ? inputs[2].getFloat32Array()[0] : MAX_CLIP;
return createAttributeWithCacheKey({min, max});
};

Expand Down
Loading