mirror of
https://github.com/saymrwulf/onnxruntime.git
synced 2026-07-03 03:58:54 +00:00
[JS/WebGPU] fix an error in Clip (#18799)
### Description <!-- Describe your changes. --> Check whether the min/max inputs are provided and use default values if not provided. ### Motivation and Context <!-- - Why is this change required? What problem does it solve? - If it fixes an open issue, please link to the issue here. -->
This commit is contained in:
parent
32fcf73740
commit
98510fb8fb
5 changed files with 10 additions and 11 deletions
|
|
@ -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>`;
|
||||
const globalIdxDefinition = is1DimensionDispatch ?
|
||||
'let global_idx = global_id.x;' :
|
||||
'let global_idx = global_id.x; let local_idx = local_id.x;' :
|
||||
`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}) {
|
||||
|
|
|
|||
|
|
@ -97,8 +97,8 @@ const createGemmProgramInfo = (inputs: readonly TensorView[], attributes: GemmAt
|
|||
${shaderHelper.mainStart()}
|
||||
${shaderHelper.guardAgainstOutOfBoundsWorkgroupSizes(outputSize)}
|
||||
|
||||
let m = global_id.x / N;
|
||||
let n = global_id.x % N;
|
||||
let m = global_idx / N;
|
||||
let n = global_idx % N;
|
||||
|
||||
var value = ${dataType}(0);
|
||||
for (var k: u32 = 0u; k<${K}u; k++) {
|
||||
|
|
@ -107,7 +107,7 @@ const createGemmProgramInfo = (inputs: readonly TensorView[], attributes: GemmAt
|
|||
|
||||
${calculateAlpha}
|
||||
${calculateC}
|
||||
output[global_id.x] = value;
|
||||
output[global_idx] = value;
|
||||
|
||||
}`;
|
||||
return {
|
||||
|
|
|
|||
|
|
@ -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;
|
||||
|
|
|
|||
|
|
@ -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);
|
||||
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;
|
||||
|
|
|
|||
|
|
@ -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});
|
||||
};
|
||||
|
||||
|
|
|
|||
Loading…
Reference in a new issue