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

Fix hlsl storage format generation (#6993)

Browse Source
This commit is contained in:
Vecvec 2025-01-31 23:31:03 +08:00 committed by GitHub
parent 8caefceb8e
commit 7cde4707ec
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
15 changed files with 1303 additions and 6 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
CHANGELOG.md
View File

@ -85,6 +85,10 @@ By @brodycj in [#6924](https://github.com/gfx-rs/wgpu/pull/6924).
- Stop naga causing undefined behavior when a ray query misses. By @Vecvec in [#6752](https://github.com/gfx-rs/wgpu/pull/6752).
#### Dx12
- Fix HLSL storage format generation. By @Vecvec in [#6993](https://github.com/gfx-rs/wgpu/pull/6993)
#### WebGPU
- Improve efficiency of dropping read-only buffer mappings. By @kpreid in [#7007](https://github.com/gfx-rs/wgpu/pull/7007).

12
naga/src/back/hlsl/conv.rs
View File

@ -127,14 +127,14 @@ impl crate::StorageFormat {
Self::R8Sint | Self::R16Sint | Self::R32Sint => "int",
Self::R64Uint => "uint64_t",
Self::Rg16Float | Self::Rg32Float => "float2",
Self::Rg8Unorm | Self::Rg16Unorm => "unorm float2",
Self::Rg8Snorm | Self::Rg16Snorm => "snorm float2",
Self::Rg16Float | Self::Rg32Float => "float4",
Self::Rg8Unorm | Self::Rg16Unorm => "unorm float4",
Self::Rg8Snorm | Self::Rg16Snorm => "snorm float4",
Self::Rg8Sint | Self::Rg16Sint | Self::Rg32Uint => "int2",
Self::Rg8Uint | Self::Rg16Uint | Self::Rg32Sint => "uint2",
Self::Rg8Sint | Self::Rg16Sint | Self::Rg32Uint => "int4",
Self::Rg8Uint | Self::Rg16Uint | Self::Rg32Sint => "uint4",
Self::Rg11b10Ufloat => "float3",
Self::Rg11b10Ufloat => "float4",
Self::Rgba16Float | Self::Rgba32Float => "float4",
Self::Rgba8Unorm | Self::Bgra8Unorm | Self::Rgba16Unorm | Self::Rgb10a2Unorm => {

17
naga/tests/in/storage-textures.wgsl Normal file
View File

@ -0,0 +1,17 @@
@group(0) @binding(0) var s_r_r: texture_storage_2d<r32float, read>;
@group(0) @binding(1) var s_rg_r: texture_storage_2d<rg32float, read>;
@group(0) @binding(2) var s_rgba_r: texture_storage_2d<rgba32float, read>;
@compute @workgroup_size(1) fn csLoad() {
_ = textureLoad(s_r_r, vec2u(0));
_ = textureLoad(s_rg_r, vec2u(0));
_ = textureLoad(s_rgba_r, vec2u(0));
}
@group(1) @binding(0) var s_r_w: texture_storage_2d<r32float, write>;
@group(1) @binding(1) var s_rg_w: texture_storage_2d<rg32float, write>;
@group(1) @binding(2) var s_rgba_w: texture_storage_2d<rgba32float, write>;
@compute @workgroup_size(1) fn csStore() {
textureStore(s_r_w, vec2u(0), vec4f(0.0));
textureStore(s_rg_w, vec2u(0), vec4f(0.0));
textureStore(s_rgba_w, vec2u(0), vec4f(0.0));
}

402
naga/tests/out/analysis/storage-textures.info.ron Normal file
View File

@ -0,0 +1,402 @@
(
type_flags: [
("CREATION_RESOLVED | ARGUMENT"),
("CREATION_RESOLVED | ARGUMENT"),
("CREATION_RESOLVED | ARGUMENT"),
("CREATION_RESOLVED | ARGUMENT"),
("CREATION_RESOLVED | ARGUMENT"),
("CREATION_RESOLVED | ARGUMENT"),
],
functions: [],
entry_points: [
(
flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
uniformity: (
non_uniform_result: None,
requirements: (""),
),
may_kill: false,
sampling_set: [],
global_uses: [
("READ"),
("READ"),
("READ"),
(""),
(""),
(""),
],
expressions: [
(
uniformity: (
non_uniform_result: Some(0),
requirements: (""),
),
ref_count: 1,
assignable_global: Some(0),
ty: Handle(0),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Scalar((
kind: Uint,
width: 4,
))),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Vector(
size: Bi,
scalar: (
kind: Uint,
width: 4,
),
)),
),
(
uniformity: (
non_uniform_result: Some(0),
requirements: (""),
),
ref_count: 0,
assignable_global: None,
ty: Value(Vector(
size: Quad,
scalar: (
kind: Float,
width: 4,
),
)),
),
(
uniformity: (
non_uniform_result: Some(4),
requirements: (""),
),
ref_count: 1,
assignable_global: Some(1),
ty: Handle(1),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Scalar((
kind: Uint,
width: 4,
))),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Vector(
size: Bi,
scalar: (
kind: Uint,
width: 4,
),
)),
),
(
uniformity: (
non_uniform_result: Some(4),
requirements: (""),
),
ref_count: 0,
assignable_global: None,
ty: Value(Vector(
size: Quad,
scalar: (
kind: Float,
width: 4,
),
)),
),
(
uniformity: (
non_uniform_result: Some(8),
requirements: (""),
),
ref_count: 1,
assignable_global: Some(2),
ty: Handle(2),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Scalar((
kind: Uint,
width: 4,
))),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Vector(
size: Bi,
scalar: (
kind: Uint,
width: 4,
),
)),
),
(
uniformity: (
non_uniform_result: Some(8),
requirements: (""),
),
ref_count: 0,
assignable_global: None,
ty: Value(Vector(
size: Quad,
scalar: (
kind: Float,
width: 4,
),
)),
),
],
sampling: [],
dual_source_blending: false,
diagnostic_filter_leaf: None,
),
(
flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
uniformity: (
non_uniform_result: None,
requirements: (""),
),
may_kill: false,
sampling_set: [],
global_uses: [
(""),
(""),
(""),
("WRITE"),
("WRITE"),
("WRITE"),
],
expressions: [
(
uniformity: (
non_uniform_result: Some(0),
requirements: (""),
),
ref_count: 1,
assignable_global: Some(3),
ty: Handle(3),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Scalar((
kind: Uint,
width: 4,
))),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Vector(
size: Bi,
scalar: (
kind: Uint,
width: 4,
),
)),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Scalar((
kind: Float,
width: 4,
))),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Vector(
size: Quad,
scalar: (
kind: Float,
width: 4,
),
)),
),
(
uniformity: (
non_uniform_result: Some(5),
requirements: (""),
),
ref_count: 1,
assignable_global: Some(4),
ty: Handle(4),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Scalar((
kind: Uint,
width: 4,
))),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Vector(
size: Bi,
scalar: (
kind: Uint,
width: 4,
),
)),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Scalar((
kind: Float,
width: 4,
))),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Vector(
size: Quad,
scalar: (
kind: Float,
width: 4,
),
)),
),
(
uniformity: (
non_uniform_result: Some(10),
requirements: (""),
),
ref_count: 1,
assignable_global: Some(5),
ty: Handle(5),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Scalar((
kind: Uint,
width: 4,
))),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Vector(
size: Bi,
scalar: (
kind: Uint,
width: 4,
),
)),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Scalar((
kind: Float,
width: 4,
))),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Vector(
size: Quad,
scalar: (
kind: Float,
width: 4,
),
)),
),
],
sampling: [],
dual_source_blending: false,
diagnostic_filter_leaf: None,
),
],
const_expression_types: [],
)

