Skip to content

Commit

Permalink
Fix WebGPU warnings
Browse files Browse the repository at this point in the history
  • Loading branch information
huningxin committed Mar 23, 2022
1 parent 1b6c37f commit 1ed8076
Showing 1 changed file with 36 additions and 35 deletions.
Original file line number Diff line number Diff line change
Expand Up @@ -9,16 +9,16 @@
'use strict';

const preprocessWGSL = `
[[block]] struct Tensor {
struct Tensor {
values: array<f32>;
};
[[group(0), binding(0)]] var samp : sampler;
[[group(0), binding(1)]] var<storage, write> inputTensor : Tensor;
[[group(0), binding(2)]] var inputTex : texture_2d<f32>;
@group(0) @binding(0) var samp : sampler;
@group(0) @binding(1) var<storage, write> inputTensor : Tensor;
@group(0) @binding(2) var inputTex : texture_2d<f32>;
[[stage(compute), workgroup_size(8, 8)]]
fn main([[builtin(global_invocation_id)]] globalID : vec3<u32>) {
@stage(compute) @workgroup_size(8, 8)
fn main(@builtin(global_invocation_id) globalID : vec3<u32>) {
let dims : vec2<i32> = textureDimensions(inputTex, 0);
var inputValue = textureSampleLevel(inputTex, samp, vec2<f32>(globalID.xy) / vec2<f32>(dims), 0.0).rgb;
var normalizedInput = (inputValue * vec3<f32>(255.0, 255.0, 255.0) - vec3<f32>(127.5, 127.5, 127.5)) / vec3<f32>(127.5, 127.5, 127.5);
Expand All @@ -31,18 +31,18 @@ fn main([[builtin(global_invocation_id)]] globalID : vec3<u32>) {
`;

const segmentationWGSL = `
[[block]] struct SegMap {
struct SegMap {
labels: array<i32>;
};
[[group(0), binding(0)]] var samp : sampler;
[[group(0), binding(1)]] var<storage, read> segmap : SegMap;
[[group(0), binding(2)]] var inputTex : texture_2d<f32>;
[[group(0), binding(3)]] var blurredInputTex : texture_2d<f32>;
[[group(0), binding(4)]] var outputTex : texture_storage_2d<rgba8unorm, write>;
@group(0) @binding(0) var samp : sampler;
@group(0) @binding(1) var<storage, read> segmap : SegMap;
@group(0) @binding(2) var inputTex : texture_2d<f32>;
@group(0) @binding(3) var blurredInputTex : texture_2d<f32>;
@group(0) @binding(4) var outputTex : texture_storage_2d<rgba8unorm, write>;
[[stage(compute), workgroup_size(8, 8)]]
fn main([[builtin(global_invocation_id)]] globalID : vec3<u32>) {
@stage(compute) @workgroup_size(8, 8)
fn main(@builtin(global_invocation_id) globalID : vec3<u32>) {
let dims : vec2<i32> = textureDimensions(inputTex, 0);
var input = textureSampleLevel(inputTex, samp, vec2<f32>(globalID.xy) / vec2<f32>(dims), 0.0).rgb;
Expand All @@ -60,20 +60,20 @@ fn main([[builtin(global_invocation_id)]] globalID : vec3<u32>) {
`;

const blurWGSL = `
[[block]] struct Params {
struct Params {
filterDim : u32;
blockDim : u32;
};
[[group(0), binding(0)]] var samp : sampler;
[[group(0), binding(1)]] var<uniform> params : Params;
[[group(1), binding(1)]] var inputTex : texture_2d<f32>;
[[group(1), binding(2)]] var outputTex : texture_storage_2d<rgba8unorm, write>;
@group(0) @binding(0) var samp : sampler;
@group(0) @binding(1) var<uniform> params : Params;
@group(1) @binding(1) var inputTex : texture_2d<f32>;
@group(1) @binding(2) var outputTex : texture_storage_2d<rgba8unorm, write>;
[[block]] struct Flip {
struct Flip {
value : u32;
};
[[group(1), binding(3)]] var<uniform> flip : Flip;
@group(1) @binding(3) var<uniform> flip : Flip;
// This shader blurs the input texture in one direction, depending on whether
// |flip.value| is 0 or 1.
Expand All @@ -91,10 +91,10 @@ const blurWGSL = `
var<workgroup> tile : array<array<vec3<f32>, 128>, 4>;
[[stage(compute), workgroup_size(32, 1, 1)]]
@stage(compute) @workgroup_size(32, 1, 1)
fn main(
[[builtin(workgroup_id)]] WorkGroupID : vec3<u32>,
[[builtin(local_invocation_id)]] LocalInvocationID : vec3<u32>
@builtin(workgroup_id) WorkGroupID : vec3<u32>,
@builtin(local_invocation_id) LocalInvocationID : vec3<u32>
) {
let filterOffset : u32 = (params.filterDim - 1u) / 2u;
let dims : vec2<i32> = textureDimensions(inputTex, 0);
Expand Down Expand Up @@ -143,16 +143,16 @@ fn main(
`;

const fullscreenTexturedQuadWGSL = `
[[group(0), binding(0)]] var mySampler : sampler;
[[group(0), binding(1)]] var myTexture : texture_2d<f32>;
@group(0) @binding(0) var mySampler : sampler;
@group(0) @binding(1) var myTexture : texture_2d<f32>;
struct VertexOutput {
[[builtin(position)]] Position : vec4<f32>;
[[location(0)]] fragUV : vec2<f32>;
@builtin(position) Position : vec4<f32>;
@location(0) fragUV : vec2<f32>;
};
[[stage(vertex)]]
fn vert_main([[builtin(vertex_index)]] VertexIndex : u32) -> VertexOutput {
@stage(vertex)
fn vert_main(@builtin(vertex_index) VertexIndex : u32) -> VertexOutput {
var pos = array<vec2<f32>, 6>(
vec2<f32>( 1.0, 1.0),
vec2<f32>( 1.0, -1.0),
Expand All @@ -175,8 +175,8 @@ fn vert_main([[builtin(vertex_index)]] VertexIndex : u32) -> VertexOutput {
return output;
}
[[stage(fragment)]]
fn frag_main([[location(0)]] fragUV : vec2<f32>) -> [[location(0)]] vec4<f32> {
@stage(fragment)
fn frag_main(@location(0) fragUV : vec2<f32>) -> @location(0) vec4<f32> {
return textureSample(myTexture, mySampler, fragUV);
}
`;
Expand Down Expand Up @@ -641,13 +641,14 @@ const batch = [4, 4];
);
}

computePass.endPass();
computePass.end();

const passEncoder = commandEncoder.beginRenderPass({
colorAttachments: [
{
view: this.context_.getCurrentTexture().createView(),
loadValue: { r: 0.0, g: 0.0, b: 0.0, a: 1.0 },
clearValue: { r: 0.0, g: 0.0, b: 0.0, a: 1.0 },
loadOp: 'clear',
storeOp: 'store',
},
],
Expand All @@ -670,7 +671,7 @@ const batch = [4, 4];
passEncoder.setPipeline(this.fullscreenQuadPipeline_);
passEncoder.setBindGroup(0, showResultBindGroup);
passEncoder.draw(6, 1, 0, 0);
passEncoder.endPass();
passEncoder.end();
device.queue.submit([commandEncoder.finish()]);

await device.queue.onSubmittedWorkDone();
Expand Down

0 comments on commit 1ed8076

Please sign in to comment.