Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

[glsl/hlsl-out] Write sizes of arrays behind pointers in function arguments #2250

Merged
merged 4 commits into from
Feb 17, 2023
Merged

[glsl/hlsl-out] Write sizes of arrays behind pointers in function arguments #2250

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
e168c62
[glsl-out] write sizes of arrays behind pointers in function arguments
pluiedev Feb 17, 2023
93fedaf
add assign array through pointer fn test
pluiedev Feb 17, 2023
618834b
update snapshots
pluiedev Feb 17, 2023
b43b0bd
[hlsl-out] write sizes of arrays behind pointers in function arguments
pluiedev Feb 17, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
12 changes: 10 additions & 2 deletions src/back/glsl/mod.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1439,8 +1439,16 @@ impl<'a, W: Write> Writer<'a, W> {
write!(this.out, " {}", &this.names[&ctx.argument_key(i as u32)])?;

// Write array size
if let TypeInner::Array { base, size, .. } = this.module.types[arg.ty].inner {
this.write_array_size(base, size)?;
match this.module.types[arg.ty].inner {
TypeInner::Array { base, size, .. } => {
this.write_array_size(base, size)?;
}
TypeInner::Pointer { base, .. } => {
if let TypeInner::Array { base, size, .. } = this.module.types[base].inner {
this.write_array_size(base, size)?;
}
}
_ => {}
}

Ok(())
Expand Down
2 changes: 1 addition & 1 deletion src/back/hlsl/writer.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1143,7 +1143,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {

// Write argument name. Space is important.
write!(self.out, " {argument_name}")?;
if let TypeInner::Array { base, size, .. } = module.types[arg.ty].inner {
if let TypeInner::Array { base, size, .. } = module.types[arg_ty].inner {
jimblandy marked this conversation as resolved.
Show resolved Hide resolved
self.write_array_size(module, base, size)?;
}
}
Expand Down
9 changes: 8 additions & 1 deletion tests/in/access.wgsl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -174,7 +174,14 @@ fn assign_through_ptr_fn(p: ptr<workgroup, u32>) {
*p = 42u;
}

fn assign_array_through_ptr_fn(foo: ptr<function, array<vec4<f32>, 2>>) {
*foo = array<vec4<f32>, 2>(vec4(1.0), vec4(2.0));
}

@compute @workgroup_size(1)
fn assign_through_ptr() {
var arr = array<vec4<f32>, 2>(vec4(6.0), vec4(7.0));

assign_through_ptr_fn(&val);
}
assign_array_through_ptr_fn(&arr);
}
208 changes: 208 additions & 0 deletions tests/out/analysis/access.info.ron
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -93,6 +93,12 @@
(
bits: 70,
),
(
bits: 215,
),
(
bits: 70,
),
],
functions: [
(
Expand Down Expand Up @@ -3588,6 +3594,125 @@
],
sampling: [],
),
(
flags: (
bits: 63,
),
available_stages: (
bits: 7,
),
uniformity: (
non_uniform_result: None,
requirements: (
bits: 0,
),
),
may_kill: false,
sampling_set: [],
global_uses: [
(
bits: 0,
),
(
bits: 0,
),
(
bits: 0,
),
(
bits: 0,
),
(
bits: 0,
),
(
bits: 0,
),
],
expressions: [
(
uniformity: (
non_uniform_result: Some(1),
requirements: (
bits: 0,
),
),
ref_count: 1,
assignable_global: None,
ty: Handle(33),
),
(
uniformity: (
non_uniform_result: None,
requirements: (
bits: 0,
),
),
ref_count: 1,
assignable_global: None,
ty: Value(Scalar(
kind: Float,
width: 4,
)),
),
(
uniformity: (
non_uniform_result: None,
requirements: (
bits: 0,
),
),
ref_count: 1,
assignable_global: None,
ty: Value(Vector(
size: Quad,
kind: Float,
width: 4,
)),
),
(
uniformity: (
non_uniform_result: None,
requirements: (
bits: 0,
),
),
ref_count: 1,
assignable_global: None,
ty: Value(Scalar(
kind: Float,
width: 4,
)),
),
(
uniformity: (
non_uniform_result: None,
requirements: (
bits: 0,
),
),
ref_count: 1,
assignable_global: None,
ty: Value(Vector(
size: Quad,
kind: Float,
width: 4,
)),
),
(
uniformity: (
non_uniform_result: None,
requirements: (
bits: 0,
),
),
ref_count: 1,
assignable_global: None,
ty: Handle(32),
),
],
sampling: [],
),
],
entry_points: [
(
Expand Down Expand Up @@ -5661,6 +5786,89 @@
),
],
expressions: [
(
uniformity: (
non_uniform_result: None,
requirements: (
bits: 0,
),
),
ref_count: 1,
assignable_global: None,
ty: Value(Scalar(
kind: Float,
width: 4,
)),
),
(
uniformity: (
non_uniform_result: None,
requirements: (
bits: 0,
),
),
ref_count: 1,
assignable_global: None,
ty: Value(Vector(
size: Quad,
kind: Float,
width: 4,
)),
),
(
uniformity: (
non_uniform_result: None,
requirements: (
bits: 0,
),
),
ref_count: 1,
assignable_global: None,
ty: Value(Scalar(
kind: Float,
width: 4,
)),
),
(
uniformity: (
non_uniform_result: None,
requirements: (
bits: 0,
),
),
ref_count: 1,
assignable_global: None,
ty: Value(Vector(
size: Quad,
kind: Float,
width: 4,
)),
),
(
uniformity: (
non_uniform_result: None,
requirements: (
bits: 0,
),
),
ref_count: 1,
assignable_global: None,
ty: Handle(32),
),
(
uniformity: (
non_uniform_result: Some(6),
requirements: (
bits: 0,
),
),
ref_count: 2,
assignable_global: None,
ty: Value(Pointer(
base: 32,
space: Function,
)),
),
(
uniformity: (
non_uniform_result: None,
Expand Down
8 changes: 8 additions & 0 deletions tests/out/glsl/access.assign_through_ptr.Compute.glsl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -36,13 +36,21 @@ void assign_through_ptr_fn(inout uint p) {
return;
}

void assign_array_through_ptr_fn(inout vec4 foo_2[2]) {
foo_2 = vec4[2](vec4(1.0), vec4(2.0));
return;
}

void main() {
if (gl_GlobalInvocationID == uvec3(0u)) {
val = 0u;
}
memoryBarrierShared();
barrier();
vec4 arr[2] = vec4[2](vec4(0.0), vec4(0.0));
arr = vec4[2](vec4(6.0), vec4(7.0));
assign_through_ptr_fn(val);
assign_array_through_ptr_fn(arr);
return;
}

5 changes: 5 additions & 0 deletions tests/out/glsl/access.atomics.Compute.glsl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -43,6 +43,11 @@ void assign_through_ptr_fn(inout uint p) {
return;
}

void assign_array_through_ptr_fn(inout vec4 foo_2[2]) {
foo_2 = vec4[2](vec4(1.0), vec4(2.0));
return;
}

void main() {
int tmp = 0;
int value = _group_0_binding_0_cs.atom;
Expand Down
5 changes: 5 additions & 0 deletions tests/out/glsl/access.foo_frag.Fragment.glsl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -44,6 +44,11 @@ void assign_through_ptr_fn(inout uint p) {
return;
}

void assign_array_through_ptr_fn(inout vec4 foo_2[2]) {
foo_2 = vec4[2](vec4(1.0), vec4(2.0));
return;
}

void main() {
_group_0_binding_0_fs._matrix[1][2] = 1.0;
_group_0_binding_0_fs._matrix = mat4x3(vec3(0.0), vec3(1.0), vec3(2.0), vec3(3.0));
Expand Down
7 changes: 6 additions & 1 deletion tests/out/glsl/access.foo_vert.Vertex.glsl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -121,6 +121,11 @@ void assign_through_ptr_fn(inout uint p) {
return;
}

void assign_array_through_ptr_fn(inout vec4 foo_2[2]) {
foo_2 = vec4[2](vec4(1.0), vec4(2.0));
return;
}

void main() {
uint vi = uint(gl_VertexID);
float foo = 0.0;
Expand All @@ -131,7 +136,7 @@ void main() {
test_matrix_within_struct_accesses();
test_matrix_within_array_within_struct_accesses();
mat4x3 _matrix = _group_0_binding_0_vs._matrix;
uvec2 arr[2] = _group_0_binding_0_vs.arr;
uvec2 arr_1[2] = _group_0_binding_0_vs.arr;
float b = _group_0_binding_0_vs._matrix[3][0];
int a_1 = _group_0_binding_0_vs.data[(uint(_group_0_binding_0_vs.data.length()) - 2u)].value;
ivec2 c = _group_0_binding_2_vs;
Expand Down
18 changes: 17 additions & 1 deletion tests/out/hlsl/access.hlsl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -237,6 +237,18 @@ void assign_through_ptr_fn(inout uint p)
return;
}

typedef float4 ret_Constructarray2_float4_[2];
ret_Constructarray2_float4_ Constructarray2_float4_(float4 arg0, float4 arg1) {
float4 ret[2] = { arg0, arg1 };
return ret;
}

void assign_array_through_ptr_fn(inout float4 foo_2[2])
{
foo_2 = Constructarray2_float4_((1.0).xxxx, (2.0).xxxx);
return;
}

uint NagaBufferLengthRW(RWByteAddressBuffer buffer)
{
uint ret;
Expand All @@ -261,7 +273,7 @@ float4 foo_vert(uint vi : SV_VertexID) : SV_Position
test_matrix_within_struct_accesses();
test_matrix_within_array_within_struct_accesses();
float4x3 _matrix = float4x3(asfloat(bar.Load3(0+0)), asfloat(bar.Load3(0+16)), asfloat(bar.Load3(0+32)), asfloat(bar.Load3(0+48)));
uint2 arr[2] = {asuint(bar.Load2(144+0)), asuint(bar.Load2(144+8))};
uint2 arr_1[2] = {asuint(bar.Load2(144+0)), asuint(bar.Load2(144+8))};
float b = asfloat(bar.Load(0+48+0));
int a_1 = asint(bar.Load(0+(((NagaBufferLengthRW(bar) - 160) / 8) - 2u)*8+160));
int2 c = asint(qux.Load2(0));
Expand Down Expand Up @@ -332,6 +344,10 @@ void assign_through_ptr(uint3 __global_invocation_id : SV_DispatchThreadID)
val = (uint)0;
}
GroupMemoryBarrierWithGroupSync();
float4 arr[2] = (float4[2])0;

arr = Constructarray2_float4_((6.0).xxxx, (7.0).xxxx);
assign_through_ptr_fn(val);
assign_array_through_ptr_fn(arr);
return;
}
Loading