nordic-dev.net/wgpu
nordic-dev.net/wgpu
2
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2024-11-22 23:04:07 +00:00
Code Issues Packages Projects Releases Wiki Activity

[glsl-out] Handle vector bitcasts (#1966)

Browse Source 
* [spv-in] Fix bitcasts on non-scalars

* Fix ::As handling in glsl

* Cargo fmt

* Make cargo clippy happy

* Add snapshot tests

* Use write_type_value

* target_vector_kind -> target_vector_type
This commit is contained in:
Ashley 2022-06-21 07:25:13 +02:00 committed by GitHub
parent 6d78f1c06d
commit ea832a9eec
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
8 changed files with 327 additions and 30 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -2931,38 +2931,50 @@ impl<'a, W: Write> Writer<'a, W> {
None => {
use crate::ScalarKind as Sk;
let source_kind = inner.scalar_kind().unwrap();
let conv_op = match (source_kind, target_kind) {
(Sk::Float, Sk::Sint) => "floatBitsToInt",
(Sk::Float, Sk::Uint) => "floatBitsToUint",
(Sk::Sint, Sk::Float) => "intBitsToFloat",
(Sk::Uint, Sk::Float) => "uintBitsToFloat",
// There is no way to bitcast between Uint/Sint in glsl. Use constructor conversion
(Sk::Uint, Sk::Sint) => "int",
(Sk::Sint, Sk::Uint) => "uint",
(Sk::Bool, Sk::Sint) => "int",
(Sk::Bool, Sk::Uint) => "uint",
(Sk::Bool, Sk::Float) => "float",
(Sk::Sint, Sk::Bool) => "bool",
(Sk::Uint, Sk::Bool) => "bool",
(Sk::Float, Sk::Bool) => "bool",
// No conversion needed
(Sk::Sint, Sk::Sint) => "",
(Sk::Uint, Sk::Uint) => "",
(Sk::Float, Sk::Float) => "",
(Sk::Bool, Sk::Bool) => "",
let target_vector_type = match *inner {
TypeInner::Vector { size, width, .. } => Some(TypeInner::Vector {
size,
width,
kind: target_kind,
}),
_ => None,
};
write!(self.out, "{}", conv_op)?;
if !conv_op.is_empty() {
write!(self.out, "(")?;
}
let source_kind = inner.scalar_kind().unwrap();
match (source_kind, target_kind, target_vector_type) {
// No conversion needed
(Sk::Sint, Sk::Sint, _)
| (Sk::Uint, Sk::Uint, _)
| (Sk::Float, Sk::Float, _)
| (Sk::Bool, Sk::Bool, _) => {
self.write_expr(expr, ctx)?;
return Ok(());
}
// Cast to/from floats
(Sk::Float, Sk::Sint, _) => write!(self.out, "floatBitsToInt")?,
(Sk::Float, Sk::Uint, _) => write!(self.out, "floatBitsToUint")?,
(Sk::Sint, Sk::Float, _) => write!(self.out, "intBitsToFloat")?,
(Sk::Uint, Sk::Float, _) => write!(self.out, "uintBitsToFloat")?,
// Cast between vector types
(_, _, Some(vector)) => {
self.write_value_type(&vector)?;
}
// There is no way to bitcast between Uint/Sint in glsl. Use constructor conversion
(Sk::Uint | Sk::Bool, Sk::Sint, None) => write!(self.out, "int")?,
(Sk::Sint | Sk::Bool, Sk::Uint, None) => write!(self.out, "uint")?,
(Sk::Bool, Sk::Float, None) => write!(self.out, "float")?,
(Sk::Sint | Sk::Uint | Sk::Float, Sk::Bool, None) => {
write!(self.out, "bool")?
}
};
write!(self.out, "(")?;
self.write_expr(expr, ctx)?;
if !conv_op.is_empty() {
write!(self.out, ")")?
}
write!(self.out, ")")?;
}
}
}

15
tests/in/bitcast.params.ron Normal file
View File

@ -0,0 +1,15 @@
(
msl: (
lang_version: (1, 2),
per_stage_map: (
cs: (
resources: {
},
sizes_buffer: Some(0),
)
),
inline_samplers: [],
spirv_cross_compatibility: false,
fake_missing_bindings: false,
),
)

26
tests/in/bitcast.wgsl Normal file
View File

@ -0,0 +1,26 @@
@compute @workgroup_size(1)
fn main() {
var i2 = vec2<i32>(0);
var i3 = vec3<i32>(0);
var i4 = vec4<i32>(0);
var u2 = vec2<u32>(0u);
var u3 = vec3<u32>(0u);
var u4 = vec4<u32>(0u);
var f2 = vec2<f32>(0.0);
var f3 = vec3<f32>(0.0);
var f4 = vec4<f32>(0.0);
u2 = bitcast<vec2<u32>>(i2);
u3 = bitcast<vec3<u32>>(i3);
u4 = bitcast<vec4<u32>>(i4);
i2 = bitcast<vec2<i32>>(u2);
i3 = bitcast<vec3<i32>>(u3);
i4 = bitcast<vec4<i32>>(u4);
f2 = bitcast<vec2<f32>>(i2);
f3 = bitcast<vec3<f32>>(i3);
f4 = bitcast<vec4<f32>>(i4);
}

48
tests/out/glsl/bitcast.main.Compute.glsl Normal file
View File

@ -0,0 +1,48 @@
#version 310 es
precision highp float;
precision highp int;
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
ivec2 i2_ = ivec2(0);
ivec3 i3_ = ivec3(0);
ivec4 i4_ = ivec4(0);
uvec2 u2_ = uvec2(0u);
uvec3 u3_ = uvec3(0u);
uvec4 u4_ = uvec4(0u);
vec2 f2_ = vec2(0.0);
vec3 f3_ = vec3(0.0);
vec4 f4_ = vec4(0.0);
i2_ = ivec2(0);
i3_ = ivec3(0);
i4_ = ivec4(0);
u2_ = uvec2(0u);
u3_ = uvec3(0u);
u4_ = uvec4(0u);
f2_ = vec2(0.0);
f3_ = vec3(0.0);
f4_ = vec4(0.0);
ivec2 _e27 = i2_;
u2_ = uvec2(_e27);
ivec3 _e29 = i3_;
u3_ = uvec3(_e29);
ivec4 _e31 = i4_;
u4_ = uvec4(_e31);
uvec2 _e33 = u2_;
i2_ = ivec2(_e33);
uvec3 _e35 = u3_;
i3_ = ivec3(_e35);
uvec4 _e37 = u4_;
i4_ = ivec4(_e37);
ivec2 _e39 = i2_;
f2_ = intBitsToFloat(_e39);
ivec3 _e41 = i3_;
f3_ = intBitsToFloat(_e41);
ivec4 _e43 = i4_;
f4_ = intBitsToFloat(_e43);
return;
}

