mirror of
https://github.com/gfx-rs/wgpu.git
synced 2024-11-22 14:55:05 +00:00
fc85e4f970
* add parsing for spirv::Op::AtomicLoad and spirv::Op::AtomicStore * spv-in parse AtomicExchange and AtomicCompareExchange * add atomic i decrement * bookend atomic store statement with emmitter.finish/emitter.start to suppress a double load expression bookend atomic result expressions with emitter.finish/start to prevent double defs * add atomic iadd, isub, smin, umin, smax, umax, and, or, xor * parse atomic flag test and set, parse atomic flag clear * remove atomic compare exchange work * changelog * moved spirv tests into front/spv/mod.rs * feature gate atomic spv tests because they require wgsl-[in,out] * BlockContext::get_contained_global_variable returns Result * Generate spans covering the entire instruction. Granted, there is pre-existing code in the SPIR-V front end that gets this wrong, but: It doesn't make sense to read `self.data_offset`, and then immediately pass that to `self.span_from_with_op`. The point of that function is to make the span cover the entire instruction, operands included. * Move `From` implementation into spv front end * doc comments, minor cleanups * remove parsing of OpAtomicFlagClear and OpAtomicFlagTestAndSet * sync atomic spvasm files --------- Co-authored-by: Jim Blandy <jimb@red-bean.com>
307 lines
7.7 KiB
Plaintext
307 lines
7.7 KiB
Plaintext
(
|
|
types: [
|
|
(
|
|
name: None,
|
|
inner: Scalar((
|
|
kind: Uint,
|
|
width: 4,
|
|
)),
|
|
),
|
|
(
|
|
name: None,
|
|
inner: Scalar((
|
|
kind: Bool,
|
|
width: 1,
|
|
)),
|
|
),
|
|
(
|
|
name: None,
|
|
inner: Pointer(
|
|
base: 0,
|
|
space: Storage(
|
|
access: ("LOAD | STORE"),
|
|
),
|
|
),
|
|
),
|
|
(
|
|
name: None,
|
|
inner: Struct(
|
|
members: [
|
|
(
|
|
name: None,
|
|
ty: 0,
|
|
binding: None,
|
|
offset: 0,
|
|
),
|
|
],
|
|
span: 4,
|
|
),
|
|
),
|
|
(
|
|
name: None,
|
|
inner: Pointer(
|
|
base: 3,
|
|
space: Storage(
|
|
access: ("LOAD | STORE"),
|
|
),
|
|
),
|
|
),
|
|
(
|
|
name: None,
|
|
inner: Atomic((
|
|
kind: Uint,
|
|
width: 4,
|
|
)),
|
|
),
|
|
(
|
|
name: None,
|
|
inner: Struct(
|
|
members: [
|
|
(
|
|
name: None,
|
|
ty: 5,
|
|
binding: None,
|
|
offset: 0,
|
|
),
|
|
],
|
|
span: 4,
|
|
),
|
|
),
|
|
],
|
|
special_types: (
|
|
ray_desc: None,
|
|
ray_intersection: None,
|
|
predeclared_types: {},
|
|
),
|
|
constants: [
|
|
(
|
|
name: None,
|
|
ty: 0,
|
|
init: 0,
|
|
),
|
|
(
|
|
name: None,
|
|
ty: 0,
|
|
init: 1,
|
|
),
|
|
(
|
|
name: None,
|
|
ty: 1,
|
|
init: 2,
|
|
),
|
|
(
|
|
name: None,
|
|
ty: 0,
|
|
init: 3,
|
|
),
|
|
(
|
|
name: None,
|
|
ty: 1,
|
|
init: 4,
|
|
),
|
|
(
|
|
name: None,
|
|
ty: 0,
|
|
init: 5,
|
|
),
|
|
],
|
|
overrides: [],
|
|
global_variables: [
|
|
(
|
|
name: None,
|
|
space: Storage(
|
|
access: ("LOAD | STORE"),
|
|
),
|
|
binding: Some((
|
|
group: 0,
|
|
binding: 0,
|
|
)),
|
|
ty: 6,
|
|
init: None,
|
|
),
|
|
(
|
|
name: None,
|
|
space: Storage(
|
|
access: ("LOAD"),
|
|
),
|
|
binding: Some((
|
|
group: 0,
|
|
binding: 1,
|
|
)),
|
|
ty: 3,
|
|
init: None,
|
|
),
|
|
],
|
|
global_expressions: [
|
|
Literal(U32(0)),
|
|
Literal(U32(2)),
|
|
Literal(Bool(false)),
|
|
Literal(U32(1)),
|
|
Literal(Bool(true)),
|
|
ZeroValue(0),
|
|
],
|
|
functions: [
|
|
(
|
|
name: None,
|
|
arguments: [],
|
|
result: None,
|
|
local_variables: [
|
|
(
|
|
name: Some("phi_23"),
|
|
ty: 0,
|
|
init: None,
|
|
),
|
|
(
|
|
name: Some("phi_24"),
|
|
ty: 0,
|
|
init: None,
|
|
),
|
|
],
|
|
expressions: [
|
|
GlobalVariable(0),
|
|
GlobalVariable(1),
|
|
Constant(4),
|
|
Constant(2),
|
|
Constant(5),
|
|
Constant(3),
|
|
Constant(1),
|
|
Constant(0),
|
|
AccessIndex(
|
|
base: 0,
|
|
index: 0,
|
|
),
|
|
AccessIndex(
|
|
base: 1,
|
|
index: 0,
|
|
),
|
|
LocalVariable(0),
|
|
Load(
|
|
pointer: 10,
|
|
),
|
|
Load(
|
|
pointer: 9,
|
|
),
|
|
Binary(
|
|
op: GreaterEqual,
|
|
left: 11,
|
|
right: 12,
|
|
),
|
|
AtomicResult(
|
|
ty: 0,
|
|
comparison: false,
|
|
),
|
|
Literal(U32(1)),
|
|
Binary(
|
|
op: Add,
|
|
left: 11,
|
|
right: 5,
|
|
),
|
|
LocalVariable(1),
|
|
Load(
|
|
pointer: 17,
|
|
),
|
|
Select(
|
|
condition: 13,
|
|
accept: 3,
|
|
reject: 2,
|
|
),
|
|
Unary(
|
|
op: LogicalNot,
|
|
expr: 19,
|
|
),
|
|
LocalVariable(0),
|
|
LocalVariable(1),
|
|
],
|
|
named_expressions: {},
|
|
body: [
|
|
Emit((
|
|
start: 8,
|
|
end: 10,
|
|
)),
|
|
Store(
|
|
pointer: 21,
|
|
value: 7,
|
|
),
|
|
Loop(
|
|
body: [
|
|
Emit((
|
|
start: 11,
|
|
end: 12,
|
|
)),
|
|
Emit((
|
|
start: 12,
|
|
end: 14,
|
|
)),
|
|
If(
|
|
condition: 13,
|
|
accept: [
|
|
Store(
|
|
pointer: 22,
|
|
value: 4,
|
|
),
|
|
],
|
|
reject: [
|
|
Atomic(
|
|
pointer: 8,
|
|
fun: Add,
|
|
value: 15,
|
|
result: Some(14),
|
|
),
|
|
Emit((
|
|
start: 16,
|
|
end: 17,
|
|
)),
|
|
Store(
|
|
pointer: 22,
|
|
value: 16,
|
|
),
|
|
],
|
|
),
|
|
Emit((
|
|
start: 18,
|
|
end: 20,
|
|
)),
|
|
Continue,
|
|
],
|
|
continuing: [
|
|
Emit((
|
|
start: 20,
|
|
end: 21,
|
|
)),
|
|
Store(
|
|
pointer: 21,
|
|
value: 18,
|
|
),
|
|
],
|
|
break_if: Some(20),
|
|
),
|
|
Return(
|
|
value: None,
|
|
),
|
|
],
|
|
),
|
|
],
|
|
entry_points: [
|
|
(
|
|
name: "stage::test_atomic_i_increment",
|
|
stage: Compute,
|
|
early_depth_test: None,
|
|
workgroup_size: (32, 1, 1),
|
|
function: (
|
|
name: Some("stage::test_atomic_i_increment_wrap"),
|
|
arguments: [],
|
|
result: None,
|
|
local_variables: [],
|
|
expressions: [],
|
|
named_expressions: {},
|
|
body: [
|
|
Call(
|
|
function: 0,
|
|
arguments: [],
|
|
result: None,
|
|
),
|
|
],
|
|
),
|
|
),
|
|
],
|
|
) |