2023-09-07 20:22:44 +00:00
|
|
|
(
|
|
|
|
types: [
|
|
|
|
(
|
|
|
|
name: None,
|
[naga] Introduce `Scalar` type to IR.
Introduce a new struct type, `Scalar`, combining a `ScalarKind` and a
`Bytes` width, and use this whenever such pairs of values are passed
around.
In particular, use `Scalar` in `TypeInner` variants `Scalar`, `Vector`,
`Atomic`, and `ValuePointer`.
Introduce associated `Scalar` constants `I32`, `U32`, `F32`, `BOOL`
and `F64`, for common cases.
Introduce a helper function `Scalar::float` for constructing `Float`
scalars of a given width, for dealing with `TypeInner::Matrix`, which
only supplies the scalar width of its elements, not a kind.
Introduce helper functions on `Literal` and `TypeInner`, to produce
the `Scalar` describing elements' values.
Use `Scalar` in `wgpu_core::validation::NumericType` as well.
2023-11-10 20:31:29 +00:00
|
|
|
inner: Scalar((
|
2023-09-07 20:22:44 +00:00
|
|
|
kind: Uint,
|
|
|
|
width: 4,
|
[naga] Introduce `Scalar` type to IR.
Introduce a new struct type, `Scalar`, combining a `ScalarKind` and a
`Bytes` width, and use this whenever such pairs of values are passed
around.
In particular, use `Scalar` in `TypeInner` variants `Scalar`, `Vector`,
`Atomic`, and `ValuePointer`.
Introduce associated `Scalar` constants `I32`, `U32`, `F32`, `BOOL`
and `F64`, for common cases.
Introduce a helper function `Scalar::float` for constructing `Float`
scalars of a given width, for dealing with `TypeInner::Matrix`, which
only supplies the scalar width of its elements, not a kind.
Introduce helper functions on `Literal` and `TypeInner`, to produce
the `Scalar` describing elements' values.
Use `Scalar` in `wgpu_core::validation::NumericType` as well.
2023-11-10 20:31:29 +00:00
|
|
|
)),
|
2023-09-07 20:22:44 +00:00
|
|
|
),
|
|
|
|
(
|
|
|
|
name: None,
|
|
|
|
inner: Array(
|
2024-06-20 06:11:11 +00:00
|
|
|
base: 0,
|
2023-09-07 20:22:44 +00:00
|
|
|
size: Dynamic,
|
|
|
|
stride: 4,
|
|
|
|
),
|
|
|
|
),
|
|
|
|
(
|
|
|
|
name: Some("PrimeIndices"),
|
|
|
|
inner: Struct(
|
|
|
|
members: [
|
|
|
|
(
|
|
|
|
name: Some("data"),
|
2024-06-20 06:11:11 +00:00
|
|
|
ty: 1,
|
2023-09-07 20:22:44 +00:00
|
|
|
binding: None,
|
|
|
|
offset: 0,
|
|
|
|
),
|
|
|
|
],
|
|
|
|
span: 4,
|
|
|
|
),
|
|
|
|
),
|
|
|
|
(
|
|
|
|
name: None,
|
|
|
|
inner: Vector(
|
|
|
|
size: Tri,
|
[naga] Introduce `Scalar` type to IR.
Introduce a new struct type, `Scalar`, combining a `ScalarKind` and a
`Bytes` width, and use this whenever such pairs of values are passed
around.
In particular, use `Scalar` in `TypeInner` variants `Scalar`, `Vector`,
`Atomic`, and `ValuePointer`.
Introduce associated `Scalar` constants `I32`, `U32`, `F32`, `BOOL`
and `F64`, for common cases.
Introduce a helper function `Scalar::float` for constructing `Float`
scalars of a given width, for dealing with `TypeInner::Matrix`, which
only supplies the scalar width of its elements, not a kind.
Introduce helper functions on `Literal` and `TypeInner`, to produce
the `Scalar` describing elements' values.
Use `Scalar` in `wgpu_core::validation::NumericType` as well.
2023-11-10 20:31:29 +00:00
|
|
|
scalar: (
|
|
|
|
kind: Uint,
|
|
|
|
width: 4,
|
|
|
|
),
|
2023-09-07 20:22:44 +00:00
|
|
|
),
|
|
|
|
),
|
|
|
|
],
|
|
|
|
special_types: (
|
|
|
|
ray_desc: None,
|
|
|
|
ray_intersection: None,
|
|
|
|
predeclared_types: {},
|
|
|
|
),
|
|
|
|
constants: [],
|
2023-12-07 19:19:43 +00:00
|
|
|
overrides: [],
|
2023-09-07 20:22:44 +00:00
|
|
|
global_variables: [
|
|
|
|
(
|
|
|
|
name: Some("v_indices"),
|
|
|
|
space: Storage(
|
|
|
|
access: ("LOAD | STORE"),
|
|
|
|
),
|
|
|
|
binding: Some((
|
|
|
|
group: 0,
|
|
|
|
binding: 0,
|
|
|
|
)),
|
2024-06-20 06:11:11 +00:00
|
|
|
ty: 2,
|
2023-09-07 20:22:44 +00:00
|
|
|
init: None,
|
|
|
|
),
|
|
|
|
],
|
2024-02-14 14:25:23 +00:00
|
|
|
global_expressions: [],
|
2023-09-07 20:22:44 +00:00
|
|
|
functions: [
|
|
|
|
(
|
|
|
|
name: Some("collatz_iterations"),
|
|
|
|
arguments: [
|
|
|
|
(
|
|
|
|
name: Some("n_base"),
|
2024-06-20 06:11:11 +00:00
|
|
|
ty: 0,
|
2023-09-07 20:22:44 +00:00
|
|
|
binding: None,
|
|
|
|
),
|
|
|
|
],
|
|
|
|
result: Some((
|
2024-06-20 06:11:11 +00:00
|
|
|
ty: 0,
|
2023-09-07 20:22:44 +00:00
|
|
|
binding: None,
|
|
|
|
)),
|
|
|
|
local_variables: [
|
|
|
|
(
|
|
|
|
name: Some("n"),
|
2024-06-20 06:11:11 +00:00
|
|
|
ty: 0,
|
2023-09-07 20:22:44 +00:00
|
|
|
init: None,
|
|
|
|
),
|
|
|
|
(
|
|
|
|
name: Some("i"),
|
2024-06-20 06:11:11 +00:00
|
|
|
ty: 0,
|
|
|
|
init: Some(2),
|
2023-09-07 20:22:44 +00:00
|
|
|
),
|
|
|
|
],
|
|
|
|
expressions: [
|
|
|
|
FunctionArgument(0),
|
2024-06-20 06:11:11 +00:00
|
|
|
LocalVariable(0),
|
2023-09-07 20:22:44 +00:00
|
|
|
Literal(U32(0)),
|
2024-06-20 06:11:11 +00:00
|
|
|
LocalVariable(1),
|
2023-09-07 20:22:44 +00:00
|
|
|
Load(
|
2024-06-20 06:11:11 +00:00
|
|
|
pointer: 1,
|
2023-09-07 20:22:44 +00:00
|
|
|
),
|
|
|
|
Literal(U32(1)),
|
|
|
|
Binary(
|
|
|
|
op: Greater,
|
2024-06-20 06:11:11 +00:00
|
|
|
left: 4,
|
|
|
|
right: 5,
|
2023-09-07 20:22:44 +00:00
|
|
|
),
|
|
|
|
Load(
|
2024-06-20 06:11:11 +00:00
|
|
|
pointer: 1,
|
2023-09-07 20:22:44 +00:00
|
|
|
),
|
|
|
|
Literal(U32(2)),
|
|
|
|
Binary(
|
|
|
|
op: Modulo,
|
2024-06-20 06:11:11 +00:00
|
|
|
left: 7,
|
|
|
|
right: 8,
|
2023-09-07 20:22:44 +00:00
|
|
|
),
|
|
|
|
Literal(U32(0)),
|
|
|
|
Binary(
|
|
|
|
op: Equal,
|
2024-06-20 06:11:11 +00:00
|
|
|
left: 9,
|
|
|
|
right: 10,
|
2023-09-07 20:22:44 +00:00
|
|
|
),
|
|
|
|
Load(
|
2024-06-20 06:11:11 +00:00
|
|
|
pointer: 1,
|
2023-09-07 20:22:44 +00:00
|
|
|
),
|
|
|
|
Literal(U32(2)),
|
|
|
|
Binary(
|
|
|
|
op: Divide,
|
2024-06-20 06:11:11 +00:00
|
|
|
left: 12,
|
|
|
|
right: 13,
|
2023-09-07 20:22:44 +00:00
|
|
|
),
|
|
|
|
Literal(U32(3)),
|
|
|
|
Load(
|
2024-06-20 06:11:11 +00:00
|
|
|
pointer: 1,
|
2023-09-07 20:22:44 +00:00
|
|
|
),
|
|
|
|
Binary(
|
|
|
|
op: Multiply,
|
2024-06-20 06:11:11 +00:00
|
|
|
left: 15,
|
|
|
|
right: 16,
|
2023-09-07 20:22:44 +00:00
|
|
|
),
|
|
|
|
Literal(U32(1)),
|
|
|
|
Binary(
|
|
|
|
op: Add,
|
2024-06-20 06:11:11 +00:00
|
|
|
left: 17,
|
|
|
|
right: 18,
|
2023-09-07 20:22:44 +00:00
|
|
|
),
|
|
|
|
Load(
|
2024-06-20 06:11:11 +00:00
|
|
|
pointer: 3,
|
2023-09-07 20:22:44 +00:00
|
|
|
),
|
|
|
|
Literal(U32(1)),
|
|
|
|
Binary(
|
|
|
|
op: Add,
|
2024-06-20 06:11:11 +00:00
|
|
|
left: 20,
|
|
|
|
right: 21,
|
2023-09-07 20:22:44 +00:00
|
|
|
),
|
|
|
|
Load(
|
2024-06-20 06:11:11 +00:00
|
|
|
pointer: 3,
|
2023-09-07 20:22:44 +00:00
|
|
|
),
|
|
|
|
],
|
|
|
|
named_expressions: {
|
2024-06-20 06:11:11 +00:00
|
|
|
0: "n_base",
|
2023-09-07 20:22:44 +00:00
|
|
|
},
|
|
|
|
body: [
|
|
|
|
Store(
|
2024-06-20 06:11:11 +00:00
|
|
|
pointer: 1,
|
|
|
|
value: 0,
|
2023-09-07 20:22:44 +00:00
|
|
|
),
|
|
|
|
Loop(
|
|
|
|
body: [
|
|
|
|
Emit((
|
|
|
|
start: 4,
|
|
|
|
end: 5,
|
|
|
|
)),
|
|
|
|
Emit((
|
|
|
|
start: 6,
|
|
|
|
end: 7,
|
|
|
|
)),
|
|
|
|
If(
|
2024-06-20 06:11:11 +00:00
|
|
|
condition: 6,
|
2023-09-07 20:22:44 +00:00
|
|
|
accept: [],
|
|
|
|
reject: [
|
|
|
|
Break,
|
|
|
|
],
|
|
|
|
),
|
|
|
|
Block([
|
|
|
|
Emit((
|
|
|
|
start: 7,
|
|
|
|
end: 8,
|
|
|
|
)),
|
|
|
|
Emit((
|
|
|
|
start: 9,
|
|
|
|
end: 10,
|
|
|
|
)),
|
|
|
|
Emit((
|
|
|
|
start: 11,
|
|
|
|
end: 12,
|
|
|
|
)),
|
|
|
|
If(
|
2024-06-20 06:11:11 +00:00
|
|
|
condition: 11,
|
2023-09-07 20:22:44 +00:00
|
|
|
accept: [
|
|
|
|
Emit((
|
|
|
|
start: 12,
|
|
|
|
end: 13,
|
|
|
|
)),
|
|
|
|
Emit((
|
|
|
|
start: 14,
|
|
|
|
end: 15,
|
|
|
|
)),
|
|
|
|
Store(
|
2024-06-20 06:11:11 +00:00
|
|
|
pointer: 1,
|
|
|
|
value: 14,
|
2023-09-07 20:22:44 +00:00
|
|
|
),
|
|
|
|
],
|
|
|
|
reject: [
|
|
|
|
Emit((
|
|
|
|
start: 16,
|
|
|
|
end: 18,
|
|
|
|
)),
|
|
|
|
Emit((
|
|
|
|
start: 19,
|
|
|
|
end: 20,
|
|
|
|
)),
|
|
|
|
Store(
|
2024-06-20 06:11:11 +00:00
|
|
|
pointer: 1,
|
|
|
|
value: 19,
|
2023-09-07 20:22:44 +00:00
|
|
|
),
|
|
|
|
],
|
|
|
|
),
|
|
|
|
Emit((
|
|
|
|
start: 20,
|
|
|
|
end: 21,
|
|
|
|
)),
|
|
|
|
Emit((
|
|
|
|
start: 22,
|
|
|
|
end: 23,
|
|
|
|
)),
|
|
|
|
Store(
|
2024-06-20 06:11:11 +00:00
|
|
|
pointer: 3,
|
|
|
|
value: 22,
|
2023-09-07 20:22:44 +00:00
|
|
|
),
|
|
|
|
]),
|
|
|
|
],
|
|
|
|
continuing: [],
|
|
|
|
break_if: None,
|
|
|
|
),
|
|
|
|
Emit((
|
|
|
|
start: 23,
|
|
|
|
end: 24,
|
|
|
|
)),
|
|
|
|
Return(
|
2024-06-20 06:11:11 +00:00
|
|
|
value: Some(23),
|
2023-09-07 20:22:44 +00:00
|
|
|
),
|
|
|
|
],
|
|
|
|
),
|
|
|
|
],
|
|
|
|
entry_points: [
|
|
|
|
(
|
|
|
|
name: "main",
|
|
|
|
stage: Compute,
|
|
|
|
early_depth_test: None,
|
|
|
|
workgroup_size: (1, 1, 1),
|
|
|
|
function: (
|
|
|
|
name: Some("main"),
|
|
|
|
arguments: [
|
|
|
|
(
|
|
|
|
name: Some("global_id"),
|
2024-06-20 06:11:11 +00:00
|
|
|
ty: 3,
|
2023-09-07 20:22:44 +00:00
|
|
|
binding: Some(BuiltIn(GlobalInvocationId)),
|
|
|
|
),
|
|
|
|
],
|
|
|
|
result: None,
|
|
|
|
local_variables: [],
|
|
|
|
expressions: [
|
|
|
|
FunctionArgument(0),
|
2024-06-20 06:11:11 +00:00
|
|
|
GlobalVariable(0),
|
2023-09-07 20:22:44 +00:00
|
|
|
AccessIndex(
|
2024-06-20 06:11:11 +00:00
|
|
|
base: 1,
|
2023-09-07 20:22:44 +00:00
|
|
|
index: 0,
|
|
|
|
),
|
|
|
|
AccessIndex(
|
2024-06-20 06:11:11 +00:00
|
|
|
base: 0,
|
2023-09-07 20:22:44 +00:00
|
|
|
index: 0,
|
|
|
|
),
|
|
|
|
Access(
|
2024-06-20 06:11:11 +00:00
|
|
|
base: 2,
|
|
|
|
index: 3,
|
2023-09-07 20:22:44 +00:00
|
|
|
),
|
2024-06-20 06:11:11 +00:00
|
|
|
GlobalVariable(0),
|
2023-09-07 20:22:44 +00:00
|
|
|
AccessIndex(
|
2024-06-20 06:11:11 +00:00
|
|
|
base: 5,
|
2023-09-07 20:22:44 +00:00
|
|
|
index: 0,
|
|
|
|
),
|
|
|
|
AccessIndex(
|
2024-06-20 06:11:11 +00:00
|
|
|
base: 0,
|
2023-09-07 20:22:44 +00:00
|
|
|
index: 0,
|
|
|
|
),
|
|
|
|
Access(
|
2024-06-20 06:11:11 +00:00
|
|
|
base: 6,
|
|
|
|
index: 7,
|
2023-09-07 20:22:44 +00:00
|
|
|
),
|
|
|
|
Load(
|
2024-06-20 06:11:11 +00:00
|
|
|
pointer: 8,
|
2023-09-07 20:22:44 +00:00
|
|
|
),
|
2024-06-20 06:11:11 +00:00
|
|
|
CallResult(0),
|
2023-09-07 20:22:44 +00:00
|
|
|
],
|
|
|
|
named_expressions: {
|
2024-06-20 06:11:11 +00:00
|
|
|
0: "global_id",
|
2023-09-07 20:22:44 +00:00
|
|
|
},
|
|
|
|
body: [
|
|
|
|
Emit((
|
|
|
|
start: 2,
|
|
|
|
end: 5,
|
|
|
|
)),
|
|
|
|
Emit((
|
|
|
|
start: 6,
|
|
|
|
end: 10,
|
|
|
|
)),
|
|
|
|
Call(
|
2024-06-20 06:11:11 +00:00
|
|
|
function: 0,
|
2023-09-07 20:22:44 +00:00
|
|
|
arguments: [
|
2024-06-20 06:11:11 +00:00
|
|
|
9,
|
2023-09-07 20:22:44 +00:00
|
|
|
],
|
2024-06-20 06:11:11 +00:00
|
|
|
result: Some(10),
|
2023-09-07 20:22:44 +00:00
|
|
|
),
|
|
|
|
Store(
|
2024-06-20 06:11:11 +00:00
|
|
|
pointer: 4,
|
|
|
|
value: 10,
|
2023-09-07 20:22:44 +00:00
|
|
|
),
|
|
|
|
Return(
|
|
|
|
value: None,
|
|
|
|
),
|
|
|
|
],
|
|
|
|
),
|
|
|
|
),
|
|
|
|
],
|
2024-08-23 17:44:48 +00:00
|
|
|
diagnostic_filters: [],
|
|
|
|
diagnostic_filter_leaf: None,
|
2023-09-07 20:22:44 +00:00
|
|
|
)
|