47
tests/out/msl/bitcast.msl Normal file
View File

@ -0,0 +1,47 @@
// language: metal2.0
#include <metal_stdlib>
#include <simd/simd.h>
using metal::uint;
kernel void main_(
) {
metal::int2 i2_ = {};
metal::int3 i3_ = {};
metal::int4 i4_ = {};
metal::uint2 u2_ = {};
metal::uint3 u3_ = {};
metal::uint4 u4_ = {};
metal::float2 f2_ = {};
metal::float3 f3_ = {};
metal::float4 f4_ = {};
i2_ = metal::int2(0);
i3_ = metal::int3(0);
i4_ = metal::int4(0);
u2_ = metal::uint2(0u);
u3_ = metal::uint3(0u);
u4_ = metal::uint4(0u);
f2_ = metal::float2(0.0);
f3_ = metal::float3(0.0);
f4_ = metal::float4(0.0);
metal::int2 _e27 = i2_;
u2_ = as_type<metal::uint2>(_e27);
metal::int3 _e29 = i3_;
u3_ = as_type<metal::uint3>(_e29);
metal::int4 _e31 = i4_;
u4_ = as_type<metal::uint4>(_e31);
metal::uint2 _e33 = u2_;
i2_ = as_type<metal::int2>(_e33);
metal::uint3 _e35 = u3_;
i3_ = as_type<metal::int3>(_e35);
metal::uint4 _e37 = u4_;
i4_ = as_type<metal::int4>(_e37);
metal::int2 _e39 = i2_;
f2_ = as_type<metal::float2>(_e39);
metal::int3 _e41 = i3_;
f3_ = as_type<metal::float3>(_e41);
metal::int4 _e43 = i4_;
f4_ = as_type<metal::float4>(_e43);
return;
}

104
tests/out/spv/bitcast.spvasm Normal file
View File

