Barriers #817
Barriers #817
Conversation
|
I just tried out this branch rather than the barrier support I've been using, and it seems to be working well! One question: why are the WGSL built-in functions in this PR named differently than the WGSL spec? |
|
@acowley thank you for diligence! Fixed now :) |
There was a problem hiding this comment.
It looks good overall. Three concerns:
- It seems like control-only barriers will be dropped in GLSL and MSL output, which is inconsistent with the handling of SPIR-V.
- I think
Statement::Barrierneeds better docs. The way it's being used by the front and back ends seems inconsistent, so I wasn't confident that I could infer the intent from the code. - The signedness of cached index constants in SPIR-V output changed, which is a bit surprising.
| let mut flags = crate::Barrier::empty(); | ||
| flags.set( | ||
| crate::Barrier::STORAGE, | ||
| semantics & spirv::MemorySemantics::UNIFORM_MEMORY.bits() != 0, | ||
| ); | ||
| flags.set( | ||
| crate::Barrier::WORK_GROUP, | ||
| semantics & spirv::MemorySemantics::WORKGROUP_MEMORY.bits() != 0, | ||
| ); | ||
| block.push(crate::Statement::Barrier(flags)); |
There was a problem hiding this comment.
I would expect this code to be promoting all "smaller" boundaries to the next larger boundary we support. That is, if the OpControlBarrier requests MemorySemantics::SUBGROUP_MEMORY, which I believe is a smaller subdivision than a workgroup, then that should also set Barrier::WORK_GROUP, because that's the only way we have to guarantee the sematics requested.
There was a problem hiding this comment.
Yes, we definitely want a promotion like this
There was a problem hiding this comment.
actually, I don't know what SubgroupMemory is, and SPIR-V spec isn't clear on this
| write!(self.out, "{}", INDENT.repeat(indent))?; | ||
| writeln!(self.out, "discard;")? | ||
| Statement::Kill => writeln!(self.out, "{}discard;", INDENT.repeat(indent))?, | ||
| // Just issue a memory barrier, ignoring the flags. |
There was a problem hiding this comment.
We're not 'ignoring the flags' here, are we?
| /// Synchronize accesses within the work group. | ||
| Barrier(Barrier), |
There was a problem hiding this comment.
Given the variety of barriers supported by our various target languages, I think this could use much more detailed documentation. For example, a SPIR-V OpControlBarrier instruction always generates a Statement::Barrier, but we emit no code in GLSL or MSL for a Statement::Barrier unless it has flags set, which makes it hard to someone reading the code to infer the intent. A comment here would help the reader understand which end is incorrect, or at least explain why the disagreement is acceptable.
There was a problem hiding this comment.
Agreed. And also, the GLSL and MSL are going to be fixed now to actually issue execution barriers.
| if flags.contains(crate::Barrier::WORK_GROUP) { | ||
| writeln!(self.out, "{}barrier();", INDENT.repeat(indent))?; | ||
| } |
There was a problem hiding this comment.
If none of the flags are set, this doesn't emit anything. But it seems to me that SPIR-V OpControlBarrier, which produces a Statement::Barrier with no flags set, should be a GLSL barrier();, at least.
| crate::Statement::Barrier(flags) => { | ||
| if flags.contains(crate::Barrier::STORAGE) { | ||
| writeln!( | ||
| self.out, | ||
| "{}{}::threadgroup_barrier({}::mem_flags::mem_device);", | ||
| level, NAMESPACE, NAMESPACE, | ||
| )?; | ||
| } | ||
| if flags.contains(crate::Barrier::WORK_GROUP) { | ||
| writeln!( | ||
| self.out, | ||
| "{}{}::threadgroup_barrier({}::mem_flags::mem_threadgroup);", | ||
| level, NAMESPACE, NAMESPACE, | ||
| )?; | ||
| } | ||
| } |
There was a problem hiding this comment.
This, too, is inconsistent with the SPIR-V behavior. With no flags set, it seems like this should be
threadgroup_barrier(mem_flags::mem_none);
| fn get_index_constant( | ||
| &mut self, | ||
| index: Word, | ||
| types: &Arena<crate::Type>, | ||
| ) -> Result<Word, Error> { | ||
| self.get_constant_scalar(crate::ScalarValue::Uint(index as _), 4, types) | ||
| } |
There was a problem hiding this comment.
This changes the types of indexes used as LOD selectors and OpAccessChain operands from Sint to Uint. I couldn't figure out the rules for LOD selection, but OpAccessChain says the indices are "treated as signed", so is this what we want?
There was a problem hiding this comment.
SPIR-V doesn't care, basically. It will reinterpret them as signed, and it's still fine.
Fixes #583