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

[glsl-out] improve array types

Browse Source
This commit is contained in:
Dzmitry Malyshau
2021-05-28 22:11:39 -04:00
committed by Dzmitry Malyshau
parent d1e2ac57fe
commit 1366546645
10 changed files with 66 additions and 57 deletions
Download Patch File Download Diff File Expand all files Collapse all files

69
src/back/glsl/mod.rs
View File

@@ -564,6 +564,27 @@ impl<'a, W: Write> Writer<'a, W> {
self.collect_reflection_info()
}
fn write_array_size(&mut self, size: ArraySize) -> BackendResult {
write!(self.out, "[")?;
// Write the array size
// Writes nothing if `ArraySize::Dynamic`
// Panics if `ArraySize::Constant` has a constant that isn't an uint
match size {
ArraySize::Constant(const_handle) => match self.module.constants[const_handle].inner {
ConstantInner::Scalar {
width: _,
value: ScalarValue::Uint(size),
} => write!(self.out, "{}", size)?,
_ => unreachable!(),
},
ArraySize::Dynamic => (),
}
write!(self.out, "]")?;
Ok(())
}
/// Helper method used to write value types
///
/// # Notes
@@ -615,27 +636,7 @@ impl<'a, W: Write> Writer<'a, W> {
// GLSL arrays are written as `type name[size]`
// Current code is written arrays only as `[size]`
// Base `type` and `name` should be written outside
TypeInner::Array { base: _, size, .. } => {
write!(self.out, "[")?;
// Write the array size
// Writes nothing if `ArraySize::Dynamic`
// Panics if `ArraySize::Constant` has a constant that isn't an uint
match size {
ArraySize::Constant(const_handle) => {
match self.module.constants[const_handle].inner {
ConstantInner::Scalar {
width: _,
value: ScalarValue::Uint(size),
} => write!(self.out, "{}", size)?,
_ => unreachable!(),
}
}
ArraySize::Dynamic => (),
}
write!(self.out, "]")?
}
TypeInner::Array { size, .. } => self.write_array_size(size)?,
// Panic if either Image, Sampler, Pointer, or a Struct is being written
//
// Write all variants instead of `_` so that if new variants are added a
@@ -674,6 +675,8 @@ impl<'a, W: Write> Writer<'a, W> {
write!(self.out, "{}", name)?;
Ok(())
}
// glsl array has the size separated from the base type
TypeInner::Array { base, .. } => self.write_type(base),
ref other => self.write_value_type(other),
}
}
@@ -756,14 +759,14 @@ impl<'a, W: Write> Writer<'a, W> {
// Finally write the global name and end the global with a `;` and a newline
// Leading space is important
let global_name = self.get_global_name(handle, global);
let global_str =
if let Some(default_value) = zero_init_value_str(&self.module.types[global.ty].inner) {
format!("{} = {}", global_name, default_value)
} else {
global_name
};
writeln!(self.out, " {};", global_str)?;
writeln!(self.out)?;
write!(self.out, " {}", global_name)?;
if let TypeInner::Array { size, .. } = self.module.types[global.ty].inner {
self.write_array_size(size)?;
}
if let Some(default_value) = zero_init_value_str(&self.module.types[global.ty].inner) {
write!(self.out, " = {}", default_value)?;
};
writeln!(self.out, ";")?;
Ok(())
}
@@ -1117,7 +1120,11 @@ impl<'a, W: Write> Writer<'a, W> {
write!(self.out, "{}", INDENT)?;
match self.module.types[member.ty].inner {
TypeInner::Array { base, .. } => {
TypeInner::Array {
base,
size,
stride: _,
} => {
// GLSL arrays are written as `type name[size]`
let ty_name = match self.module.types[base].inner {
// Write scalar type by backend so as not to depend on the front-end implementation
@@ -1134,7 +1141,7 @@ impl<'a, W: Write> Writer<'a, W> {
&self.names[&NameKey::StructMember(handle, idx as u32)]
)?;
// Write [size]
self.write_type(member.ty)?;
self.write_array_size(size)?;
// Newline is important
writeln!(self.out, ";")?;
}

3
tests/in/globals.wgsl
View File

@@ -2,8 +2,9 @@
let Foo: bool = true;
var<workgroup> wg : array<f32, 10>;
var<workgroup> wg : array<f32, 10u>;
[[stage(compute), workgroup_size(1)]]
fn main() {
wg[3] = 1.0;
}

3
tests/out/boids.Compute.glsl
View File

@@ -18,16 +18,13 @@ uniform SimParams_block_0 {
float rule2Scale;
float rule3Scale;
} _group_0_binding_0;
readonly buffer Particles_block_1 {
Particle particles[];
} _group_0_binding_1;
buffer Particles_block_2 {
Particle particles[];
} _group_0_binding_2;
void main() {
uvec3 global_invocation_id = gl_GlobalInvocationID;
vec2 vPos;

2
tests/out/globals.Compute.glsl
View File

@@ -4,8 +4,10 @@ precision highp float;
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
shared float wg[10];
void main() {
wg[3] = 1.0;
return;
}

4
tests/out/globals.msl
View File

@@ -3,10 +3,12 @@
constexpr constant bool Foo = true;
struct type2 {
float inner[10];
float inner[10u];
};
kernel void main1(
threadgroup type2& wg
) {
wg.inner[3] = 1.0;
return;
}

32
tests/out/globals.spvasm
View File

@@ -1,26 +1,32 @@
; SPIR-V
; Version: 1.0
; Generator: rspirv
; Bound: 15
; Bound: 20
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %12 "main"
OpExecutionMode %12 LocalSize 1 1 1
OpDecorate %8 ArrayStride 4
OpEntryPoint GLCompute %15 "main"
OpExecutionMode %15 LocalSize 1 1 1
OpDecorate %11 ArrayStride 4
%2 = OpTypeVoid
%4 = OpTypeBool
%3 = OpConstantTrue %4
%6 = OpTypeInt 32 1
%6 = OpTypeInt 32 0
%5 = OpConstant %6 10
%7 = OpTypeFloat 32
%8 = OpTypeArray %7 %5
%10 = OpTypePointer Workgroup %8
%9 = OpVariable %10 Workgroup
%13 = OpTypeFunction %2
%12 = OpFunction %2 None %13
%11 = OpLabel
OpBranch %14
%8 = OpTypeInt 32 1
%7 = OpConstant %8 3
%10 = OpTypeFloat 32
%9 = OpConstant %10 1.0
%11 = OpTypeArray %10 %5
%13 = OpTypePointer Workgroup %11
%12 = OpVariable %13 Workgroup
%16 = OpTypeFunction %2
%18 = OpTypePointer Workgroup %10
%15 = OpFunction %2 None %16
%14 = OpLabel
OpBranch %17
%17 = OpLabel
%19 = OpAccessChain %18 %12 %7
OpStore %19 %9
OpReturn
OpFunctionEnd

3
tests/out/globals.wgsl
View File

@@ -1,8 +1,9 @@
let Foo: bool = true;
var<workgroup> wg: array<f32,10>;
var<workgroup> wg: array<f32,10u>;
[[stage(compute), workgroup_size(1, 1, 1)]]
fn main() {
wg[3] = 1.0;
return;
}

4
tests/out/quad-vert.Vertex.glsl
View File

@@ -11,18 +11,14 @@ struct type10 {
};
vec2 v_uv = vec2(0, 0);
vec2 a_uv = vec2(0, 0);
struct gen_gl_PerVertex_block_0 {
vec4 gen_gl_Position;
float gen_gl_PointSize;
float gen_gl_ClipDistance[1];
float gen_gl_CullDistance[1];
} perVertexStruct;
vec2 a_pos = vec2(0, 0);
layout(location = 1) in vec2 _p2vs_location1;
layout(location = 0) in vec2 _p2vs_location0;
smooth layout(location = 0) out vec2 _vs2fs_location0;

2
tests/out/shadow.Fragment.glsl
View File

@@ -11,11 +11,9 @@ struct Light {
uniform Globals_block_0 {
uvec4 num_lights;
} _group_0_binding_0;
readonly buffer Lights_block_1 {
Light data[];
} _group_0_binding_1;
uniform highp sampler2DArrayShadow _group_0_binding_2;
smooth layout(location = 0) in vec3 _vs2fs_location0;

1
tests/out/skybox.Vertex.glsl
View File

@@ -11,7 +11,6 @@ uniform Data_block_0 {
mat4x4 proj_inv;
mat4x4 view;
} _group_0_binding_0;
smooth layout(location = 0) out vec3 _vs2fs_location0;
void main() {