nordic-dev.net/wgpu
nordic-dev.net/wgpu
2
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2025-02-16 08:53:20 +00:00
Code Issues Packages Projects Releases Wiki Activity

allow private variables to have an override-expression initializer

Browse Source
This commit is contained in:
teoxoy 2024-03-06 12:23:16 +01:00 committed by Teodor Tanasoaia
parent 3abdfde0ba
commit ca252b9e74
9 changed files with 199 additions and 23 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
naga/src/front/wgsl/lower/mod.rs
View File

@ -916,7 +916,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
let init;
if let Some(init_ast) = v.init {
let mut ectx = ctx.as_const();
let mut ectx = ctx.as_override();
let lowered = self.expression_for_abstract(init_ast, &mut ectx)?;
let ty_res = crate::proc::TypeResolution::Handle(ty);
let converted = ectx

4
naga/src/valid/interface.rs
View File

@ -31,7 +31,7 @@ pub enum GlobalVariableError {
Handle<crate::Type>,
#[source] Disalignment,
),
#[error("Initializer must be a const-expression")]
#[error("Initializer must be an override-expression")]
InitializerExprType,
#[error("Initializer doesn't match the variable type")]
InitializerType,
@ -529,7 +529,7 @@ impl super::Validator {
}
}
if !global_expr_kind.is_const(init) {
if !global_expr_kind.is_const_or_override(init) {
return Err(GlobalVariableError::InitializerExprType);
}

4
naga/tests/in/overrides.wgsl
View File

@ -13,9 +13,13 @@
override inferred_f32 = 2.718;
var<private> gain_x_10: f32 = gain * 10.;
@compute @workgroup_size(1)
fn main() {
var t = height * 5;
let a = !has_point_light;
var x = a;
var gain_x_100 = gain_x_10 * 10.;
}

70
naga/tests/out/analysis/overrides.info.ron
View File

@ -14,7 +14,9 @@
),
may_kill: false,
sampling_set: [],
global_uses: [],
global_uses: [
("READ"),
],
expressions: [
(
uniformity: (
@ -91,6 +93,63 @@
space: Function,
)),
),
(
uniformity: (
non_uniform_result: Some(8),
requirements: (""),
),
ref_count: 1,
assignable_global: Some(1),
ty: Value(Pointer(
base: 2,
space: Private,
)),
),
(
uniformity: (
non_uniform_result: Some(8),
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Handle(2),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Scalar((
kind: Float,
width: 4,
))),
),
(
uniformity: (
non_uniform_result: Some(8),
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Scalar((
kind: Float,
width: 4,
))),
),
(
uniformity: (
non_uniform_result: Some(12),
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Pointer(
base: 2,
space: Function,
)),
),
],
sampling: [],
dual_source_blending: false,
@ -119,5 +178,14 @@
kind: Float,
width: 4,
))),
Handle(2),
Value(Scalar((
kind: Float,
width: 4,
))),
Value(Scalar((
kind: Float,
width: 4,
))),
],
)

5
naga/tests/out/hlsl/overrides.hlsl
View File

@ -6,13 +6,18 @@ static const float depth = 2.3;
static const float height = 4.6;
static const float inferred_f32_ = 2.718;
static float gain_x_10_ = 11.0;
[numthreads(1, 1, 1)]
void main()
{
float t = (float)0;
bool x = (bool)0;
float gain_x_100_ = (float)0;
t = 23.0;
x = true;
float _expr10 = gain_x_10_;
gain_x_100_ = (_expr10 * 10.0);
return;
}

45
naga/tests/out/ir/overrides.compact.ron
View File

@ -65,7 +65,15 @@
init: Some(7),
),
],
global_variables: [],
global_variables: [
(
name: Some("gain_x_10"),
space: Private,
binding: None,
ty: 2,
init: Some(10),
),
],
global_expressions: [
Literal(Bool(true)),
Literal(F32(2.3)),
@ -78,6 +86,13 @@
right: 4,
),
Literal(F32(2.718)),
Override(3),
Literal(F32(10.0)),
Binary(
op: Multiply,
left: 8,
right: 9,
),
],
functions: [],
entry_points: [
@ -101,6 +116,11 @@
ty: 1,
init: None,
),
(
name: Some("gain_x_100"),
ty: 2,
init: None,
),
],
expressions: [
Override(6),
@ -117,6 +137,17 @@
expr: 5,
),
LocalVariable(2),
GlobalVariable(1),
Load(
pointer: 8,
),
Literal(F32(10.0)),
Binary(
op: Multiply,
left: 9,
right: 10,
),
LocalVariable(3),
],
named_expressions: {
6: "a",
@ -138,6 +169,18 @@
pointer: 7,
value: 6,
),
Emit((
start: 8,
end: 9,
)),
Emit((
start: 10,
end: 11,
)),
Store(
pointer: 12,
value: 11,
),
Return(
value: None,
),

