archive-gh-me/wgpu
1
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2025-12-08 21:26:17 +00:00
Code Issues Projects Releases Wiki Activity

Add textureBarrier (#7173)

Browse Source
This commit is contained in:
Devon 2025-02-27 12:27:15 -07:00 committed by GitHub
parent 41693bc8bd
commit d6ca412732
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
19 changed files with 100 additions and 49 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
CHANGELOG.md
View File

@ -151,6 +151,7 @@ By @jamienicol in [#6929](https://github.com/gfx-rs/wgpu/pull/6929) and [#7080](
- Support @must_use attribute on function declarations. By @turbocrime in [#6801](https://github.com/gfx-rs/wgpu/pull/6801).
- Support for generating the candidate intersections from AABB geometry, and confirming the hits. By @kvark in [#7047](https://github.com/gfx-rs/wgpu/pull/7047).
- Make naga::back::spv::Function::to_words write the OpFunctionEnd instruction in itself, instead of making another call after it. By @junjunjd in [#7156](https://github.com/gfx-rs/wgpu/pull/7156).
- Add support for texture memory barriers. By @Devon7925 in [#7173](https://github.com/gfx-rs/wgpu/pull/7173).
### Changes

3
naga/src/back/glsl/mod.rs
View File

@ -4628,6 +4628,9 @@ impl<'a, W: Write> Writer<'a, W> {
if flags.contains(crate::Barrier::SUB_GROUP) {
writeln!(self.out, "{level}subgroupMemoryBarrier();")?;
}
if flags.contains(crate::Barrier::TEXTURE) {
writeln!(self.out, "{level}memoryBarrierImage();")?;
}
writeln!(self.out, "{level}barrier();")?;
Ok(())
}

3
naga/src/back/hlsl/writer.rs
View File

@ -4130,6 +4130,9 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
if barrier.contains(crate::Barrier::SUB_GROUP) {
// Does not exist in DirectX
}
if barrier.contains(crate::Barrier::TEXTURE) {
writeln!(self.out, "{level}DeviceMemoryBarrierWithGroupSync();")?;
}
Ok(())
}
}

6
naga/src/back/msl/writer.rs
View File

@ -6513,6 +6513,12 @@ template <typename A>
"{level}{NAMESPACE}::simdgroup_barrier({NAMESPACE}::mem_flags::mem_threadgroup);",
)?;
}
if flags.contains(crate::Barrier::TEXTURE) {
writeln!(
self.out,
"{level}{NAMESPACE}::threadgroup_barrier({NAMESPACE}::mem_flags::mem_texture);",
)?;
}
Ok(())
}
}

4
naga/src/back/spv/writer.rs
View File

@ -1571,6 +1571,10 @@ impl Writer {
spirv::MemorySemantics::WORKGROUP_MEMORY,
flags.contains(crate::Barrier::WORK_GROUP),
);
semantics.set(
spirv::MemorySemantics::IMAGE_MEMORY,
flags.contains(crate::Barrier::TEXTURE),
);
let exec_scope_id = if flags.contains(crate::Barrier::SUB_GROUP) {
self.get_index_constant(spirv::Scope::Subgroup as u32)
} else {

4
naga/src/back/wgsl/writer.rs
View File

@ -973,6 +973,10 @@ impl<W: Write> Writer<W> {
if barrier.contains(crate::Barrier::SUB_GROUP) {
writeln!(self.out, "{level}subgroupBarrier();")?;
}
if barrier.contains(crate::Barrier::TEXTURE) {
writeln!(self.out, "{level}textureBarrier();")?;
}
}
Statement::RayQuery { .. } => unreachable!(),
Statement::SubgroupBallot { result, predicate } => {

4
naga/src/front/spv/mod.rs
View File

@ -3846,6 +3846,10 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
.bits()
!= 0,
);
flags.set(
crate::Barrier::TEXTURE,
semantics & spirv::MemorySemantics::IMAGE_MEMORY.bits() != 0,
);
block.push(crate::Statement::Barrier(flags), span);
} else {
log::warn!("Unsupported barrier execution scope: {}", exec_scope);

8
naga/src/front/wgsl/lower/mod.rs
View File

@ -2534,6 +2534,14 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
.push(crate::Statement::Barrier(crate::Barrier::SUB_GROUP), span);
return Ok(None);
}
"textureBarrier" => {
ctx.prepare_args(arguments, 0, span).finish()?;
let rctx = ctx.runtime_expression_ctx(span)?;
rctx.block
.push(crate::Statement::Barrier(crate::Barrier::TEXTURE), span);
return Ok(None);
}
"workgroupUniformLoad" => {
let mut args = ctx.prepare_args(arguments, 1, span);
let expr = args.next()?;

2
naga/src/lib.rs
View File

@ -1371,6 +1371,8 @@ bitflags::bitflags! {
const WORK_GROUP = 1 << 1;
/// Barrier synchronizes execution across all invocations within a subgroup that execute this instruction.
const SUB_GROUP = 1 << 2;
/// Barrier synchronizes texture memory accesses in a workgroup.
const TEXTURE = 1 << 3;
}
}