24
naga/tests/out/hlsl/storage-textures.hlsl Normal file
View File

@ -0,0 +1,24 @@
RWTexture2D<float> s_r_r : register(u0);
RWTexture2D<float4> s_rg_r : register(u1);
RWTexture2D<float4> s_rgba_r : register(u2);
RWTexture2D<float> s_r_w : register(u0, space1);
RWTexture2D<float4> s_rg_w : register(u1, space1);
RWTexture2D<float4> s_rgba_w : register(u2, space1);
[numthreads(1, 1, 1)]
void csLoad()
{
float4 phony = s_r_r.Load((0u).xx);
float4 phony_1 = s_rg_r.Load((0u).xx);
float4 phony_2 = s_rgba_r.Load((0u).xx);
return;
}
[numthreads(1, 1, 1)]
void csStore()
{
s_r_w[(0u).xx] = (0.0).xxxx;
s_rg_w[(0u).xx] = (0.0).xxxx;
s_rgba_w[(0u).xx] = (0.0).xxxx;
return;
}

16
naga/tests/out/hlsl/storage-textures.ron Normal file
View File

@ -0,0 +1,16 @@
(
vertex:[
],
fragment:[
],
compute:[
(
entry_point:"csLoad",
target_profile:"cs_5_1",
),
(
entry_point:"csStore",
target_profile:"cs_5_1",
),
],
)

