dawn-cmake/test/bug/tint/942.wgsl.expected.wgsl
Ben Clayton 01e4b6fc18 wgsl: Replace [[decoration]] with @decoration
Deprecate the old syntax. Migrate everything to the new syntax.

Bug: tint:1382
Change-Id: Ide12b2e927b17dc93b9714c7049090864cc568d3
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/77260
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
Commit-Queue: David Neto <dneto@google.com>
2022-01-19 22:46:57 +00:00

55 lines
2.0 KiB
WebGPU Shading Language

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>;
struct Flip {
value : u32;
}
@group(1) @binding(3) var<uniform> flip : Flip;
var<workgroup> tile : array<array<vec3<f32>, 256>, 4>;
@stage(compute) @workgroup_size(64, 1, 1)
fn main(@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);
let baseIndex = (vec2<i32>(((WorkGroupID.xy * vec2<u32>(params.blockDim, 4u)) + (LocalInvocationID.xy * vec2<u32>(4u, 1u)))) - vec2<i32>(i32(filterOffset), 0));
for(var r : u32 = 0u; (r < 4u); r = (r + 1u)) {
for(var c : u32 = 0u; (c < 4u); c = (c + 1u)) {
var loadIndex = (baseIndex + vec2<i32>(i32(c), i32(r)));
if ((flip.value != 0u)) {
loadIndex = loadIndex.yx;
}
tile[r][((4u * LocalInvocationID.x) + c)] = textureSampleLevel(inputTex, samp, ((vec2<f32>(loadIndex) + vec2<f32>(0.25, 0.25)) / vec2<f32>(dims)), 0.0).rgb;
}
}
workgroupBarrier();
for(var r : u32 = 0u; (r < 4u); r = (r + 1u)) {
for(var c : u32 = 0u; (c < 4u); c = (c + 1u)) {
var writeIndex = (baseIndex + vec2<i32>(i32(c), i32(r)));
if ((flip.value != 0u)) {
writeIndex = writeIndex.yx;
}
let center : u32 = ((4u * LocalInvocationID.x) + c);
if ((((center >= filterOffset) && (center < (256u - filterOffset))) && all((writeIndex < dims)))) {
var acc : vec3<f32> = vec3<f32>(0.0, 0.0, 0.0);
for(var f : u32 = 0u; (f < params.filterDim); f = (f + 1u)) {
var i : u32 = ((center + f) - filterOffset);
acc = (acc + ((1.0 / f32(params.filterDim)) * tile[r][i]));
}
textureStore(outputTex, writeIndex, vec4<f32>(acc, 1.0));
}
}
}
}