Skip to content

Commit

Permalink
Fix hlsl storage format generation (#6993)
Browse files Browse the repository at this point in the history
@Vecvec
Vecvec authored Jan 31, 2025
1 parent 8caefce commit 7cde470
Showing 15 changed files with 1,303 additions and 6 deletions.
4 changes: 4 additions & 0 deletions CHANGELOG.md
View file
Original file line number Diff line number Diff line change
@@ -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 changes: 6 additions & 6 deletions naga/src/back/hlsl/conv.rs
View file
Original file line number Diff line number Diff line change
@@ -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 changes: 17 additions & 0 deletions naga/tests/in/storage-textures.wgsl
View file
Original file line number Diff line number Diff line change
@@ -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 changes: 402 additions & 0 deletions naga/tests/out/analysis/storage-textures.info.ron
View file

Large diffs are not rendered by default.

24 changes: 24 additions & 0 deletions naga/tests/out/hlsl/storage-textures.hlsl
View file
Original file line number Diff line number Diff line change
@@ -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 changes: 16 additions & 0 deletions naga/tests/out/hlsl/storage-textures.ron
View file
Original file line number Diff line number Diff line change
@@ -0,0 +1,16 @@
(
vertex:[
],
fragment:[
],
compute:[
(
entry_point:"csLoad",
target_profile:"cs_5_1",
),
(
entry_point:"csStore",
target_profile:"cs_5_1",
),
],
)
319 changes: 319 additions & 0 deletions naga/tests/out/ir/storage-textures.compact.ron
View file
Original file line number Diff line number Diff line change
@@ -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 changes: 319 additions & 0 deletions naga/tests/out/ir/storage-textures.ron
View file
Original file line number Diff line number Diff line change
@@ -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 changes: 29 additions & 0 deletions naga/tests/out/msl/storage-textures.msl
View file
Original file line number Diff line number Diff line change
@@ -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 changes: 79 additions & 0 deletions naga/tests/out/spv/storage-textures.spvasm
View file
Original file line number Diff line number Diff line change
@@ -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 changes: 13 additions & 0 deletions naga/tests/out/wgsl/storage-textures.wgsl
View file
Original file line number Diff line number Diff line change
@@ -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 changes: 4 additions & 0 deletions naga/tests/snapshots.rs
View file
Original file line number Diff line number Diff line change
@@ -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 change: 1 addition & 0 deletions tests/tests/root.rs
View file
Original file line number Diff line number Diff line change
@@ -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 changes: 64 additions & 0 deletions tests/tests/texture_binding/mod.rs
View file
Original file line number Diff line number Diff line change
@@ -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 changes: 6 additions & 0 deletions tests/tests/texture_binding/shader.wgsl
View file
Original file line number Diff line number Diff line change
@@ -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));
}

0 comments on commit 7cde470

Please sign in to comment.