mirror of
https://github.com/gfx-rs/wgpu.git
synced 2025-02-27 14:23:40 +00:00

Fixes #1745: Support out-of-order module scope declarations in WGSL Fixes #1044: Forbid local variable shadowing in WGSL Fixes #2076: [wgsl-in] no error for duplicated type definition Fixes #2071: Global item does not support 'const' Fixes #2105: [wgsl-in] Type aliases for a vecN<T> doesn't work when constructing vec from a single argument Fixes #1775: Referencing a function without a return type yields an unknown identifier error. Fixes #2089: Error span reported on the declaration of a variable instead of its use Fixes #1996: [wgsl-in] Confusing error: "expected unsigned/signed integer literal, found '1'" Separate parsing from lowering by generating an AST, which desugars as much as possible down to something like Naga IR. The AST is then used to resolve identifiers while lowering to Naga IR. Co-authored-by: Teodor Tanasoaia <28601907+teoxoy@users.noreply.github.com> Co-authored-by: Jim Blandy <jimb@red-bean.com>
35 lines
1.0 KiB
WebGPU Shading Language
35 lines
1.0 KiB
WebGPU Shading Language
const SIZE: u32 = 128u;
|
|
|
|
@group(0) @binding(0)
|
|
var<storage,read_write> arr_i32: array<atomic<i32>, SIZE>;
|
|
@group(0) @binding(1)
|
|
var<storage,read_write> arr_u32: array<atomic<u32>, SIZE>;
|
|
|
|
@compute @workgroup_size(1)
|
|
fn test_atomic_compare_exchange_i32() {
|
|
for(var i = 0u; i < SIZE; i++) {
|
|
var old = atomicLoad(&arr_i32[i]);
|
|
var exchanged = false;
|
|
while(!exchanged) {
|
|
let new_ = bitcast<i32>(bitcast<f32>(old) + 1.0);
|
|
let result = atomicCompareExchangeWeak(&arr_i32[i], old, new_);
|
|
old = result.old_value;
|
|
exchanged = result.exchanged;
|
|
}
|
|
}
|
|
}
|
|
|
|
@compute @workgroup_size(1)
|
|
fn test_atomic_compare_exchange_u32() {
|
|
for(var i = 0u; i < SIZE; i++) {
|
|
var old = atomicLoad(&arr_u32[i]);
|
|
var exchanged = false;
|
|
while(!exchanged) {
|
|
let new_ = bitcast<u32>(bitcast<f32>(old) + 1.0);
|
|
let result = atomicCompareExchangeWeak(&arr_u32[i], old, new_);
|
|
old = result.old_value;
|
|
exchanged = result.exchanged;
|
|
}
|
|
}
|
|
}
|