45
naga/tests/out/ir/overrides.ron
View File

@ -65,7 +65,15 @@
init: Some(7),
),
],
global_variables: [],
global_variables: [
(
name: Some("gain_x_10"),
space: Private,
binding: None,
ty: 2,
init: Some(10),
),
],
global_expressions: [
Literal(Bool(true)),
Literal(F32(2.3)),
@ -78,6 +86,13 @@
right: 4,
),
Literal(F32(2.718)),
Override(3),
Literal(F32(10.0)),
Binary(
op: Multiply,
left: 8,
right: 9,
),
],
functions: [],
entry_points: [
@ -101,6 +116,11 @@
ty: 1,
init: None,
),
(
name: Some("gain_x_100"),
ty: 2,
init: None,
),
],
expressions: [
Override(6),
@ -117,6 +137,17 @@
expr: 5,
),
LocalVariable(2),
GlobalVariable(1),
Load(
pointer: 8,
),
Literal(F32(10.0)),
Binary(
op: Multiply,
left: 9,
right: 10,
),
LocalVariable(3),
],
named_expressions: {
6: "a",
@ -138,6 +169,18 @@
pointer: 7,
value: 6,
),
Emit((
start: 8,
end: 9,
)),
Emit((
start: 10,
end: 11,
)),
Store(
pointer: 12,
value: 11,
),
Return(
value: None,
),

4
naga/tests/out/msl/overrides.msl
View File

@ -14,9 +14,13 @@ constant float inferred_f32_ = 2.718;
kernel void main_(
) {
float gain_x_10_ = 11.0;
float t = {};
bool x = {};
float gain_x_100_ = {};
t = 23.0;
x = true;
float _e10 = gain_x_10_;
gain_x_100_ = _e10 * 10.0;
return;
}

43
naga/tests/out/spv/overrides.main.spvasm
View File

@ -1,12 +1,12 @@
; SPIR-V
; Version: 1.0
; Generator: rspirv
; Bound: 24
; Bound: 32
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %14 "main"
OpExecutionMode %14 LocalSize 1 1 1
OpEntryPoint GLCompute %18 "main"
OpExecutionMode %18 LocalSize 1 1 1
%2 = OpTypeVoid
%3 = OpTypeBool
%4 = OpTypeFloat 32
@ -18,19 +18,28 @@ OpExecutionMode %14 LocalSize 1 1 1
%10 = OpConstant %4 2.0
%11 = OpConstant %4 4.6
%12 = OpConstant %4 2.718
%15 = OpTypeFunction %2
%16 = OpConstant %4 23.0
%18 = OpTypePointer Function %4
%19 = OpConstantNull %4
%21 = OpTypePointer Function %3
%22 = OpConstantNull %3
%14 = OpFunction %2 None %15
%13 = OpLabel
%17 = OpVariable %18 Function %19
%20 = OpVariable %21 Function %22
OpBranch %23
%23 = OpLabel
OpStore %17 %16
OpStore %20 %5
%13 = OpConstant %4 10.0
%14 = OpConstant %4 11.0
%16 = OpTypePointer Private %4
%15 = OpVariable %16 Private %14
%19 = OpTypeFunction %2
%20 = OpConstant %4 23.0
%22 = OpTypePointer Function %4
%23 = OpConstantNull %4
%25 = OpTypePointer Function %3
%26 = OpConstantNull %3
%28 = OpConstantNull %4
%18 = OpFunction %2 None %19
%17 = OpLabel
%21 = OpVariable %22 Function %23
%24 = OpVariable %25 Function %26
%27 = OpVariable %22 Function %28
OpBranch %29
%29 = OpLabel
OpStore %21 %20
OpStore %24 %5
%30 = OpLoad %4 %15
%31 = OpFMul %4 %30 %13
OpStore %27 %31
OpReturn
OpFunctionEnd