mirror of
https://github.com/gfx-rs/wgpu.git
synced 2024-11-22 14:55:05 +00:00
abba12ae4e
Add the following flags to `wgpu_types::Features`: - `SHADER_INT64_ATOMIC_ALL_OPS` enables all atomic operations on `atomic<i64>` and `atomic<u64>` values. - `SHADER_INT64_ATOMIC_MIN_MAX` is a subset of the above, enabling only `AtomicFunction::Min` and `AtomicFunction::Max` operations on `atomic<i64>` and `atomic<u64>` values in the `Storage` address space. These are the only 64-bit atomic operations available on Metal as of 3.1. Add corresponding flags to `naga::valid::Capabilities`. These are supported by the WGSL front end, and all Naga backends. Platform support: - On Direct3d 12, in `D3D12_FEATURE_DATA_D3D12_OPTIONS9`, if `AtomicInt64OnTypedResourceSupported` and `AtomicInt64OnGroupSharedSupported` are both available, then both wgpu features described above are available. - On Metal, `SHADER_INT64_ATOMIC_MIN_MAX` is available on Apple9 hardware, and on hardware that advertises both Apple8 and Mac2 support. This also requires Metal Shading Language 2.4 or later. Metal does not yet support the more general `SHADER_INT64_ATOMIC_ALL_OPS`. - On Vulkan, if the `VK_KHR_shader_atomic_int64` extension is available with both the `shader_buffer_int64_atomics` and `shader_shared_int64_atomics` features, then both wgpu features described above are available.
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_i64: array<atomic<i64>, SIZE>;
|
|
@group(0) @binding(1)
|
|
var<storage,read_write> arr_u64: array<atomic<u64>, SIZE>;
|
|
|
|
@compute @workgroup_size(1)
|
|
fn test_atomic_compare_exchange_i64() {
|
|
for(var i = 0u; i < SIZE; i++) {
|
|
var old : i64 = atomicLoad(&arr_i64[i]);
|
|
var exchanged = false;
|
|
while(!exchanged) {
|
|
let new_ : i64 = bitcast<i64>(old + 10li);
|
|
let result = atomicCompareExchangeWeak(&arr_i64[i], old, new_);
|
|
old = result.old_value;
|
|
exchanged = result.exchanged;
|
|
}
|
|
}
|
|
}
|
|
|
|
@compute @workgroup_size(1)
|
|
fn test_atomic_compare_exchange_u64() {
|
|
for(var i = 0u; i < SIZE; i++) {
|
|
var old : u64 = atomicLoad(&arr_u64[i]);
|
|
var exchanged = false;
|
|
while(!exchanged) {
|
|
let new_ : u64 = bitcast<u64>(old + 10lu);
|
|
let result = atomicCompareExchangeWeak(&arr_u64[i], old, new_);
|
|
old = result.old_value;
|
|
exchanged = result.exchanged;
|
|
}
|
|
}
|
|
}
|