github/wgpu
1
1
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2026-04-22 03:02:01 -04:00
Code Issues Packages Projects Releases Wiki Activity

chore: [spv-in] atomics snapshot tests (#6692)

Browse Source 
* chore: [spv-in] clean up atomic upgrade tests

* add output test files

* update changelog

* remove extraneous snapshot out files
This commit is contained in:
Schell Carl Scivally
2024-12-10 16:54:02 +13:00
committed by GitHub
parent 90859b4ea1
commit 8b93a71129
11 changed files with 330 additions and 686 deletions
Download Patch File Download Diff File Expand all files Collapse all files

1
CHANGELOG.md
View File

@@ -107,6 +107,7 @@ By @ErichDonGubler in [#6456](https://github.com/gfx-rs/wgpu/pull/6456), [#6148]
#### Naga
- Clean up tests for atomic operations support in SPIR-V frontend. By @schell in [#6692](https://github.com/gfx-rs/wgpu/pull/6692)
- Fix an issue where `naga` CLI would incorrectly skip the first positional argument when `--stdin-file-path` was specified. By @ErichDonGubler in [#6480](https://github.com/gfx-rs/wgpu/pull/6480).
- Fix textureNumLevels in the GLSL backend. By @magcius in [#6483](https://github.com/gfx-rs/wgpu/pull/6483).
- Support 64-bit hex literals and unary operations in constants [#6616](https://github.com/gfx-rs/wgpu/pull/6616).

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

@@ -6057,89 +6057,3 @@ mod test {
let _ = super::parse_u8_slice(&bin, &Default::default()).unwrap();
}
}
#[cfg(all(test, feature = "wgsl-in", wgsl_out))]
mod test_atomic {
fn atomic_test(bytes: &[u8]) {
let _ = env_logger::builder().is_test(true).try_init();
let m = crate::front::spv::parse_u8_slice(bytes, &Default::default()).unwrap();
let mut wgsl = String::new();
for (vflags, name) in [
(crate::valid::ValidationFlags::empty(), "empty"),
(crate::valid::ValidationFlags::all(), "all"),
] {
log::info!("validating with flags - {name}");
let mut validator = crate::valid::Validator::new(vflags, Default::default());
match validator.validate(&m) {
Err(e) => {
log::error!("SPIR-V validation {}", e.emit_to_string(""));
log::info!("types: {:#?}", m.types);
panic!("validation error");
}
Ok(i) => {
wgsl = crate::back::wgsl::write_string(
&m,
&i,
crate::back::wgsl::WriterFlags::empty(),
)
.unwrap();
log::info!("wgsl-out:\n{wgsl}");
}
};
}
let m = match crate::front::wgsl::parse_str(&wgsl) {
Ok(m) => m,
Err(e) => {
log::error!("round trip WGSL validation {}", e.emit_to_string(&wgsl));
panic!("invalid module");
}
};
let mut validator =
crate::valid::Validator::new(crate::valid::ValidationFlags::all(), Default::default());
if let Err(e) = validator.validate(&m) {
log::error!("{}", e.emit_to_string(&wgsl));
panic!("invalid generated wgsl");
}
}
#[test]
fn atomic_i_inc() {
atomic_test(include_bytes!(
"../../../tests/in/spv/atomic_i_increment.spv"
));
}
#[test]
fn atomic_load_and_store() {
atomic_test(include_bytes!(
"../../../tests/in/spv/atomic_load_and_store.spv"
));
}
#[test]
fn atomic_exchange() {
atomic_test(include_bytes!("../../../tests/in/spv/atomic_exchange.spv"));
}
#[test]
fn atomic_compare_exchange() {
atomic_test(include_bytes!(
"../../../tests/in/spv/atomic_compare_exchange.spv"
));
}
#[test]
fn atomic_i_decrement() {
atomic_test(include_bytes!(
"../../../tests/in/spv/atomic_i_decrement.spv"
));
}
#[test]
fn atomic_i_add_and_sub() {
atomic_test(include_bytes!("../../../tests/in/spv/atomic_i_add_sub.spv"));
}
}

287
naga/tests/out/ir/atomic_i_increment.compact.ron
View File

@@ -1,287 +0,0 @@
(
types: [
(
name: None,
inner: Scalar((
kind: Uint,
width: 4,
)),
),
(
name: None,
inner: Scalar((
kind: Bool,
width: 1,
)),
),
(
name: None,
inner: Struct(
members: [
(
name: None,
ty: 0,
binding: None,
offset: 0,
),
],
span: 4,
),
),
(
name: None,
inner: Atomic((
kind: Uint,
width: 4,
)),
),
(
name: None,
inner: Struct(
members: [
(
name: None,
ty: 3,
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: 1,
init: 1,
),
(
name: None,
ty: 0,
init: 2,
),
(
name: None,
ty: 1,
init: 3,
),
(
name: None,
ty: 0,
init: 4,
),
],
overrides: [],
global_variables: [
(
name: None,
space: Storage(
access: ("LOAD | STORE"),
),
binding: Some((
group: 0,
binding: 0,
)),
ty: 4,
init: None,
),
(
name: None,
space: Storage(
access: ("LOAD"),
),
binding: Some((
group: 0,
binding: 1,
)),
ty: 2,
init: None,
),
],
global_expressions: [
Literal(U32(0)),
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(3),
Constant(1),
Constant(4),
Constant(2),
Constant(0),
AccessIndex(
base: 0,
index: 0,
),
AccessIndex(
base: 1,
index: 0,
),
LocalVariable(0),
Load(
pointer: 9,
),
Load(
pointer: 8,
),
Binary(
op: GreaterEqual,
left: 10,
right: 11,
),
AtomicResult(
ty: 0,
comparison: false,
),
Literal(U32(1)),
Binary(
op: Add,
left: 10,
right: 5,
),
LocalVariable(1),
Load(
pointer: 16,
),
Select(
condition: 12,
accept: 3,
reject: 2,
),
Unary(
op: LogicalNot,
expr: 18,
),
LocalVariable(0),
LocalVariable(1),
],
named_expressions: {},
body: [
Emit((
start: 7,
end: 9,
)),
Store(
pointer: 20,
value: 6,
),
Loop(
body: [
Emit((
start: 10,
end: 11,
)),
Emit((
start: 11,
end: 13,
)),
If(
condition: 12,
accept: [
Store(
pointer: 21,
value: 4,
),
],
reject: [
Atomic(
pointer: 7,
fun: Add,
value: 14,
result: Some(13),
),
Emit((
start: 15,
end: 16,
)),
Store(
pointer: 21,
value: 15,
),
],
),
Emit((
start: 17,
end: 19,
)),
Continue,
],
continuing: [
Emit((
start: 19,
end: 20,
)),
Store(
pointer: 20,
value: 17,
),
],
break_if: Some(19),
),
Return(
value: None,
),
],
diagnostic_filter_leaf: None,
),
],
entry_points: [
(
name: "stage::test_atomic_i_increment",
stage: Compute,
early_depth_test: None,
workgroup_size: (32, 1, 1),
workgroup_size_overrides: None,
function: (
name: Some("stage::test_atomic_i_increment_wrap"),
arguments: [],
result: None,
local_variables: [],
expressions: [],
named_expressions: {},
body: [
Call(
function: 0,
arguments: [],
result: None,
),
],
diagnostic_filter_leaf: None,
),
),
],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

312
naga/tests/out/ir/atomic_i_increment.ron
View File

@@ -1,312 +0,0 @@
(
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,
),
],
diagnostic_filter_leaf: None,
),
],
entry_points: [
(
name: "stage::test_atomic_i_increment",
stage: Compute,
early_depth_test: None,
workgroup_size: (32, 1, 1),
workgroup_size_overrides: None,
function: (
name: Some("stage::test_atomic_i_increment_wrap"),
arguments: [],
result: None,
local_variables: [],
expressions: [],
named_expressions: {},
body: [
Call(
function: 0,
arguments: [],
result: None,
),
],
diagnostic_filter_leaf: None,
),
),
],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

66
naga/tests/out/wgsl/atomic_compare_exchange.wgsl Normal file
View File

@@ -0,0 +1,66 @@
struct type_2 {
member: u32,
member_1: u32,
}
struct type_3 {
member: u32,
}
struct type_5 {
member: atomic<u32>,
}
@group(0) @binding(0)
var<storage, read_write> global: type_5;
@group(0) @binding(1)
var<storage> global_1: type_3;
fn function() {
var phi_33_: type_2;
var phi_34_: type_2;
var phi_49_: type_2;
var phi_63_: bool;
let _e11 = global_1.member;
phi_33_ = type_2(0u, _e11);
loop {
let _e14 = phi_33_;
if (_e14.member < _e14.member_1) {
phi_34_ = type_2((_e14.member + 1u), _e14.member_1);
phi_49_ = type_2(1u, _e14.member);
} else {
phi_34_ = _e14;
phi_49_ = type_2(0u, type_2().member_1);
}
let _e25 = phi_34_;
let _e27 = phi_49_;
switch bitcast<i32>(_e27.member) {
case 0: {
phi_63_ = false;
break;
}
case 1: {
let _e31 = atomicCompareExchangeWeak((&global.member), 3u, _e27.member_1);
phi_63_ = select(true, false, (_e31.old_value == 3u));
break;
}
default: {
phi_63_ = bool();
break;
}
}
let _e36 = phi_63_;
continue;
continuing {
phi_33_ = _e25;
break if !(_e36);
}
}
return;
}
@compute @workgroup_size(32, 1, 1)
fn stagetest_atomic_compare_exchange() {
function();
}

80
naga/tests/out/wgsl/atomic_exchange.wgsl Normal file
View File

@@ -0,0 +1,80 @@
struct type_2 {
member: u32,
member_1: u32,
}
struct type_3 {
member: u32,
}
struct type_5 {
member: atomic<u32>,
}
@group(0) @binding(0)
var<storage, read_write> global: type_5;
@group(0) @binding(1)
var<storage> global_1: type_3;
fn function() {
var phi_33_: type_2;
var phi_36_: u32;
var phi_52_: type_2;
var phi_53_: type_2;
var phi_62_: bool;
var phi_34_: type_2;
var phi_37_: u32;
let _e10 = global_1.member;
phi_33_ = type_2(0u, _e10);
phi_36_ = 0u;
loop {
let _e13 = phi_33_;
let _e15 = phi_36_;
if (_e13.member < _e13.member_1) {
phi_52_ = type_2((_e13.member + 1u), _e13.member_1);
phi_53_ = type_2(1u, _e13.member);
} else {
phi_52_ = _e13;
phi_53_ = type_2(0u, type_2().member_1);
}
let _e26 = phi_52_;
let _e28 = phi_53_;
switch bitcast<i32>(_e28.member) {
case 0: {
phi_62_ = false;
phi_34_ = type_2();
phi_37_ = u32();
break;
}
case 1: {
let _e31 = atomicExchange((&global.member), _e15);
phi_62_ = true;
phi_34_ = _e26;
phi_37_ = (_e15 + _e31);
break;
}
default: {
phi_62_ = false;
phi_34_ = type_2();
phi_37_ = u32();
break;
}
}
let _e34 = phi_62_;
let _e36 = phi_34_;
let _e38 = phi_37_;
continue;
continuing {
phi_33_ = _e36;
phi_36_ = _e38;
break if !(_e34);
}
}
return;
}
@compute @workgroup_size(32, 1, 1)
fn stagetest_atomic_exchange() {
function();
}

26
naga/tests/out/wgsl/atomic_i_add_sub.wgsl Normal file
View File

@@ -0,0 +1,26 @@
struct type_2 {
member: array<u32>,
}
struct type_4 {
member: atomic<u32>,
}
@group(0) @binding(0)
var<storage, read_write> global: type_4;
@group(0) @binding(1)
var<storage, read_write> global_1: type_2;
fn function() {
let _e6 = atomicAdd((&global.member), 2u);
let _e7 = atomicSub((&global.member), _e6);
if (_e6 < arrayLength((&global_1.member))) {
global_1.member[_e6] = _e7;
}
return;
}
@compute @workgroup_size(32, 1, 1)
fn stagetest_atomic_i_add_sub() {
function();
}

37
naga/tests/out/wgsl/atomic_i_decrement.wgsl Normal file
View File

@@ -0,0 +1,37 @@
struct type_3 {
member: array<u32>,
}
struct type_5 {
member: atomic<u32>,
}
@group(0) @binding(0)
var<storage, read_write> global: type_5;
@group(0) @binding(1)
var<storage, read_write> global_1: type_3;
fn function() {
var phi_40_: bool;
loop {
let _e8 = atomicSub((&global.member), 1u);
if (_e8 < arrayLength((&global_1.member))) {
global_1.member[_e8] = _e8;
phi_40_ = select(true, false, (_e8 == 0u));
} else {
phi_40_ = false;
}
let _e16 = phi_40_;
continue;
continuing {
break if !(_e16);
}
}
return;
}
@compute @workgroup_size(32, 1, 1)
fn stagetest_atomic_i_decrement() {
function();
}

42
naga/tests/out/wgsl/atomic_i_increment.wgsl Normal file
View File

@@ -0,0 +1,42 @@
struct type_2 {
member: u32,
}
struct type_4 {
member: atomic<u32>,
}
@group(0) @binding(0)
var<storage, read_write> global: type_4;
@group(0) @binding(1)
var<storage> global_1: type_2;
fn function() {
var phi_23_: u32;
var phi_24_: u32;
phi_23_ = 0u;
loop {
let _e10 = phi_23_;
let _e11 = global_1.member;
let _e12 = (_e10 >= _e11);
if _e12 {
phi_24_ = u32();
} else {
let _e13 = atomicAdd((&global.member), 1u);
phi_24_ = (_e10 + 1u);
}
let _e17 = phi_24_;
continue;
continuing {
phi_23_ = _e17;
break if !(select(true, false, _e12));
}
}
return;
}
@compute @workgroup_size(32, 1, 1)
fn stagetest_atomic_i_increment() {
function();
}

72
naga/tests/out/wgsl/atomic_load_and_store.wgsl Normal file
View File

@@ -0,0 +1,72 @@
struct type_2 {
member: u32,
member_1: u32,
}
struct type_3 {
member: u32,
}
struct type_5 {
member: atomic<u32>,
}
@group(0) @binding(0)
var<storage, read_write> global: type_5;
@group(0) @binding(1)
var<storage> global_1: type_3;
fn function() {
var phi_32_: type_2;
var phi_49_: type_2;
var phi_50_: type_2;
var phi_59_: bool;
var phi_33_: type_2;
let _e10 = global_1.member;
phi_32_ = type_2(0u, _e10);
loop {
let _e13 = phi_32_;
if (_e13.member < _e13.member_1) {
phi_49_ = type_2((_e13.member + 1u), _e13.member_1);
phi_50_ = type_2(1u, _e13.member);
} else {
phi_49_ = _e13;
phi_50_ = type_2(0u, type_2().member_1);
}
let _e24 = phi_49_;
let _e26 = phi_50_;
switch bitcast<i32>(_e26.member) {
case 0: {
phi_59_ = false;
phi_33_ = type_2();
break;
}
case 1: {
let _e29 = atomicLoad((&global.member));
atomicStore((&global.member), (_e29 + 2u));
phi_59_ = true;
phi_33_ = _e24;
break;
}
default: {
phi_59_ = false;
phi_33_ = type_2();
break;
}
}
let _e32 = phi_59_;
let _e34 = phi_33_;
continue;
continuing {
phi_32_ = _e34;
break if !(_e32);
}
}
return;
}
@compute @workgroup_size(32, 1, 1)
fn stagetest_atomic_load_and_store() {
function();
}

7
naga/tests/snapshots.rs
View File

@@ -1071,7 +1071,12 @@ fn convert_spv_all() {
false,
Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL,
);
convert_spv("atomic_i_increment", false, Targets::IR);
convert_spv("atomic_i_increment", false, Targets::WGSL);
convert_spv("atomic_load_and_store", false, Targets::WGSL);
convert_spv("atomic_exchange", false, Targets::WGSL);
convert_spv("atomic_compare_exchange", false, Targets::WGSL);
convert_spv("atomic_i_decrement", false, Targets::WGSL);
convert_spv("atomic_i_add_sub", false, Targets::WGSL);
convert_spv(
"fetch_depth",
false,