Skip to content

Commit 5a583b1

Browse files
DJMcNabDJMcNab
authored
Polyfill unpackUnorm4x8 and friends on unsupported GLSL versions (gfx-rs#7408)
1 parent 4791731 commit 5a583b1

File tree

9 files changed

+379
-10
lines changed

9 files changed

+379
-10
lines changed

‎CHANGELOG.md‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -193,6 +193,7 @@ By @wumpf in [#7144](https://github.com/gfx-rs/wgpu/pull/7144)
193193
- Support for generating the candidate intersections from AABB geometry, and confirming the hits. By @kvark in [#7047](https://github.com/gfx-rs/wgpu/pull/7047).
194194
- Make naga::back::spv::Function::to_words write the OpFunctionEnd instruction in itself, instead of making another call after it. By @junjunjd in [#7156](https://github.com/gfx-rs/wgpu/pull/7156).
195195
- Add support for texture memory barriers. By @Devon7925 in [#7173](https://github.com/gfx-rs/wgpu/pull/7173).
196+
- Add polyfills for `unpackSnorm4x8`, `unpackUnorm4x8`, `unpackSnorm2x16`, `unpackUnorm2x16` for GLSL versions they aren't supported in. By @DJMcNab in [#7408](https://github.com/gfx-rs/wgpu/pull/7408).
196197
197198
### Changes
198199

‎naga/src/back/glsl/mod.rs‎

Lines changed: 157 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -230,6 +230,30 @@ impl Version {
230230
fn supports_derivative_control(&self) -> bool {
231231
*self >= Version::Desktop(450)
232232
}
233+
234+
// For supports_pack_unpack_4x8, supports_pack_unpack_snorm_2x16, supports_pack_unpack_unorm_2x16
235+
// see:
236+
// https://registry.khronos.org/OpenGL-Refpages/gl4/html/unpackUnorm.xhtml
237+
// https://registry.khronos.org/OpenGL-Refpages/es3/html/unpackUnorm.xhtml
238+
// https://registry.khronos.org/OpenGL-Refpages/gl4/html/packUnorm.xhtml
239+
// https://registry.khronos.org/OpenGL-Refpages/es3/html/packUnorm.xhtml
240+
fn supports_pack_unpack_4x8(&self) -> bool {
241+
*self >= Version::Desktop(400) || *self >= Version::new_gles(310)
242+
}
243+
fn supports_pack_unpack_snorm_2x16(&self) -> bool {
244+
*self >= Version::Desktop(420) || *self >= Version::new_gles(300)
245+
}
246+
fn supports_pack_unpack_unorm_2x16(&self) -> bool {
247+
*self >= Version::Desktop(400) || *self >= Version::new_gles(300)
248+
}
249+
250+
// https://registry.khronos.org/OpenGL-Refpages/gl4/html/unpackHalf2x16.xhtml
251+
// https://registry.khronos.org/OpenGL-Refpages/gl4/html/packHalf2x16.xhtml
252+
// https://registry.khronos.org/OpenGL-Refpages/es3/html/unpackHalf2x16.xhtml
253+
// https://registry.khronos.org/OpenGL-Refpages/es3/html/packHalf2x16.xhtml
254+
fn supports_pack_unpack_half_2x16(&self) -> bool {
255+
*self >= Version::Desktop(420) || *self >= Version::new_gles(300)
256+
}
233257
}
234258

235259
impl PartialOrd for Version {
@@ -1369,6 +1393,31 @@ impl<'a, W: Write> Writer<'a, W> {
13691393
| crate::MathFunction::QuantizeToF16 => {
13701394
self.need_bake_expressions.insert(arg);
13711395
}
1396+
/* crate::MathFunction::Pack4x8unorm | */
1397+
crate::MathFunction::Unpack4x8snorm
1398+
if !self.options.version.supports_pack_unpack_4x8() =>
1399+
{
1400+
// We have a fallback if the platform doesn't natively support these
1401+
self.need_bake_expressions.insert(arg);
1402+
}
1403+
/* crate::MathFunction::Pack4x8unorm | */
1404+
crate::MathFunction::Unpack4x8unorm
1405+
if !self.options.version.supports_pack_unpack_4x8() =>
1406+
{
1407+
self.need_bake_expressions.insert(arg);
1408+
}
1409+
/* crate::MathFunction::Pack2x16snorm | */
1410+
crate::MathFunction::Unpack2x16snorm
1411+
if !self.options.version.supports_pack_unpack_snorm_2x16() =>
1412+
{
1413+
self.need_bake_expressions.insert(arg);
1414+
}
1415+
/* crate::MathFunction::Pack2x16unorm | */
1416+
crate::MathFunction::Unpack2x16unorm
1417+
if !self.options.version.supports_pack_unpack_unorm_2x16() =>
1418+
{
1419+
self.need_bake_expressions.insert(arg);
1420+
}
13721421
crate::MathFunction::ExtractBits => {
13731422
// Only argument 1 is re-used.
13741423
self.need_bake_expressions.insert(arg1.unwrap());
@@ -3756,11 +3805,43 @@ impl<'a, W: Write> Writer<'a, W> {
37563805
Mf::FirstTrailingBit => "findLSB",
37573806
Mf::FirstLeadingBit => "findMSB",
37583807
// data packing
3759-
Mf::Pack4x8snorm => "packSnorm4x8",
3760-
Mf::Pack4x8unorm => "packUnorm4x8",
3761-
Mf::Pack2x16snorm => "packSnorm2x16",
3762-
Mf::Pack2x16unorm => "packUnorm2x16",
3763-
Mf::Pack2x16float => "packHalf2x16",
3808+
Mf::Pack4x8snorm => {
3809+
if self.options.version.supports_pack_unpack_4x8() {
3810+
"packSnorm4x8"
3811+
} else {
3812+
// polyfill should go here. Needs a corresponding entry in `need_bake_expression`
3813+
return Err(Error::UnsupportedExternal("packSnorm4x8".into()));
3814+
}
3815+
}
3816+
Mf::Pack4x8unorm => {
3817+
if self.options.version.supports_pack_unpack_4x8() {
3818+
"packUnorm4x8"
3819+
} else {
3820+
return Err(Error::UnsupportedExternal("packUnorm4x8".to_owned()));
3821+
}
3822+
}
3823+
Mf::Pack2x16snorm => {
3824+
if self.options.version.supports_pack_unpack_snorm_2x16() {
3825+
"packSnorm2x16"
3826+
} else {
3827+
return Err(Error::UnsupportedExternal("packSnorm2x16".to_owned()));
3828+
}
3829+
}
3830+
Mf::Pack2x16unorm => {
3831+
if self.options.version.supports_pack_unpack_unorm_2x16() {
3832+
"packUnorm2x16"
3833+
} else {
3834+
return Err(Error::UnsupportedExternal("packUnorm2x16".to_owned()));
3835+
}
3836+
}
3837+
Mf::Pack2x16float => {
3838+
if self.options.version.supports_pack_unpack_half_2x16() {
3839+
"packHalf2x16"
3840+
} else {
3841+
return Err(Error::UnsupportedExternal("packHalf2x16".to_owned()));
3842+
}
3843+
}
3844+
37643845
fun @ (Mf::Pack4xI8 | Mf::Pack4xU8) => {
37653846
let was_signed = match fun {
37663847
Mf::Pack4xI8 => true,
@@ -3787,11 +3868,77 @@ impl<'a, W: Write> Writer<'a, W> {
37873868
return Ok(());
37883869
}
37893870
// data unpacking
3790-
Mf::Unpack4x8snorm => "unpackSnorm4x8",
3791-
Mf::Unpack4x8unorm => "unpackUnorm4x8",
3792-
Mf::Unpack2x16snorm => "unpackSnorm2x16",
3793-
Mf::Unpack2x16unorm => "unpackUnorm2x16",
3794-
Mf::Unpack2x16float => "unpackHalf2x16",
3871+
Mf::Unpack2x16float => {
3872+
if self.options.version.supports_pack_unpack_half_2x16() {
3873+
"unpackHalf2x16"
3874+
} else {
3875+
return Err(Error::UnsupportedExternal("unpackHalf2x16".into()));
3876+
}
3877+
}
3878+
Mf::Unpack2x16snorm => {
3879+
if self.options.version.supports_pack_unpack_snorm_2x16() {
3880+
"unpackSnorm2x16"
3881+
} else {
3882+
let scale = 32767;
3883+
3884+
write!(self.out, "(vec2(ivec2(")?;
3885+
self.write_expr(arg, ctx)?;
3886+
write!(self.out, " << 16, ")?;
3887+
self.write_expr(arg, ctx)?;
3888+
write!(self.out, ") >> 16) / {scale}.0)")?;
3889+
return Ok(());
3890+
}
3891+
}
3892+
Mf::Unpack2x16unorm => {
3893+
if self.options.version.supports_pack_unpack_unorm_2x16() {
3894+
"unpackUnorm2x16"
3895+
} else {
3896+
let scale = 65535;
3897+
3898+
write!(self.out, "(vec2(")?;
3899+
self.write_expr(arg, ctx)?;
3900+
write!(self.out, " & 0xFFFFu, ")?;
3901+
self.write_expr(arg, ctx)?;
3902+
write!(self.out, " >> 16) / {scale}.0)")?;
3903+
return Ok(());
3904+
}
3905+
}
3906+
Mf::Unpack4x8snorm => {
3907+
if self.options.version.supports_pack_unpack_4x8() {
3908+
"unpackSnorm4x8"
3909+
} else {
3910+
let scale = 127;
3911+
3912+
write!(self.out, "(vec4(ivec4(")?;
3913+
self.write_expr(arg, ctx)?;
3914+
write!(self.out, " << 24, ")?;
3915+
self.write_expr(arg, ctx)?;
3916+
write!(self.out, " << 16, ")?;
3917+
self.write_expr(arg, ctx)?;
3918+
write!(self.out, " << 8, ")?;
3919+
self.write_expr(arg, ctx)?;
3920+
write!(self.out, ") >> 24) / {scale}.0)")?;
3921+
return Ok(());
3922+
}
3923+
}
3924+
Mf::Unpack4x8unorm => {
3925+
if self.options.version.supports_pack_unpack_4x8() {
3926+
"unpackUnorm4x8"
3927+
} else {
3928+
let scale = 255;
3929+
3930+
write!(self.out, "(vec4(")?;
3931+
self.write_expr(arg, ctx)?;
3932+
write!(self.out, " & 0xFFu, ")?;
3933+
self.write_expr(arg, ctx)?;
3934+
write!(self.out, " >> 8 & 0xFFu, ")?;
3935+
self.write_expr(arg, ctx)?;
3936+
write!(self.out, " >> 16 & 0xFFu, ")?;
3937+
self.write_expr(arg, ctx)?;
3938+
write!(self.out, " >> 24) / {scale}.0)")?;
3939+
return Ok(());
3940+
}
3941+
}
37953942
fun @ (Mf::Unpack4xI8 | Mf::Unpack4xU8) => {
37963943
let sign_prefix = match fun {
37973944
Mf::Unpack4xI8 => 'i',

‎naga/tests/in/wgsl/bits.wgsl‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,3 +1,5 @@
1+
// Keep in sync with `bits_downlevel` and `bits_downlevel_webgl`
2+
13
@compute @workgroup_size(1)
24
fn main() {
35
var i = 0;

‎naga/tests/in/wgsl/bits_downlevel.toml‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,4 @@
1+
targets = "GLSL"
2+
3+
[glsl]
4+
version.Desktop = 330

‎naga/tests/in/wgsl/bits_downlevel.wgsl‎

Lines changed: 76 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,76 @@
1+
// Keep in sync with bits.wgsl
2+
3+
@fragment
4+
fn main() {
5+
var i = 0;
6+
var i2 = vec2<i32>(0);
7+
var i3 = vec3<i32>(0);
8+
var i4 = vec4<i32>(0);
9+
var u = 0u;
10+
var u2 = vec2<u32>(0u);
11+
var u3 = vec3<u32>(0u);
12+
var u4 = vec4<u32>(0u);
13+
var f2 = vec2<f32>(0.0);
14+
var f4 = vec4<f32>(0.0);
15+
// No polyfill for these yet
16+
// u = pack4x8snorm(f4);
17+
// u = pack4x8unorm(f4);
18+
// u = pack2x16snorm(f2);
19+
// u = pack2x16unorm(f2);
20+
// u = pack2x16float(f2);
21+
u = pack4xI8(i4);
22+
u = pack4xU8(u4);
23+
f4 = unpack4x8snorm(u);
24+
f4 = unpack4x8unorm(u);
25+
f2 = unpack2x16snorm(u);
26+
f2 = unpack2x16unorm(u);
27+
// No polyfill for this yet
28+
// f2 = unpack2x16float(u);
29+
// Polyfill for this is broken in downlevel
30+
// i4 = unpack4xI8(u);
31+
// u4 = unpack4xU8(u);
32+
// Implementation is broken on downlevel
33+
// i = insertBits(i, i, 5u, 10u);
34+
// i2 = insertBits(i2, i2, 5u, 10u);
35+
// i3 = insertBits(i3, i3, 5u, 10u);
36+
// i4 = insertBits(i4, i4, 5u, 10u);
37+
// u = insertBits(u, u, 5u, 10u);
38+
// u2 = insertBits(u2, u2, 5u, 10u);
39+
// u3 = insertBits(u3, u3, 5u, 10u);
40+
// u4 = insertBits(u4, u4, 5u, 10u);
41+
// Implementation is broken on downlevel
42+
// i = extractBits(i, 5u, 10u);
43+
// i2 = extractBits(i2, 5u, 10u);
44+
// i3 = extractBits(i3, 5u, 10u);
45+
// i4 = extractBits(i4, 5u, 10u);
46+
// u = extractBits(u, 5u, 10u);
47+
// u2 = extractBits(u2, 5u, 10u);
48+
// u3 = extractBits(u3, 5u, 10u);
49+
// u4 = extractBits(u4, 5u, 10u);
50+
// Implementation is broken on downlevel
51+
// i = firstTrailingBit(i);
52+
// u2 = firstTrailingBit(u2);
53+
// i3 = firstLeadingBit(i3);
54+
// u3 = firstLeadingBit(u3);
55+
// Implementation is broken on downlevel
56+
// i = firstLeadingBit(i);
57+
// u = firstLeadingBit(u);
58+
// Implementation is broken on downlevel
59+
// i = countOneBits(i);
60+
// i2 = countOneBits(i2);
61+
// i3 = countOneBits(i3);
62+
// i4 = countOneBits(i4);
63+
// u = countOneBits(u);
64+
// u2 = countOneBits(u2);
65+
// u3 = countOneBits(u3);
66+
// u4 = countOneBits(u4);
67+
// Implementation is broken on downlevel
68+
// i = reverseBits(i);
69+
// i2 = reverseBits(i2);
70+
// i3 = reverseBits(i3);
71+
// i4 = reverseBits(i4);
72+
// u = reverseBits(u);
73+
// u2 = reverseBits(u2);
74+
// u3 = reverseBits(u3);
75+
// u4 = reverseBits(u4);
76+
}

‎naga/tests/in/wgsl/bits_downlevel_webgl.toml‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,4 @@
1+
targets = "GLSL"
2+
3+
[glsl]
4+
version.Embedded = { is_webgl = true, version = 300 }

‎naga/tests/in/wgsl/bits_downlevel_webgl.wgsl‎

Lines changed: 75 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,75 @@
1+
// Keep in sync with bits.wgsl
2+
3+
@fragment
4+
fn main() {
5+
var i = 0;
6+
var i2 = vec2<i32>(0);
7+
var i3 = vec3<i32>(0);
8+
var i4 = vec4<i32>(0);
9+
var u = 0u;
10+
var u2 = vec2<u32>(0u);
11+
var u3 = vec3<u32>(0u);
12+
var u4 = vec4<u32>(0u);
13+
var f2 = vec2<f32>(0.0);
14+
var f4 = vec4<f32>(0.0);
15+
// No polyfill for these yet
16+
// u = pack4x8snorm(f4);
17+
// u = pack4x8unorm(f4);
18+
// u = pack2x16snorm(f2);
19+
// u = pack2x16unorm(f2);
20+
// u = pack2x16float(f2);
21+
u = pack4xI8(i4);
22+
u = pack4xU8(u4);
23+
f4 = unpack4x8snorm(u);
24+
f4 = unpack4x8unorm(u);
25+
f2 = unpack2x16snorm(u);
26+
f2 = unpack2x16unorm(u);
27+
// No polyfill for this yet
28+
// f2 = unpack2x16float(u);
29+
// Polyfill for this is broken in downlevel
30+
// i4 = unpack4xI8(u);
31+
// u4 = unpack4xU8(u);
32+
// Implementation is broken on downlevel
33+
// i = insertBits(i, i, 5u, 10u);
34+
// i2 = insertBits(i2, i2, 5u, 10u);
35+
// i3 = insertBits(i3, i3, 5u, 10u);
36+
// i4 = insertBits(i4, i4, 5u, 10u);
37+
// u = insertBits(u, u, 5u, 10u);
38+
// u2 = insertBits(u2, u2, 5u, 10u);
39+
// u3 = insertBits(u3, u3, 5u, 10u);
40+
// u4 = insertBits(u4, u4, 5u, 10u);
41+
// Implementation is broken on downlevel
42+
// i = extractBits(i, 5u, 10u);
43+
// i2 = extractBits(i2, 5u, 10u);
44+
// i3 = extractBits(i3, 5u, 10u);
45+
// i4 = extractBits(i4, 5u, 10u);
46+
// u = extractBits(u, 5u, 10u);
47+
// u2 = extractBits(u2, 5u, 10u);
48+
// u3 = extractBits(u3, 5u, 10u);
49+
// u4 = extractBits(u4, 5u, 10u);
50+
// Implementation is broken on downlevel
51+
// i = firstTrailingBit(i);
52+
// u2 = firstTrailingBit(u2);
53+
// Implementation is broken on downlevel
54+
// i3 = firstLeadingBit(i3);
55+
// u3 = firstLeadingBit(u3);
56+
// i = firstLeadingBit(i);
57+
// u = firstLeadingBit(u);
58+
// Implementation is broken on downlevel
59+
// i = countOneBits(i);
60+
// i2 = countOneBits(i2);
61+
// i3 = countOneBits(i3);
62+
// i4 = countOneBits(i4);
63+
// u = countOneBits(u);
64+
// u2 = countOneBits(u2);
65+
// u3 = countOneBits(u3);
66+
// u4 = countOneBits(u4);
67+
// i = reverseBits(i);
68+
// i2 = reverseBits(i2);
69+
// i3 = reverseBits(i3);
70+
// i4 = reverseBits(i4);
71+
// u = reverseBits(u);
72+
// u2 = reverseBits(u2);
73+
// u3 = reverseBits(u3);
74+
// u4 = reverseBits(u4);
75+
}

0 commit comments

Comments
 (0)