3
naga/tests/in/wgsl/control-flow.toml
View File

@ -1 +1,4 @@
targets = "SPIRV | METAL | GLSL | HLSL | WGSL"
[msl]
lang_version = [1, 2]

1
naga/tests/in/wgsl/control-flow.wgsl
View File

@ -3,6 +3,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
//TODO: execution-only barrier?
storageBarrier();
workgroupBarrier();
textureBarrier();
var pos: i32;
// switch without cases

1
naga/tests/in/wgsl/separate-entry-points.wgsl
View File

@ -9,6 +9,7 @@ fn derivatives() {
fn barriers() {
storageBarrier();
workgroupBarrier();
textureBarrier();
}
@fragment

2
naga/tests/out/glsl/control-flow.main.Compute.glsl
View File

@ -149,6 +149,8 @@ void main() {
barrier();
memoryBarrierShared();
barrier();
memoryBarrierImage();
barrier();
do {
pos = 1;
} while(false);

2
naga/tests/out/glsl/separate-entry-points.compute.Compute.glsl
View File

@ -11,6 +11,8 @@ void barriers() {
barrier();
memoryBarrierShared();
barrier();
memoryBarrierImage();
barrier();
return;
}

1
naga/tests/out/hlsl/control-flow.hlsl
View File

@ -193,6 +193,7 @@ void main(uint3 global_id : SV_DispatchThreadID)
DeviceMemoryBarrierWithGroupSync();
GroupMemoryBarrierWithGroupSync();
DeviceMemoryBarrierWithGroupSync();
do {
pos = int(1);
} while(false);

3
naga/tests/out/msl/control-flow.msl
View File

@ -1,4 +1,4 @@
// language: metal1.0
// language: metal1.2
#include <metal_stdlib>
#include <simd/simd.h>
@ -178,6 +178,7 @@ kernel void main_(
int pos = {};
metal::threadgroup_barrier(metal::mem_flags::mem_device);
metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
metal::threadgroup_barrier(metal::mem_flags::mem_texture);
switch(1) {
default: {
pos = 1;

82
naga/tests/out/spv/control-flow.spvasm
View File

@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 207
; Bound: 208
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
@ -37,6 +37,7 @@ OpDecorate %176 BuiltIn GlobalInvocationId
%185 = OpConstant %4 2
%186 = OpConstant %4 72
%187 = OpConstant %4 264
%188 = OpConstant %4 2056
%8 = OpFunction %2 None %9
%7 = OpFunctionParameter %5
%6 = OpLabel
@ -318,55 +319,56 @@ OpBranch %184
%184 = OpLabel
OpControlBarrier %185 %35 %186
OpControlBarrier %185 %185 %187
OpSelectionMerge %188 None
OpSwitch %127 %189
OpControlBarrier %185 %185 %188
OpSelectionMerge %189 None
OpSwitch %127 %190
%190 = OpLabel
OpStore %182 %127
OpBranch %189
%189 = OpLabel
OpStore %182 %127
OpBranch %188
%188 = OpLabel
%190 = OpLoad %5 %182
OpSelectionMerge %191 None
OpSwitch %190 %196 1 %192 2 %193 3 %194 4 %194 5 %195 6 %196
%192 = OpLabel
OpStore %182 %16
OpBranch %191
%191 = OpLoad %5 %182
OpSelectionMerge %192 None
OpSwitch %191 %197 1 %193 2 %194 3 %195 4 %195 5 %196 6 %197
%193 = OpLabel
OpStore %182 %127
OpBranch %191
%194 = OpLabel
OpStore %182 %128
OpBranch %191
%195 = OpLabel
OpStore %182 %180
OpBranch %191
%196 = OpLabel
OpStore %182 %181
OpBranch %191
%191 = OpLabel
OpSelectionMerge %197 None
OpSwitch %33 %199 0 %198
%198 = OpLabel
OpBranch %197
%199 = OpLabel
OpBranch %197
%197 = OpLabel
%200 = OpLoad %5 %182
OpSelectionMerge %201 None
OpSwitch %200 %206 1 %202 2 %203 3 %204 4 %205
%202 = OpLabel
OpStore %182 %16
OpBranch %201
%203 = OpLabel
OpBranch %192
%194 = OpLabel
OpStore %182 %127
OpReturn
%204 = OpLabel
OpBranch %192
%195 = OpLabel
OpStore %182 %128
OpBranch %192
%196 = OpLabel
OpStore %182 %180
OpBranch %192
%197 = OpLabel
OpStore %182 %181
OpBranch %192
%192 = OpLabel
OpSelectionMerge %198 None
OpSwitch %33 %200 0 %199
%199 = OpLabel
OpBranch %198
%200 = OpLabel
OpBranch %198
%198 = OpLabel
%201 = OpLoad %5 %182
OpSelectionMerge %202 None
OpSwitch %201 %207 1 %203 2 %204 3 %205 4 %206
%203 = OpLabel
OpStore %182 %16
OpBranch %202
%204 = OpLabel
OpStore %182 %127
OpReturn
%205 = OpLabel
OpStore %182 %128
OpReturn
%206 = OpLabel
OpReturn
%207 = OpLabel
OpStore %182 %180
OpReturn
%201 = OpLabel
%202 = OpLabel
OpReturn
OpFunctionEnd

18
naga/tests/out/spv/separate-entry-points.compute.spvasm
View File

@ -1,12 +1,12 @@
; SPIR-V
; Version: 1.0
; Generator: rspirv
; Bound: 18
; Bound: 19
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %15 "compute"
OpExecutionMode %15 LocalSize 1 1 1
OpEntryPoint GLCompute %16 "compute"
OpExecutionMode %16 LocalSize 1 1 1
%2 = OpTypeVoid
%4 = OpTypeFloat 32
%3 = OpTypeVector %4 4
@ -16,18 +16,20 @@ OpExecutionMode %15 LocalSize 1 1 1
%11 = OpConstant %10 1
%12 = OpConstant %10 72
%13 = OpConstant %10 264
%14 = OpConstant %10 2056
%6 = OpFunction %2 None %7
%5 = OpLabel
OpBranch %8
%8 = OpLabel
OpControlBarrier %9 %11 %12
OpControlBarrier %9 %9 %13
OpControlBarrier %9 %9 %14
OpReturn
OpFunctionEnd
%15 = OpFunction %2 None %7
%14 = OpLabel
OpBranch %16
%16 = OpLabel
%17 = OpFunctionCall %2 %6
%16 = OpFunction %2 None %7
%15 = OpLabel
OpBranch %17
%17 = OpLabel
%18 = OpFunctionCall %2 %6
OpReturn
OpFunctionEnd

1
naga/tests/out/wgsl/control-flow.wgsl
View File

@ -122,6 +122,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
storageBarrier();
workgroupBarrier();
textureBarrier();
switch 1i {
default: {
pos = 1i;