319
naga/tests/out/ir/storage-textures.compact.ron Normal file
View File

@ -0,0 +1,319 @@
(
types: [
(
name: None,
inner: Image(
dim: D2,
arrayed: false,
class: Storage(
format: R32Float,
access: ("LOAD"),
),
),
),
(
name: None,
inner: Image(
dim: D2,
arrayed: false,
class: Storage(
format: Rg32Float,
access: ("LOAD"),
),
),
),
(
name: None,
inner: Image(
dim: D2,
arrayed: false,
class: Storage(
format: Rgba32Float,
access: ("LOAD"),
),
),
),
(
name: None,
inner: Image(
dim: D2,
arrayed: false,
class: Storage(
format: R32Float,
access: ("STORE"),
),
),
),
(
name: None,
inner: Image(
dim: D2,
arrayed: false,
class: Storage(
format: Rg32Float,
access: ("STORE"),
),
),
),
(
name: None,
inner: Image(
dim: D2,
arrayed: false,
class: Storage(
format: Rgba32Float,
access: ("STORE"),
),
),
),
],
special_types: (
ray_desc: None,
ray_intersection: None,
predeclared_types: {},
),
constants: [],
overrides: [],
global_variables: [
(
name: Some("s_r_r"),
space: Handle,
binding: Some((
group: 0,
binding: 0,
)),
ty: 0,
init: None,
),
(
name: Some("s_rg_r"),
space: Handle,
binding: Some((
group: 0,
binding: 1,
)),
ty: 1,
init: None,
),
(
name: Some("s_rgba_r"),
space: Handle,
binding: Some((
group: 0,
binding: 2,
)),
ty: 2,
init: None,
),
(
name: Some("s_r_w"),
space: Handle,
binding: Some((
group: 1,
binding: 0,
)),
ty: 3,
init: None,
),
(
name: Some("s_rg_w"),
space: Handle,
binding: Some((
group: 1,
binding: 1,
)),
ty: 4,
init: None,
),
(
name: Some("s_rgba_w"),
space: Handle,
binding: Some((
group: 1,
binding: 2,
)),
ty: 5,
init: None,
),
],
global_expressions: [],
functions: [],
entry_points: [
(
name: "csLoad",
stage: Compute,
early_depth_test: None,
workgroup_size: (1, 1, 1),
workgroup_size_overrides: None,
function: (
name: Some("csLoad"),
arguments: [],
result: None,
local_variables: [],
expressions: [
GlobalVariable(0),
Literal(U32(0)),
Splat(
size: Bi,
value: 1,
),
ImageLoad(
image: 0,
coordinate: 2,
array_index: None,
sample: None,
level: None,
),
GlobalVariable(1),
Literal(U32(0)),
Splat(
size: Bi,
value: 5,
),
ImageLoad(
image: 4,
coordinate: 6,
array_index: None,
sample: None,
level: None,
),
GlobalVariable(2),
Literal(U32(0)),
Splat(
size: Bi,
value: 9,
),
ImageLoad(
image: 8,
coordinate: 10,
array_index: None,
sample: None,
level: None,
),
],
named_expressions: {
3: "phony",
7: "phony",
11: "phony",
},
body: [
Emit((
start: 2,
end: 4,
)),
Emit((
start: 6,
end: 8,
)),
Emit((
start: 10,
end: 12,
)),
Return(
value: None,
),
],
diagnostic_filter_leaf: None,
),
),
(
name: "csStore",
stage: Compute,
early_depth_test: None,
workgroup_size: (1, 1, 1),
workgroup_size_overrides: None,
function: (
name: Some("csStore"),
arguments: [],
result: None,
local_variables: [],
expressions: [
GlobalVariable(3),
Literal(U32(0)),
Splat(
size: Bi,
value: 1,
),
Literal(F32(0.0)),
Splat(
size: Quad,
value: 3,
),
GlobalVariable(4),
Literal(U32(0)),
Splat(
size: Bi,
value: 6,
),
Literal(F32(0.0)),
Splat(
size: Quad,
value: 8,
),
GlobalVariable(5),
Literal(U32(0)),
Splat(
size: Bi,
value: 11,
),
Literal(F32(0.0)),
Splat(
size: Quad,
value: 13,
),
],
named_expressions: {},
body: [
Emit((
start: 2,
end: 3,
)),
Emit((
start: 4,
end: 5,
)),
ImageStore(
image: 0,
coordinate: 2,
array_index: None,
value: 4,
),
Emit((
start: 7,
end: 8,
)),
Emit((
start: 9,
end: 10,
)),
ImageStore(
image: 5,
coordinate: 7,
array_index: None,
value: 9,
),
Emit((
start: 12,
end: 13,
)),
Emit((
start: 14,
end: 15,
)),
ImageStore(
image: 10,
coordinate: 12,
array_index: None,
value: 14,
),
Return(
value: None,
),
],
diagnostic_filter_leaf: None,
),
),
],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

319
naga/tests/out/ir/storage-textures.ron Normal file
View File

@ -0,0 +1,319 @@
(
types: [
(
name: None,
inner: Image(
dim: D2,
arrayed: false,
class: Storage(
format: R32Float,
access: ("LOAD"),
),
),
),
(
name: None,
inner: Image(
dim: D2,
arrayed: false,
class: Storage(
format: Rg32Float,
access: ("LOAD"),
),
),
),
(
name: None,
inner: Image(
dim: D2,
arrayed: false,
class: Storage(
format: Rgba32Float,
access: ("LOAD"),
),
),
),
(
name: None,
inner: Image(
dim: D2,
arrayed: false,
class: Storage(
format: R32Float,
access: ("STORE"),
),
),
),
(
name: None,
inner: Image(
dim: D2,
arrayed: false,
class: Storage(
format: Rg32Float,
access: ("STORE"),
),
),
),
(
name: None,
inner: Image(
dim: D2,
arrayed: false,
class: Storage(
format: Rgba32Float,
access: ("STORE"),
),
),
),
],
special_types: (
ray_desc: None,
ray_intersection: None,
predeclared_types: {},
),
constants: [],
overrides: [],
global_variables: [
(
name: Some("s_r_r"),
space: Handle,
binding: Some((
group: 0,
binding: 0,
)),
ty: 0,
init: None,
),
(
name: Some("s_rg_r"),
space: Handle,
binding: Some((
group: 0,
binding: 1,
)),
ty: 1,
init: None,
),
(
name: Some("s_rgba_r"),
space: Handle,
binding: Some((
group: 0,
binding: 2,
)),
ty: 2,
init: None,
),
(
name: Some("s_r_w"),
space: Handle,
binding: Some((
group: 1,
binding: 0,
)),
ty: 3,
init: None,
),
(
name: Some("s_rg_w"),
space: Handle,
binding: Some((
group: 1,
binding: 1,
)),
ty: 4,
init: None,
),
(
name: Some("s_rgba_w"),
space: Handle,
binding: Some((
group: 1,
binding: 2,
)),
ty: 5,
init: None,
),
],
global_expressions: [],
functions: [],
entry_points: [
(
name: "csLoad",
stage: Compute,
early_depth_test: None,
workgroup_size: (1, 1, 1),
workgroup_size_overrides: None,
function: (
name: Some("csLoad"),
arguments: [],
result: None,
local_variables: [],
expressions: [
GlobalVariable(0),
Literal(U32(0)),
Splat(
size: Bi,
value: 1,
),
ImageLoad(
image: 0,
coordinate: 2,
array_index: None,
sample: None,
level: None,
),
GlobalVariable(1),
Literal(U32(0)),
Splat(
size: Bi,
value: 5,
),
ImageLoad(
image: 4,
coordinate: 6,
array_index: None,
sample: None,
level: None,
),
GlobalVariable(2),
Literal(U32(0)),
Splat(
size: Bi,
value: 9,
),
ImageLoad(
image: 8,
coordinate: 10,
array_index: None,
sample: None,
level: None,
),
],
named_expressions: {
3: "phony",
7: "phony",
11: "phony",
},
body: [
Emit((
start: 2,
end: 4,
)),
Emit((
start: 6,
end: 8,
)),
Emit((
start: 10,
end: 12,
)),
Return(
value: None,
),
],
diagnostic_filter_leaf: None,
),
),
(
name: "csStore",
stage: Compute,
early_depth_test: None,
workgroup_size: (1, 1, 1),
workgroup_size_overrides: None,
function: (
name: Some("csStore"),
arguments: [],
result: None,
local_variables: [],
expressions: [
GlobalVariable(3),
Literal(U32(0)),
Splat(
size: Bi,
value: 1,
),
Literal(F32(0.0)),
Splat(
size: Quad,
value: 3,
),
GlobalVariable(4),
Literal(U32(0)),
Splat(
size: Bi,
value: 6,
),
Literal(F32(0.0)),
Splat(
size: Quad,
value: 8,
),
GlobalVariable(5),
Literal(U32(0)),
Splat(
size: Bi,
value: 11,
),
Literal(F32(0.0)),
Splat(
size: Quad,
value: 13,
),
],
named_expressions: {},
body: [
Emit((
start: 2,
end: 3,
)),
Emit((
start: 4,
end: 5,
)),
ImageStore(
image: 0,
coordinate: 2,
array_index: None,
value: 4,
),
Emit((
start: 7,
end: 8,
)),
Emit((
start: 9,
end: 10,
)),
ImageStore(
image: 5,
coordinate: 7,
array_index: None,
value: 9,
),
Emit((
start: 12,
end: 13,
)),
Emit((
start: 14,
end: 15,
)),
ImageStore(
image: 10,
coordinate: 12,
array_index: None,
value: 14,
),
Return(
value: None,
),
],
diagnostic_filter_leaf: None,
),
),
],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

29
naga/tests/out/msl/storage-textures.msl Normal file
View File

@ -0,0 +1,29 @@
// language: metal1.0
#include <metal_stdlib>
#include <simd/simd.h>
using metal::uint;
kernel void csLoad(
metal::texture2d<float, metal::access::read> s_r_r [[user(fake0)]]
, metal::texture2d<float, metal::access::read> s_rg_r [[user(fake0)]]
, metal::texture2d<float, metal::access::read> s_rgba_r [[user(fake0)]]
) {
metal::float4 phony = s_r_r.read(metal::uint2(metal::uint2(0u)));
metal::float4 phony_1 = s_rg_r.read(metal::uint2(metal::uint2(0u)));
metal::float4 phony_2 = s_rgba_r.read(metal::uint2(metal::uint2(0u)));
return;
}
kernel void csStore(
metal::texture2d<float, metal::access::write> s_r_w [[user(fake0)]]
, metal::texture2d<float, metal::access::write> s_rg_w [[user(fake0)]]
, metal::texture2d<float, metal::access::write> s_rgba_w [[user(fake0)]]
) {
s_r_w.write(metal::float4(0.0), metal::uint2(metal::uint2(0u)));
s_rg_w.write(metal::float4(0.0), metal::uint2(metal::uint2(0u)));
s_rgba_w.write(metal::float4(0.0), metal::uint2(metal::uint2(0u)));
return;
}

79
naga/tests/out/spv/storage-textures.spvasm Normal file
View File

@ -0,0 +1,79 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 42
OpCapability Shader
OpCapability StorageImageExtendedFormats
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %20 "csLoad"
OpEntryPoint GLCompute %35 "csStore"
OpExecutionMode %20 LocalSize 1 1 1
OpExecutionMode %35 LocalSize 1 1 1
OpDecorate %7 NonWritable
OpDecorate %7 DescriptorSet 0
OpDecorate %7 Binding 0
OpDecorate %9 NonWritable
OpDecorate %9 DescriptorSet 0
OpDecorate %9 Binding 1
OpDecorate %11 NonWritable
OpDecorate %11 DescriptorSet 0
OpDecorate %11 Binding 2
OpDecorate %13 NonReadable
OpDecorate %13 DescriptorSet 1
OpDecorate %13 Binding 0
OpDecorate %15 NonReadable
OpDecorate %15 DescriptorSet 1
OpDecorate %15 Binding 1
OpDecorate %17 NonReadable
OpDecorate %17 DescriptorSet 1
OpDecorate %17 Binding 2
%2 = OpTypeVoid
%4 = OpTypeFloat 32
%3 = OpTypeImage %4 2D 0 0 0 2 R32f
%5 = OpTypeImage %4 2D 0 0 0 2 Rg32f
%6 = OpTypeImage %4 2D 0 0 0 2 Rgba32f
%8 = OpTypePointer UniformConstant %3
%7 = OpVariable %8 UniformConstant
%10 = OpTypePointer UniformConstant %5
%9 = OpVariable %10 UniformConstant
%12 = OpTypePointer UniformConstant %6
%11 = OpVariable %12 UniformConstant
%14 = OpTypePointer UniformConstant %3
%13 = OpVariable %14 UniformConstant
%16 = OpTypePointer UniformConstant %5
%15 = OpVariable %16 UniformConstant
%18 = OpTypePointer UniformConstant %6
%17 = OpVariable %18 UniformConstant
%21 = OpTypeFunction %2
%25 = OpTypeInt 32 0
%26 = OpConstant %25 0
%27 = OpTypeVector %25 2
%28 = OpConstantComposite %27 %26 %26
%30 = OpTypeVector %4 4
%39 = OpConstant %4 0.0
%40 = OpConstantComposite %30 %39 %39 %39 %39
%20 = OpFunction %2 None %21
%19 = OpLabel
%22 = OpLoad %3 %7
%23 = OpLoad %5 %9
%24 = OpLoad %6 %11
OpBranch %29
%29 = OpLabel
%31 = OpImageRead %30 %22 %28
%32 = OpImageRead %30 %23 %28
%33 = OpImageRead %30 %24 %28
OpReturn
OpFunctionEnd
%35 = OpFunction %2 None %21
%34 = OpLabel
%36 = OpLoad %3 %13
%37 = OpLoad %5 %15
%38 = OpLoad %6 %17
OpBranch %41
%41 = OpLabel
OpImageWrite %36 %28 %40
OpImageWrite %37 %28 %40
OpImageWrite %38 %28 %40
OpReturn
OpFunctionEnd

13
naga/tests/out/wgsl/storage-textures.wgsl Normal file
View File

@ -0,0 +1,13 @@
@group(0) @binding(0)
var s_r: texture_storage_2d<r32float,read>;
@group(0) @binding(1)
var s_rg: texture_storage_2d<rg32float,read>;
@group(0) @binding(2)
var s_rgba: texture_storage_2d<rgba32float,read>;
@compute @workgroup_size(1, 1, 1)
fn csWithStorageUsage() {
let phony = textureLoad(s_r, vec2(0u));
let phony_1 = textureLoad(s_rg, vec2(0u));
let phony_2 = textureLoad(s_rgba, vec2(0u));
}

4
naga/tests/snapshots.rs
View File

@ -967,6 +967,10 @@ fn convert_wgsl() {
Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL,
),
("must-use", Targets::IR),
(
"storage-textures",
Targets::IR | Targets::ANALYSIS | Targets::SPIRV | Targets::METAL | Targets::HLSL,
),
];
for &(name, targets) in inputs.iter() {

1
tests/tests/root.rs
View File

@ -50,6 +50,7 @@ mod shader;
mod shader_primitive_index;
mod shader_view_format;
mod subgroup_operations;
mod texture_binding;
mod texture_blit;
mod texture_bounds;
mod texture_view_creation;

64
tests/tests/texture_binding/mod.rs Normal file
View File

@ -0,0 +1,64 @@
use wgpu::{
include_wgsl, BindGroupDescriptor, BindGroupEntry, BindingResource, ComputePassDescriptor,
ComputePipelineDescriptor, DownlevelFlags, Extent3d, Features, TextureDescriptor,
TextureDimension, TextureFormat, TextureUsages,
};
use wgpu_macros::gpu_test;
use wgpu_test::{GpuTestConfiguration, TestParameters, TestingContext};
#[gpu_test]
static TEXTURE_BINDING: GpuTestConfiguration = GpuTestConfiguration::new()
.parameters(
TestParameters::default()
.test_features_limits()
.downlevel_flags(DownlevelFlags::WEBGPU_TEXTURE_FORMAT_SUPPORT)
.features(Features::TEXTURE_ADAPTER_SPECIFIC_FORMAT_FEATURES),
)
.run_sync(texture_binding);
fn texture_binding(ctx: TestingContext) {
let texture = ctx.device.create_texture(&TextureDescriptor {
label: None,
size: Extent3d {
width: 1,
height: 1,
depth_or_array_layers: 1,
},
mip_level_count: 1,
sample_count: 1,
dimension: TextureDimension::D2,
format: TextureFormat::Rg32Float,
usage: TextureUsages::STORAGE_BINDING,
view_formats: &[],
});
let shader = ctx
.device
.create_shader_module(include_wgsl!("shader.wgsl"));
let pipeline = ctx
.device
.create_compute_pipeline(&ComputePipelineDescriptor {
label: None,
layout: None,
module: &shader,
entry_point: None,
compilation_options: Default::default(),
cache: None,
});
let bind = ctx.device.create_bind_group(&BindGroupDescriptor {
label: None,
layout: &pipeline.get_bind_group_layout(0),
entries: &[BindGroupEntry {
binding: 0,
resource: BindingResource::TextureView(&texture.create_view(&Default::default())),
}],
});
let mut encoder = ctx.device.create_command_encoder(&Default::default());
{
let mut pass = encoder.begin_compute_pass(&ComputePassDescriptor::default());
pass.set_pipeline(&pipeline);
pass.set_bind_group(0, &bind, &[]);
pass.dispatch_workgroups(1, 1, 1);
}
ctx.queue.submit([encoder.finish()]);
}

6
tests/tests/texture_binding/shader.wgsl Normal file
View File

@ -0,0 +1,6 @@
@group(0) @binding(0)
var tex: texture_storage_2d<rg32float, read>;
@compute @workgroup_size(1) fn csStore() {
_ = textureLoad(tex, vec2u(0));
}