@ -0,0 +1,104 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 76
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %46 "main"
OpExecutionMode %46 LocalSize 1 1 1
%2 = OpTypeVoid
%4 = OpTypeInt 32 1
%3 = OpConstant %4 0
%6 = OpTypeInt 32 0
%5 = OpConstant %6 0
%8 = OpTypeFloat 32
%7 = OpConstant %8 0.0
%9 = OpTypeVector %4 2
%10 = OpTypeVector %4 3
%11 = OpTypeVector %4 4
%12 = OpTypeVector %6 2
%13 = OpTypeVector %6 3
%14 = OpTypeVector %6 4
%15 = OpTypeVector %8 2
%16 = OpTypeVector %8 3
%17 = OpTypeVector %8 4
%19 = OpTypePointer Function %9
%20 = OpConstantNull %9
%22 = OpTypePointer Function %10
%23 = OpConstantNull %10
%25 = OpTypePointer Function %11
%26 = OpConstantNull %11
%28 = OpTypePointer Function %12
%29 = OpConstantNull %12
%31 = OpTypePointer Function %13
%32 = OpConstantNull %13
%34 = OpTypePointer Function %14
%35 = OpConstantNull %14
%37 = OpTypePointer Function %15
%38 = OpConstantNull %15
%40 = OpTypePointer Function %16
%41 = OpConstantNull %16
%43 = OpTypePointer Function %17
%44 = OpConstantNull %17
%47 = OpTypeFunction %2
%46 = OpFunction %2 None %47
%45 = OpLabel
%36 = OpVariable %37 Function %38
%27 = OpVariable %28 Function %29
%18 = OpVariable %19 Function %20
%39 = OpVariable %40 Function %41
%30 = OpVariable %31 Function %32
%21 = OpVariable %22 Function %23
%42 = OpVariable %43 Function %44
%33 = OpVariable %34 Function %35
%24 = OpVariable %25 Function %26
OpBranch %48
%48 = OpLabel
%49 = OpCompositeConstruct %9 %3 %3
OpStore %18 %49
%50 = OpCompositeConstruct %10 %3 %3 %3
OpStore %21 %50
%51 = OpCompositeConstruct %11 %3 %3 %3 %3
OpStore %24 %51
%52 = OpCompositeConstruct %12 %5 %5
OpStore %27 %52
%53 = OpCompositeConstruct %13 %5 %5 %5
OpStore %30 %53
%54 = OpCompositeConstruct %14 %5 %5 %5 %5
OpStore %33 %54
%55 = OpCompositeConstruct %15 %7 %7
OpStore %36 %55
%56 = OpCompositeConstruct %16 %7 %7 %7
OpStore %39 %56
%57 = OpCompositeConstruct %17 %7 %7 %7 %7
OpStore %42 %57
%58 = OpLoad %9 %18
%59 = OpBitcast %12 %58
OpStore %27 %59
%60 = OpLoad %10 %21
%61 = OpBitcast %13 %60
OpStore %30 %61
%62 = OpLoad %11 %24
%63 = OpBitcast %14 %62
OpStore %33 %63
%64 = OpLoad %12 %27
%65 = OpBitcast %9 %64
OpStore %18 %65
%66 = OpLoad %13 %30
%67 = OpBitcast %10 %66
OpStore %21 %67
%68 = OpLoad %14 %33
%69 = OpBitcast %11 %68
OpStore %24 %69
%70 = OpLoad %9 %18
%71 = OpBitcast %15 %70
OpStore %36 %71
%72 = OpLoad %10 %21
%73 = OpBitcast %16 %72
OpStore %39 %73
%74 = OpLoad %11 %24
%75 = OpBitcast %17 %74
OpStore %42 %75
OpReturn
OpFunctionEnd

41
tests/out/wgsl/bitcast.wgsl Normal file
View File

@ -0,0 +1,41 @@
@compute @workgroup_size(1, 1, 1)
fn main() {
var i2_: vec2<i32>;
var i3_: vec3<i32>;
var i4_: vec4<i32>;
var u2_: vec2<u32>;
var u3_: vec3<u32>;
var u4_: vec4<u32>;
var f2_: vec2<f32>;
var f3_: vec3<f32>;
var f4_: vec4<f32>;
i2_ = vec2<i32>(0);
i3_ = vec3<i32>(0);
i4_ = vec4<i32>(0);
u2_ = vec2<u32>(0u);
u3_ = vec3<u32>(0u);
u4_ = vec4<u32>(0u);
f2_ = vec2<f32>(0.0);
f3_ = vec3<f32>(0.0);
f4_ = vec4<f32>(0.0);
let _e27 = i2_;
u2_ = bitcast<vec2<u32>>(_e27);
let _e29 = i3_;
u3_ = bitcast<vec3<u32>>(_e29);
let _e31 = i4_;
u4_ = bitcast<vec4<u32>>(_e31);
let _e33 = u2_;
i2_ = bitcast<vec2<i32>>(_e33);
let _e35 = u3_;
i3_ = bitcast<vec3<i32>>(_e35);
let _e37 = u4_;
i4_ = bitcast<vec4<i32>>(_e37);
let _e39 = i2_;
f2_ = bitcast<vec2<f32>>(_e39);
let _e41 = i3_;
f3_ = bitcast<vec3<f32>>(_e41);
let _e43 = i4_;
f4_ = bitcast<vec4<f32>>(_e43);
return;
}

4
tests/snapshots.rs
View File

@ -435,6 +435,10 @@ fn convert_wgsl() {
"bits",
Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::WGSL,
),
(
"bitcast",
Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::WGSL,
),
(
"boids",
Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL,