From 6c4effe654459e61792558d38d952d26c6abf8fa Mon Sep 17 00:00:00 2001 From: Ben Clayton Date: Tue, 26 Oct 2021 10:05:02 +0000 Subject: [PATCH] writer/spirv: Fix [[override]] for zero-init defaults These expressions were not emitted as spec-ops, failing new CTS tests. Change-Id: I56e8f56a22de2b8dac9a8bd7a2d694d8d81dca35 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/67480 Kokoro: Kokoro Commit-Queue: Ben Clayton Reviewed-by: James Price --- src/writer/spirv/builder.cc | 50 +++++++++++----- src/writer/spirv/builder.h | 13 ++--- .../builder_constructor_expression_test.cc | 58 +++++++++---------- .../spirv/builder_global_variable_test.cc | 36 ++++++++++++ src/writer/spirv/scalar_constant.h | 19 ++++++ test/var/override/named/no_init/bool.wgsl | 6 ++ .../named/no_init/bool.wgsl.expected.hlsl | 16 +++++ .../named/no_init/bool.wgsl.expected.msl | 8 +++ .../named/no_init/bool.wgsl.expected.spvasm | 20 +++++++ .../named/no_init/bool.wgsl.expected.wgsl | 6 ++ test/var/override/named/no_init/f32.wgsl | 6 ++ .../named/no_init/f32.wgsl.expected.hlsl | 16 +++++ .../named/no_init/f32.wgsl.expected.msl | 8 +++ .../named/no_init/f32.wgsl.expected.spvasm | 20 +++++++ .../named/no_init/f32.wgsl.expected.wgsl | 6 ++ test/var/override/named/no_init/i32.wgsl | 6 ++ .../named/no_init/i32.wgsl.expected.hlsl | 16 +++++ .../named/no_init/i32.wgsl.expected.msl | 8 +++ .../named/no_init/i32.wgsl.expected.spvasm | 20 +++++++ .../named/no_init/i32.wgsl.expected.wgsl | 6 ++ test/var/override/named/no_init/u32.wgsl | 6 ++ .../named/no_init/u32.wgsl.expected.hlsl | 16 +++++ .../named/no_init/u32.wgsl.expected.msl | 8 +++ .../named/no_init/u32.wgsl.expected.spvasm | 20 +++++++ .../named/no_init/u32.wgsl.expected.wgsl | 6 ++ test/var/override/named/val_init/bool.wgsl | 6 ++ .../named/val_init/bool.wgsl.expected.hlsl | 9 +++ .../named/val_init/bool.wgsl.expected.msl | 8 +++ .../named/val_init/bool.wgsl.expected.spvasm | 20 +++++++ .../named/val_init/bool.wgsl.expected.wgsl | 6 ++ test/var/override/named/val_init/f32.wgsl | 6 ++ .../named/val_init/f32.wgsl.expected.hlsl | 9 +++ .../named/val_init/f32.wgsl.expected.msl | 8 +++ .../named/val_init/f32.wgsl.expected.spvasm | 20 +++++++ .../named/val_init/f32.wgsl.expected.wgsl | 6 ++ test/var/override/named/val_init/i32.wgsl | 6 ++ .../named/val_init/i32.wgsl.expected.hlsl | 9 +++ .../named/val_init/i32.wgsl.expected.msl | 8 +++ .../named/val_init/i32.wgsl.expected.spvasm | 20 +++++++ .../named/val_init/i32.wgsl.expected.wgsl | 6 ++ test/var/override/named/val_init/u32.wgsl | 6 ++ .../named/val_init/u32.wgsl.expected.hlsl | 9 +++ .../named/val_init/u32.wgsl.expected.msl | 8 +++ .../named/val_init/u32.wgsl.expected.spvasm | 20 +++++++ .../named/val_init/u32.wgsl.expected.wgsl | 6 ++ test/var/override/named/zero_init/bool.wgsl | 6 ++ .../named/zero_init/bool.wgsl.expected.hlsl | 9 +++ .../named/zero_init/bool.wgsl.expected.msl | 8 +++ .../named/zero_init/bool.wgsl.expected.spvasm | 20 +++++++ .../named/zero_init/bool.wgsl.expected.wgsl | 6 ++ test/var/override/named/zero_init/f32.wgsl | 6 ++ .../named/zero_init/f32.wgsl.expected.hlsl | 9 +++ .../named/zero_init/f32.wgsl.expected.msl | 8 +++ .../named/zero_init/f32.wgsl.expected.spvasm | 20 +++++++ .../named/zero_init/f32.wgsl.expected.wgsl | 6 ++ test/var/override/named/zero_init/i32.wgsl | 6 ++ .../named/zero_init/i32.wgsl.expected.hlsl | 9 +++ .../named/zero_init/i32.wgsl.expected.msl | 8 +++ .../named/zero_init/i32.wgsl.expected.spvasm | 20 +++++++ .../named/zero_init/i32.wgsl.expected.wgsl | 6 ++ test/var/override/named/zero_init/u32.wgsl | 6 ++ .../named/zero_init/u32.wgsl.expected.hlsl | 9 +++ .../named/zero_init/u32.wgsl.expected.msl | 8 +++ .../named/zero_init/u32.wgsl.expected.spvasm | 20 +++++++ .../named/zero_init/u32.wgsl.expected.wgsl | 6 ++ test/var/override/numbered/no_init/bool.wgsl | 6 ++ .../numbered/no_init/bool.wgsl.expected.hlsl | 16 +++++ .../numbered/no_init/bool.wgsl.expected.msl | 8 +++ .../no_init/bool.wgsl.expected.spvasm | 20 +++++++ .../numbered/no_init/bool.wgsl.expected.wgsl | 6 ++ test/var/override/numbered/no_init/f32.wgsl | 6 ++ .../numbered/no_init/f32.wgsl.expected.hlsl | 16 +++++ .../numbered/no_init/f32.wgsl.expected.msl | 8 +++ .../numbered/no_init/f32.wgsl.expected.spvasm | 20 +++++++ .../numbered/no_init/f32.wgsl.expected.wgsl | 6 ++ test/var/override/numbered/no_init/i32.wgsl | 6 ++ .../numbered/no_init/i32.wgsl.expected.hlsl | 16 +++++ .../numbered/no_init/i32.wgsl.expected.msl | 8 +++ .../numbered/no_init/i32.wgsl.expected.spvasm | 20 +++++++ .../numbered/no_init/i32.wgsl.expected.wgsl | 6 ++ test/var/override/numbered/no_init/u32.wgsl | 6 ++ .../numbered/no_init/u32.wgsl.expected.hlsl | 16 +++++ .../numbered/no_init/u32.wgsl.expected.msl | 8 +++ .../numbered/no_init/u32.wgsl.expected.spvasm | 20 +++++++ .../numbered/no_init/u32.wgsl.expected.wgsl | 6 ++ test/var/override/numbered/val_init/bool.wgsl | 6 ++ .../numbered/val_init/bool.wgsl.expected.hlsl | 9 +++ .../numbered/val_init/bool.wgsl.expected.msl | 8 +++ .../val_init/bool.wgsl.expected.spvasm | 20 +++++++ .../numbered/val_init/bool.wgsl.expected.wgsl | 6 ++ test/var/override/numbered/val_init/f32.wgsl | 6 ++ .../numbered/val_init/f32.wgsl.expected.hlsl | 9 +++ .../numbered/val_init/f32.wgsl.expected.msl | 8 +++ .../val_init/f32.wgsl.expected.spvasm | 20 +++++++ .../numbered/val_init/f32.wgsl.expected.wgsl | 6 ++ test/var/override/numbered/val_init/i32.wgsl | 6 ++ .../numbered/val_init/i32.wgsl.expected.hlsl | 9 +++ .../numbered/val_init/i32.wgsl.expected.msl | 8 +++ .../val_init/i32.wgsl.expected.spvasm | 20 +++++++ .../numbered/val_init/i32.wgsl.expected.wgsl | 6 ++ test/var/override/numbered/val_init/u32.wgsl | 6 ++ .../numbered/val_init/u32.wgsl.expected.hlsl | 9 +++ .../numbered/val_init/u32.wgsl.expected.msl | 8 +++ .../val_init/u32.wgsl.expected.spvasm | 20 +++++++ .../numbered/val_init/u32.wgsl.expected.wgsl | 6 ++ .../var/override/numbered/zero_init/bool.wgsl | 6 ++ .../zero_init/bool.wgsl.expected.hlsl | 9 +++ .../numbered/zero_init/bool.wgsl.expected.msl | 8 +++ .../zero_init/bool.wgsl.expected.spvasm | 20 +++++++ .../zero_init/bool.wgsl.expected.wgsl | 6 ++ test/var/override/numbered/zero_init/f32.wgsl | 6 ++ .../numbered/zero_init/f32.wgsl.expected.hlsl | 9 +++ .../numbered/zero_init/f32.wgsl.expected.msl | 8 +++ .../zero_init/f32.wgsl.expected.spvasm | 20 +++++++ .../numbered/zero_init/f32.wgsl.expected.wgsl | 6 ++ test/var/override/numbered/zero_init/i32.wgsl | 6 ++ .../numbered/zero_init/i32.wgsl.expected.hlsl | 9 +++ .../numbered/zero_init/i32.wgsl.expected.msl | 8 +++ .../zero_init/i32.wgsl.expected.spvasm | 20 +++++++ .../numbered/zero_init/i32.wgsl.expected.wgsl | 6 ++ test/var/override/numbered/zero_init/u32.wgsl | 6 ++ .../numbered/zero_init/u32.wgsl.expected.hlsl | 9 +++ .../numbered/zero_init/u32.wgsl.expected.msl | 8 +++ .../zero_init/u32.wgsl.expected.spvasm | 20 +++++++ .../numbered/zero_init/u32.wgsl.expected.wgsl | 6 ++ 125 files changed, 1356 insertions(+), 52 deletions(-) create mode 100644 test/var/override/named/no_init/bool.wgsl create mode 100644 test/var/override/named/no_init/bool.wgsl.expected.hlsl create mode 100644 test/var/override/named/no_init/bool.wgsl.expected.msl create mode 100644 test/var/override/named/no_init/bool.wgsl.expected.spvasm create mode 100644 test/var/override/named/no_init/bool.wgsl.expected.wgsl create mode 100644 test/var/override/named/no_init/f32.wgsl create mode 100644 test/var/override/named/no_init/f32.wgsl.expected.hlsl create mode 100644 test/var/override/named/no_init/f32.wgsl.expected.msl create mode 100644 test/var/override/named/no_init/f32.wgsl.expected.spvasm create mode 100644 test/var/override/named/no_init/f32.wgsl.expected.wgsl create mode 100644 test/var/override/named/no_init/i32.wgsl create mode 100644 test/var/override/named/no_init/i32.wgsl.expected.hlsl create mode 100644 test/var/override/named/no_init/i32.wgsl.expected.msl create mode 100644 test/var/override/named/no_init/i32.wgsl.expected.spvasm create mode 100644 test/var/override/named/no_init/i32.wgsl.expected.wgsl create mode 100644 test/var/override/named/no_init/u32.wgsl create mode 100644 test/var/override/named/no_init/u32.wgsl.expected.hlsl create mode 100644 test/var/override/named/no_init/u32.wgsl.expected.msl create mode 100644 test/var/override/named/no_init/u32.wgsl.expected.spvasm create mode 100644 test/var/override/named/no_init/u32.wgsl.expected.wgsl create mode 100644 test/var/override/named/val_init/bool.wgsl create mode 100644 test/var/override/named/val_init/bool.wgsl.expected.hlsl create mode 100644 test/var/override/named/val_init/bool.wgsl.expected.msl create mode 100644 test/var/override/named/val_init/bool.wgsl.expected.spvasm create mode 100644 test/var/override/named/val_init/bool.wgsl.expected.wgsl create mode 100644 test/var/override/named/val_init/f32.wgsl create mode 100644 test/var/override/named/val_init/f32.wgsl.expected.hlsl create mode 100644 test/var/override/named/val_init/f32.wgsl.expected.msl create mode 100644 test/var/override/named/val_init/f32.wgsl.expected.spvasm create mode 100644 test/var/override/named/val_init/f32.wgsl.expected.wgsl create mode 100644 test/var/override/named/val_init/i32.wgsl create mode 100644 test/var/override/named/val_init/i32.wgsl.expected.hlsl create mode 100644 test/var/override/named/val_init/i32.wgsl.expected.msl create mode 100644 test/var/override/named/val_init/i32.wgsl.expected.spvasm create mode 100644 test/var/override/named/val_init/i32.wgsl.expected.wgsl create mode 100644 test/var/override/named/val_init/u32.wgsl create mode 100644 test/var/override/named/val_init/u32.wgsl.expected.hlsl create mode 100644 test/var/override/named/val_init/u32.wgsl.expected.msl create mode 100644 test/var/override/named/val_init/u32.wgsl.expected.spvasm create mode 100644 test/var/override/named/val_init/u32.wgsl.expected.wgsl create mode 100644 test/var/override/named/zero_init/bool.wgsl create mode 100644 test/var/override/named/zero_init/bool.wgsl.expected.hlsl create mode 100644 test/var/override/named/zero_init/bool.wgsl.expected.msl create mode 100644 test/var/override/named/zero_init/bool.wgsl.expected.spvasm create mode 100644 test/var/override/named/zero_init/bool.wgsl.expected.wgsl create mode 100644 test/var/override/named/zero_init/f32.wgsl create mode 100644 test/var/override/named/zero_init/f32.wgsl.expected.hlsl create mode 100644 test/var/override/named/zero_init/f32.wgsl.expected.msl create mode 100644 test/var/override/named/zero_init/f32.wgsl.expected.spvasm create mode 100644 test/var/override/named/zero_init/f32.wgsl.expected.wgsl create mode 100644 test/var/override/named/zero_init/i32.wgsl create mode 100644 test/var/override/named/zero_init/i32.wgsl.expected.hlsl create mode 100644 test/var/override/named/zero_init/i32.wgsl.expected.msl create mode 100644 test/var/override/named/zero_init/i32.wgsl.expected.spvasm create mode 100644 test/var/override/named/zero_init/i32.wgsl.expected.wgsl create mode 100644 test/var/override/named/zero_init/u32.wgsl create mode 100644 test/var/override/named/zero_init/u32.wgsl.expected.hlsl create mode 100644 test/var/override/named/zero_init/u32.wgsl.expected.msl create mode 100644 test/var/override/named/zero_init/u32.wgsl.expected.spvasm create mode 100644 test/var/override/named/zero_init/u32.wgsl.expected.wgsl create mode 100644 test/var/override/numbered/no_init/bool.wgsl create mode 100644 test/var/override/numbered/no_init/bool.wgsl.expected.hlsl create mode 100644 test/var/override/numbered/no_init/bool.wgsl.expected.msl create mode 100644 test/var/override/numbered/no_init/bool.wgsl.expected.spvasm create mode 100644 test/var/override/numbered/no_init/bool.wgsl.expected.wgsl create mode 100644 test/var/override/numbered/no_init/f32.wgsl create mode 100644 test/var/override/numbered/no_init/f32.wgsl.expected.hlsl create mode 100644 test/var/override/numbered/no_init/f32.wgsl.expected.msl create mode 100644 test/var/override/numbered/no_init/f32.wgsl.expected.spvasm create mode 100644 test/var/override/numbered/no_init/f32.wgsl.expected.wgsl create mode 100644 test/var/override/numbered/no_init/i32.wgsl create mode 100644 test/var/override/numbered/no_init/i32.wgsl.expected.hlsl create mode 100644 test/var/override/numbered/no_init/i32.wgsl.expected.msl create mode 100644 test/var/override/numbered/no_init/i32.wgsl.expected.spvasm create mode 100644 test/var/override/numbered/no_init/i32.wgsl.expected.wgsl create mode 100644 test/var/override/numbered/no_init/u32.wgsl create mode 100644 test/var/override/numbered/no_init/u32.wgsl.expected.hlsl create mode 100644 test/var/override/numbered/no_init/u32.wgsl.expected.msl create mode 100644 test/var/override/numbered/no_init/u32.wgsl.expected.spvasm create mode 100644 test/var/override/numbered/no_init/u32.wgsl.expected.wgsl create mode 100644 test/var/override/numbered/val_init/bool.wgsl create mode 100644 test/var/override/numbered/val_init/bool.wgsl.expected.hlsl create mode 100644 test/var/override/numbered/val_init/bool.wgsl.expected.msl create mode 100644 test/var/override/numbered/val_init/bool.wgsl.expected.spvasm create mode 100644 test/var/override/numbered/val_init/bool.wgsl.expected.wgsl create mode 100644 test/var/override/numbered/val_init/f32.wgsl create mode 100644 test/var/override/numbered/val_init/f32.wgsl.expected.hlsl create mode 100644 test/var/override/numbered/val_init/f32.wgsl.expected.msl create mode 100644 test/var/override/numbered/val_init/f32.wgsl.expected.spvasm create mode 100644 test/var/override/numbered/val_init/f32.wgsl.expected.wgsl create mode 100644 test/var/override/numbered/val_init/i32.wgsl create mode 100644 test/var/override/numbered/val_init/i32.wgsl.expected.hlsl create mode 100644 test/var/override/numbered/val_init/i32.wgsl.expected.msl create mode 100644 test/var/override/numbered/val_init/i32.wgsl.expected.spvasm create mode 100644 test/var/override/numbered/val_init/i32.wgsl.expected.wgsl create mode 100644 test/var/override/numbered/val_init/u32.wgsl create mode 100644 test/var/override/numbered/val_init/u32.wgsl.expected.hlsl create mode 100644 test/var/override/numbered/val_init/u32.wgsl.expected.msl create mode 100644 test/var/override/numbered/val_init/u32.wgsl.expected.spvasm create mode 100644 test/var/override/numbered/val_init/u32.wgsl.expected.wgsl create mode 100644 test/var/override/numbered/zero_init/bool.wgsl create mode 100644 test/var/override/numbered/zero_init/bool.wgsl.expected.hlsl create mode 100644 test/var/override/numbered/zero_init/bool.wgsl.expected.msl create mode 100644 test/var/override/numbered/zero_init/bool.wgsl.expected.spvasm create mode 100644 test/var/override/numbered/zero_init/bool.wgsl.expected.wgsl create mode 100644 test/var/override/numbered/zero_init/f32.wgsl create mode 100644 test/var/override/numbered/zero_init/f32.wgsl.expected.hlsl create mode 100644 test/var/override/numbered/zero_init/f32.wgsl.expected.msl create mode 100644 test/var/override/numbered/zero_init/f32.wgsl.expected.spvasm create mode 100644 test/var/override/numbered/zero_init/f32.wgsl.expected.wgsl create mode 100644 test/var/override/numbered/zero_init/i32.wgsl create mode 100644 test/var/override/numbered/zero_init/i32.wgsl.expected.hlsl create mode 100644 test/var/override/numbered/zero_init/i32.wgsl.expected.msl create mode 100644 test/var/override/numbered/zero_init/i32.wgsl.expected.spvasm create mode 100644 test/var/override/numbered/zero_init/i32.wgsl.expected.wgsl create mode 100644 test/var/override/numbered/zero_init/u32.wgsl create mode 100644 test/var/override/numbered/zero_init/u32.wgsl.expected.hlsl create mode 100644 test/var/override/numbered/zero_init/u32.wgsl.expected.msl create mode 100644 test/var/override/numbered/zero_init/u32.wgsl.expected.spvasm create mode 100644 test/var/override/numbered/zero_init/u32.wgsl.expected.wgsl diff --git a/src/writer/spirv/builder.cc b/src/writer/spirv/builder.cc index 5481a71039..1ca479f9ba 100644 --- a/src/writer/spirv/builder.cc +++ b/src/writer/spirv/builder.cc @@ -577,7 +577,7 @@ uint32_t Builder::GenerateExpression(const ast::Expression* expr) { return GenerateCallExpression(c); } if (auto* c = expr->As()) { - return GenerateConstructorExpression(nullptr, c, false); + return GenerateConstructorExpression(nullptr, c); } if (auto* i = expr->As()) { return GenerateIdentifierExpression(i); @@ -763,7 +763,7 @@ bool Builder::GenerateGlobalVariable(const ast::Variable* var) { } init_id = GenerateConstructorExpression( - var, var->constructor->As(), true); + var, var->constructor->As()); if (init_id == 0) { return false; } @@ -1268,13 +1268,12 @@ uint32_t Builder::GetGLSLstd450Import() { uint32_t Builder::GenerateConstructorExpression( const ast::Variable* var, - const ast::ConstructorExpression* expr, - bool is_global_init) { + const ast::ConstructorExpression* expr) { if (auto* scalar = expr->As()) { return GenerateLiteralIfNeeded(var, scalar->literal); } if (auto* type = expr->As()) { - return GenerateTypeConstructorExpression(type, is_global_init); + return GenerateTypeConstructorExpression(var, type); } error_ = "unknown constructor expression"; @@ -1338,14 +1337,35 @@ bool Builder::is_constructor_const(const ast::Expression* expr, } uint32_t Builder::GenerateTypeConstructorExpression( - const ast::TypeConstructorExpression* init, - bool is_global_init) { + const ast::Variable* var, + const ast::TypeConstructorExpression* init) { + auto* global_var = builder_.Sem().Get(var); + auto& values = init->values; auto* result_type = TypeOf(init); // Generate the zero initializer if there are no values provided. if (values.empty()) { + if (global_var && global_var->IsPipelineConstant()) { + auto constant_id = global_var->ConstantId(); + if (result_type->Is()) { + return GenerateConstantIfNeeded( + ScalarConstant::I32(0).AsSpecOp(constant_id)); + } + if (result_type->Is()) { + return GenerateConstantIfNeeded( + ScalarConstant::U32(0).AsSpecOp(constant_id)); + } + if (result_type->Is()) { + return GenerateConstantIfNeeded( + ScalarConstant::F32(0).AsSpecOp(constant_id)); + } + if (result_type->Is()) { + return GenerateConstantIfNeeded( + ScalarConstant::Bool(false).AsSpecOp(constant_id)); + } + } return GenerateConstantNullIfNeeded(result_type->UnwrapRef()); } @@ -1353,7 +1373,7 @@ uint32_t Builder::GenerateTypeConstructorExpression( out << "__const_" << init->type->FriendlyName(builder_.Symbols()) << "_"; result_type = result_type->UnwrapRef(); - bool constructor_is_const = is_constructor_const(init, is_global_init); + bool constructor_is_const = is_constructor_const(init, global_var); if (has_error()) { return 0; } @@ -1372,8 +1392,7 @@ uint32_t Builder::GenerateTypeConstructorExpression( } if (can_cast_or_copy) { - return GenerateCastOrCopyOrPassthrough(result_type, values[0], - is_global_init); + return GenerateCastOrCopyOrPassthrough(result_type, values[0], global_var); } auto type_id = GenerateTypeIfNeeded(result_type); @@ -1392,8 +1411,8 @@ uint32_t Builder::GenerateTypeConstructorExpression( for (auto* e : values) { uint32_t id = 0; if (constructor_is_const) { - id = GenerateConstructorExpression( - nullptr, e->As(), is_global_init); + id = GenerateConstructorExpression(nullptr, + e->As()); } else { id = GenerateExpression(e); id = GenerateLoadIfNeeded(TypeOf(e), id); @@ -1417,8 +1436,7 @@ uint32_t Builder::GenerateTypeConstructorExpression( // Both scalars, but not the same type so we need to generate a conversion // of the value. if (value_type->is_scalar() && result_type->is_scalar()) { - id = GenerateCastOrCopyOrPassthrough(result_type, values[0], - is_global_init); + id = GenerateCastOrCopyOrPassthrough(result_type, values[0], global_var); out << "_" << id; ops.push_back(Operand::Int(id)); continue; @@ -1445,7 +1463,7 @@ uint32_t Builder::GenerateTypeConstructorExpression( auto extract = result_op(); auto extract_id = extract.to_i(); - if (!is_global_init) { + if (!global_var) { // A non-global initializer. Case 2. if (!push_function_inst(spv::Op::OpCompositeExtract, {Operand::Int(value_type_id), extract, @@ -2816,7 +2834,7 @@ bool Builder::GenerateTextureIntrinsic(const ast::CallExpression* call, if (auto* array_index = arg(Usage::kArrayIndex)) { // Array index needs to be appended to the coordinates. auto* packed = AppendVector(&builder_, arg(Usage::kCoords), array_index); - auto param = GenerateTypeConstructorExpression(packed, false); + auto param = GenerateTypeConstructorExpression(nullptr, packed); if (param == 0) { return false; } diff --git a/src/writer/spirv/builder.h b/src/writer/spirv/builder.h index 0aad2ddc9a..77d878f980 100644 --- a/src/writer/spirv/builder.h +++ b/src/writer/spirv/builder.h @@ -338,18 +338,17 @@ class Builder { /// Generates a constructor expression /// @param var the variable generated for, nullptr if no variable associated. /// @param expr the expression to generate - /// @param is_global_init set true if this is a global variable constructor /// @returns the ID of the expression or 0 on failure. - uint32_t GenerateConstructorExpression(const ast::Variable* var, - const ast::ConstructorExpression* expr, - bool is_global_init); + uint32_t GenerateConstructorExpression( + const ast::Variable* var, + const ast::ConstructorExpression* expr); /// Generates a type constructor expression + /// @param var the variable generated for, nullptr if no variable associated. /// @param init the expression to generate - /// @param is_global_init set true if this is a global variable constructor /// @returns the ID of the expression or 0 on failure. uint32_t GenerateTypeConstructorExpression( - const ast::TypeConstructorExpression* init, - bool is_global_init); + const ast::Variable* var, + const ast::TypeConstructorExpression* init); /// Generates a literal constant if needed /// @param var the variable generated for, nullptr if no variable associated. /// @param lit the literal to generate diff --git a/src/writer/spirv/builder_constructor_expression_test.cc b/src/writer/spirv/builder_constructor_expression_test.cc index 0f65afef28..3488b2b91f 100644 --- a/src/writer/spirv/builder_constructor_expression_test.cc +++ b/src/writer/spirv/builder_constructor_expression_test.cc @@ -24,11 +24,11 @@ using SpvBuilderConstructorTest = TestHelper; TEST_F(SpvBuilderConstructorTest, Const) { auto* c = Expr(42.2f); - WrapInFunction(c); + auto* g = Global("g", ty.f32(), c, ast::StorageClass::kPrivate); spirv::Builder& b = Build(); - EXPECT_EQ(b.GenerateConstructorExpression(nullptr, c, true), 2u); + EXPECT_EQ(b.GenerateConstructorExpression(g, c), 2u); ASSERT_FALSE(b.has_error()) << b.error(); EXPECT_EQ(DumpInstructions(b.types()), R"(%1 = OpTypeFloat 32 @@ -55,7 +55,7 @@ TEST_F(SpvBuilderConstructorTest, Type) { spirv::Builder& b = Build(); - EXPECT_EQ(b.GenerateConstructorExpression(nullptr, t, true), 5u); + EXPECT_EQ(b.GenerateConstructorExpression(nullptr, t), 5u); ASSERT_FALSE(b.has_error()) << b.error(); EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeFloat 32 @@ -170,11 +170,11 @@ TEST_F(SpvBuilderConstructorTest, Type_NonConst_Value_Fails) { Expr(3.0f)); auto* t = vec2(1.0f, rel); - WrapInFunction(t); + auto* g = Global("g", ty.vec2(), t, ast::StorageClass::kPrivate); spirv::Builder& b = Build(); - EXPECT_EQ(b.GenerateConstructorExpression(nullptr, t, true), 0u); + EXPECT_EQ(b.GenerateConstructorExpression(g, t), 0u); EXPECT_TRUE(b.has_error()); EXPECT_EQ(b.error(), R"(constructor must be a constant expression)"); } @@ -670,12 +670,12 @@ TEST_F(SpvBuilderConstructorTest, Type_Vec4_With_Vec4) { TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec2_With_F32) { auto* cast = vec2(2.0f); - WrapInFunction(cast); + auto* g = Global("g", ty.vec2(), cast, ast::StorageClass::kPrivate); spirv::Builder& b = Build(); b.push_function(Function{}); - EXPECT_EQ(b.GenerateConstructorExpression(nullptr, cast, true), 4u); + EXPECT_EQ(b.GenerateConstructorExpression(g, cast), 4u); EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeFloat 32 %1 = OpTypeVector %2 2 @@ -740,12 +740,12 @@ TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec4_With_Vec4) { TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec3_With_F32) { auto* cast = vec3(2.0f); - WrapInFunction(cast); + auto* g = Global("g", ty.vec3(), cast, ast::StorageClass::kPrivate); spirv::Builder& b = Build(); b.push_function(Function{}); - EXPECT_EQ(b.GenerateConstructorExpression(nullptr, cast, true), 4u); + EXPECT_EQ(b.GenerateConstructorExpression(g, cast), 4u); EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeFloat 32 %1 = OpTypeVector %2 3 @@ -756,12 +756,12 @@ TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec3_With_F32) { TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec3_With_F32_Vec2) { auto* cast = vec3(2.0f, vec2(2.0f, 2.0f)); - WrapInFunction(cast); + auto* g = Global("g", ty.vec3(), cast, ast::StorageClass::kPrivate); spirv::Builder& b = Build(); b.push_function(Function{}); - EXPECT_EQ(b.GenerateConstructorExpression(nullptr, cast, true), 11u); + EXPECT_EQ(b.GenerateConstructorExpression(g, cast), 11u); EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeFloat 32 %1 = OpTypeVector %2 3 @@ -779,12 +779,12 @@ TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec3_With_F32_Vec2) { TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec3_With_Vec2_F32) { auto* cast = vec3(vec2(2.0f, 2.0f), 2.0f); - WrapInFunction(cast); + auto* g = Global("g", ty.vec3(), cast, ast::StorageClass::kPrivate); spirv::Builder& b = Build(); b.push_function(Function{}); - EXPECT_EQ(b.GenerateConstructorExpression(nullptr, cast, true), 11u); + EXPECT_EQ(b.GenerateConstructorExpression(g, cast), 11u); EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeFloat 32 %1 = OpTypeVector %2 3 @@ -802,12 +802,12 @@ TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec3_With_Vec2_F32) { TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec4_With_F32) { auto* cast = vec4(2.0f); - WrapInFunction(cast); + auto* g = Global("g", ty.vec4(), cast, ast::StorageClass::kPrivate); spirv::Builder& b = Build(); b.push_function(Function{}); - EXPECT_EQ(b.GenerateConstructorExpression(nullptr, cast, true), 4u); + EXPECT_EQ(b.GenerateConstructorExpression(g, cast), 4u); EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeFloat 32 %1 = OpTypeVector %2 4 @@ -818,12 +818,12 @@ TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec4_With_F32) { TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec4_With_F32_F32_Vec2) { auto* cast = vec4(2.0f, 2.0f, vec2(2.0f, 2.0f)); - WrapInFunction(cast); + auto* g = Global("g", ty.vec4(), cast, ast::StorageClass::kPrivate); spirv::Builder& b = Build(); b.push_function(Function{}); - EXPECT_EQ(b.GenerateConstructorExpression(nullptr, cast, true), 11u); + EXPECT_EQ(b.GenerateConstructorExpression(g, cast), 11u); EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeFloat 32 %1 = OpTypeVector %2 4 @@ -841,12 +841,12 @@ TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec4_With_F32_F32_Vec2) { TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec4_With_F32_Vec2_F32) { auto* cast = vec4(2.0f, vec2(2.0f, 2.0f), 2.0f); - WrapInFunction(cast); + auto* g = Global("g", ty.vec4(), cast, ast::StorageClass::kPrivate); spirv::Builder& b = Build(); b.push_function(Function{}); - EXPECT_EQ(b.GenerateConstructorExpression(nullptr, cast, true), 11u); + EXPECT_EQ(b.GenerateConstructorExpression(g, cast), 11u); EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeFloat 32 %1 = OpTypeVector %2 4 @@ -864,12 +864,12 @@ TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec4_With_F32_Vec2_F32) { TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec4_With_Vec2_F32_F32) { auto* cast = vec4(vec2(2.0f, 2.0f), 2.0f, 2.0f); - WrapInFunction(cast); + auto* g = Global("g", ty.vec4(), cast, ast::StorageClass::kPrivate); spirv::Builder& b = Build(); b.push_function(Function{}); - EXPECT_EQ(b.GenerateConstructorExpression(nullptr, cast, true), 11u); + EXPECT_EQ(b.GenerateConstructorExpression(g, cast), 11u); EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeFloat 32 %1 = OpTypeVector %2 4 @@ -887,12 +887,12 @@ TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec4_With_Vec2_F32_F32) { TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec4_With_Vec2_Vec2) { auto* cast = vec4(vec2(2.0f, 2.0f), vec2(2.0f, 2.0f)); - WrapInFunction(cast); + auto* g = Global("g", ty.vec4(), cast, ast::StorageClass::kPrivate); spirv::Builder& b = Build(); b.push_function(Function{}); - EXPECT_EQ(b.GenerateConstructorExpression(nullptr, cast, true), 13u); + EXPECT_EQ(b.GenerateConstructorExpression(g, cast), 13u); EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeFloat 32 %1 = OpTypeVector %2 4 @@ -912,12 +912,12 @@ TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec4_With_Vec2_Vec2) { TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec4_With_F32_Vec3) { auto* cast = vec4(2.0f, vec3(2.0f, 2.0f, 2.0f)); - WrapInFunction(cast); + auto* g = Global("g", ty.vec4(), cast, ast::StorageClass::kPrivate); spirv::Builder& b = Build(); b.push_function(Function{}); - EXPECT_EQ(b.GenerateConstructorExpression(nullptr, cast, true), 13u); + EXPECT_EQ(b.GenerateConstructorExpression(g, cast), 13u); EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeFloat 32 %1 = OpTypeVector %2 4 @@ -937,12 +937,12 @@ TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec4_With_F32_Vec3) { TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec4_With_Vec3_F32) { auto* cast = vec4(vec3(2.0f, 2.0f, 2.0f), 2.0f); - WrapInFunction(cast); + auto* g = Global("g", ty.vec4(), cast, ast::StorageClass::kPrivate); spirv::Builder& b = Build(); b.push_function(Function{}); - EXPECT_EQ(b.GenerateConstructorExpression(nullptr, cast, true), 13u); + EXPECT_EQ(b.GenerateConstructorExpression(g, cast), 13u); EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeFloat 32 %1 = OpTypeVector %2 4 @@ -1275,7 +1275,7 @@ TEST_F(SpvBuilderConstructorTest, Type_Struct) { } TEST_F(SpvBuilderConstructorTest, Type_ZeroInit_F32) { - auto* t = Construct(ty.f32()); + auto* t = Construct(); WrapInFunction(t); @@ -1326,7 +1326,7 @@ TEST_F(SpvBuilderConstructorTest, Type_ZeroInit_U32) { } TEST_F(SpvBuilderConstructorTest, Type_ZeroInit_Bool) { - auto* t = Construct(ty.bool_()); + auto* t = Construct(); WrapInFunction(t); diff --git a/src/writer/spirv/builder_global_variable_test.cc b/src/writer/spirv/builder_global_variable_test.cc index f464bee12c..4160617106 100644 --- a/src/writer/spirv/builder_global_variable_test.cc +++ b/src/writer/spirv/builder_global_variable_test.cc @@ -167,6 +167,24 @@ TEST_F(BuilderTest, GlobalVar_Override_Bool) { )"); } +TEST_F(BuilderTest, GlobalVar_Override_Bool_ZeroValue) { + auto* v = GlobalConst("var", ty.bool_(), Construct(), + ast::DecorationList{ + create(1200), + }); + + spirv::Builder& b = Build(); + + EXPECT_TRUE(b.GenerateGlobalVariable(v)) << b.error(); + EXPECT_EQ(DumpInstructions(b.debug()), R"(OpName %2 "var" +)"); + EXPECT_EQ(DumpInstructions(b.annots()), R"(OpDecorate %2 SpecId 1200 +)"); + EXPECT_EQ(DumpInstructions(b.types()), R"(%1 = OpTypeBool +%2 = OpSpecConstantFalse %1 +)"); +} + TEST_F(BuilderTest, GlobalVar_Override_Bool_NoConstructor) { auto* v = GlobalConst("var", ty.bool_(), nullptr, ast::DecorationList{ @@ -203,6 +221,24 @@ TEST_F(BuilderTest, GlobalVar_Override_Scalar) { )"); } +TEST_F(BuilderTest, GlobalVar_Override_Scalar_ZeroValue) { + auto* v = GlobalConst("var", ty.f32(), Construct(), + ast::DecorationList{ + create(0), + }); + + spirv::Builder& b = Build(); + + EXPECT_TRUE(b.GenerateGlobalVariable(v)) << b.error(); + EXPECT_EQ(DumpInstructions(b.debug()), R"(OpName %2 "var" +)"); + EXPECT_EQ(DumpInstructions(b.annots()), R"(OpDecorate %2 SpecId 0 +)"); + EXPECT_EQ(DumpInstructions(b.types()), R"(%1 = OpTypeFloat 32 +%2 = OpSpecConstant %1 0 +)"); +} + TEST_F(BuilderTest, GlobalVar_Override_Scalar_F32_NoConstructor) { auto* v = GlobalConst("var", ty.f32(), nullptr, ast::DecorationList{ diff --git a/src/writer/spirv/scalar_constant.h b/src/writer/spirv/scalar_constant.h index 2911600673..d8deece8f6 100644 --- a/src/writer/spirv/scalar_constant.h +++ b/src/writer/spirv/scalar_constant.h @@ -83,6 +83,15 @@ struct ScalarConstant { return c; } + /// @param value the value of the constant + /// @returns a new ScalarConstant with the provided value and kind Kind::kBool + static inline ScalarConstant Bool(bool value) { + ScalarConstant c; + c.value.b = value; + c.kind = Kind::kBool; + return c; + } + /// Equality operator /// @param rhs the ScalarConstant to compare against /// @returns true if this ScalarConstant is equal to `rhs` @@ -98,6 +107,16 @@ struct ScalarConstant { return !(*this == rhs); } + /// @returns this ScalarConstant as a specialization op with the given + /// specialization constant identifier + /// @param id the constant identifier + ScalarConstant AsSpecOp(uint32_t id) const { + auto ret = *this; + ret.is_spec_op = true; + ret.constant_id = id; + return ret; + } + /// The constant value Value value; /// The constant value kind diff --git a/test/var/override/named/no_init/bool.wgsl b/test/var/override/named/no_init/bool.wgsl new file mode 100644 index 0000000000..4b58e68d40 --- /dev/null +++ b/test/var/override/named/no_init/bool.wgsl @@ -0,0 +1,6 @@ +[[override]] let o : bool; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/named/no_init/bool.wgsl.expected.hlsl b/test/var/override/named/no_init/bool.wgsl.expected.hlsl new file mode 100644 index 0000000000..977045cb2b --- /dev/null +++ b/test/var/override/named/no_init/bool.wgsl.expected.hlsl @@ -0,0 +1,16 @@ +SKIP: FAILED + +#ifndef WGSL_SPEC_CONSTANT_0 +#error spec constant required for constant id 0 +#endif +static const bool o = WGSL_SPEC_CONSTANT_0; + +[numthreads(1, 1, 1)] +void main() { + return; +} +/tmp/tint_N41FiO:2:2: error: spec constant required for constant id 0 +#error spec constant required for constant id 0 + ^ + + diff --git a/test/var/override/named/no_init/bool.wgsl.expected.msl b/test/var/override/named/no_init/bool.wgsl.expected.msl new file mode 100644 index 0000000000..95ce385236 --- /dev/null +++ b/test/var/override/named/no_init/bool.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +using namespace metal; +constant bool o [[function_constant(0)]]; +kernel void tint_symbol() { + return; +} + diff --git a/test/var/override/named/no_init/bool.wgsl.expected.spvasm b/test/var/override/named/no_init/bool.wgsl.expected.spvasm new file mode 100644 index 0000000000..ff12e004f7 --- /dev/null +++ b/test/var/override/named/no_init/bool.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 7 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %o "o" + OpName %main "main" + OpDecorate %o SpecId 0 + %bool = OpTypeBool + %o = OpSpecConstantFalse %bool + %void = OpTypeVoid + %3 = OpTypeFunction %void + %main = OpFunction %void None %3 + %6 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/var/override/named/no_init/bool.wgsl.expected.wgsl b/test/var/override/named/no_init/bool.wgsl.expected.wgsl new file mode 100644 index 0000000000..8fcab9b24e --- /dev/null +++ b/test/var/override/named/no_init/bool.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[override]] let o : bool; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/named/no_init/f32.wgsl b/test/var/override/named/no_init/f32.wgsl new file mode 100644 index 0000000000..c09196d385 --- /dev/null +++ b/test/var/override/named/no_init/f32.wgsl @@ -0,0 +1,6 @@ +[[override]] let o : f32; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/named/no_init/f32.wgsl.expected.hlsl b/test/var/override/named/no_init/f32.wgsl.expected.hlsl new file mode 100644 index 0000000000..2ae175b34e --- /dev/null +++ b/test/var/override/named/no_init/f32.wgsl.expected.hlsl @@ -0,0 +1,16 @@ +SKIP: FAILED + +#ifndef WGSL_SPEC_CONSTANT_0 +#error spec constant required for constant id 0 +#endif +static const float o = WGSL_SPEC_CONSTANT_0; + +[numthreads(1, 1, 1)] +void main() { + return; +} +/tmp/tint_gYmCvK:2:2: error: spec constant required for constant id 0 +#error spec constant required for constant id 0 + ^ + + diff --git a/test/var/override/named/no_init/f32.wgsl.expected.msl b/test/var/override/named/no_init/f32.wgsl.expected.msl new file mode 100644 index 0000000000..2ed922c18f --- /dev/null +++ b/test/var/override/named/no_init/f32.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +using namespace metal; +constant float o [[function_constant(0)]]; +kernel void tint_symbol() { + return; +} + diff --git a/test/var/override/named/no_init/f32.wgsl.expected.spvasm b/test/var/override/named/no_init/f32.wgsl.expected.spvasm new file mode 100644 index 0000000000..7ae9060492 --- /dev/null +++ b/test/var/override/named/no_init/f32.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 7 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %o "o" + OpName %main "main" + OpDecorate %o SpecId 0 + %float = OpTypeFloat 32 + %o = OpSpecConstant %float 0 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %main = OpFunction %void None %3 + %6 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/var/override/named/no_init/f32.wgsl.expected.wgsl b/test/var/override/named/no_init/f32.wgsl.expected.wgsl new file mode 100644 index 0000000000..5e60be6a91 --- /dev/null +++ b/test/var/override/named/no_init/f32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[override]] let o : f32; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/named/no_init/i32.wgsl b/test/var/override/named/no_init/i32.wgsl new file mode 100644 index 0000000000..790c56de03 --- /dev/null +++ b/test/var/override/named/no_init/i32.wgsl @@ -0,0 +1,6 @@ +[[override]] let o : i32; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/named/no_init/i32.wgsl.expected.hlsl b/test/var/override/named/no_init/i32.wgsl.expected.hlsl new file mode 100644 index 0000000000..7685972332 --- /dev/null +++ b/test/var/override/named/no_init/i32.wgsl.expected.hlsl @@ -0,0 +1,16 @@ +SKIP: FAILED + +#ifndef WGSL_SPEC_CONSTANT_0 +#error spec constant required for constant id 0 +#endif +static const int o = WGSL_SPEC_CONSTANT_0; + +[numthreads(1, 1, 1)] +void main() { + return; +} +/tmp/tint_kYjxDo:2:2: error: spec constant required for constant id 0 +#error spec constant required for constant id 0 + ^ + + diff --git a/test/var/override/named/no_init/i32.wgsl.expected.msl b/test/var/override/named/no_init/i32.wgsl.expected.msl new file mode 100644 index 0000000000..b5b1fae715 --- /dev/null +++ b/test/var/override/named/no_init/i32.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +using namespace metal; +constant int o [[function_constant(0)]]; +kernel void tint_symbol() { + return; +} + diff --git a/test/var/override/named/no_init/i32.wgsl.expected.spvasm b/test/var/override/named/no_init/i32.wgsl.expected.spvasm new file mode 100644 index 0000000000..18c8840440 --- /dev/null +++ b/test/var/override/named/no_init/i32.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 7 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %o "o" + OpName %main "main" + OpDecorate %o SpecId 0 + %int = OpTypeInt 32 1 + %o = OpSpecConstant %int 0 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %main = OpFunction %void None %3 + %6 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/var/override/named/no_init/i32.wgsl.expected.wgsl b/test/var/override/named/no_init/i32.wgsl.expected.wgsl new file mode 100644 index 0000000000..56a81dd6d3 --- /dev/null +++ b/test/var/override/named/no_init/i32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[override]] let o : i32; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/named/no_init/u32.wgsl b/test/var/override/named/no_init/u32.wgsl new file mode 100644 index 0000000000..b4826d8c4c --- /dev/null +++ b/test/var/override/named/no_init/u32.wgsl @@ -0,0 +1,6 @@ +[[override]] let o : u32; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/named/no_init/u32.wgsl.expected.hlsl b/test/var/override/named/no_init/u32.wgsl.expected.hlsl new file mode 100644 index 0000000000..4d5ed7a33f --- /dev/null +++ b/test/var/override/named/no_init/u32.wgsl.expected.hlsl @@ -0,0 +1,16 @@ +SKIP: FAILED + +#ifndef WGSL_SPEC_CONSTANT_0 +#error spec constant required for constant id 0 +#endif +static const uint o = WGSL_SPEC_CONSTANT_0; + +[numthreads(1, 1, 1)] +void main() { + return; +} +/tmp/tint_7FSY6m:2:2: error: spec constant required for constant id 0 +#error spec constant required for constant id 0 + ^ + + diff --git a/test/var/override/named/no_init/u32.wgsl.expected.msl b/test/var/override/named/no_init/u32.wgsl.expected.msl new file mode 100644 index 0000000000..918348cdd0 --- /dev/null +++ b/test/var/override/named/no_init/u32.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +using namespace metal; +constant uint o [[function_constant(0)]]; +kernel void tint_symbol() { + return; +} + diff --git a/test/var/override/named/no_init/u32.wgsl.expected.spvasm b/test/var/override/named/no_init/u32.wgsl.expected.spvasm new file mode 100644 index 0000000000..e5486b0dfb --- /dev/null +++ b/test/var/override/named/no_init/u32.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 7 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %o "o" + OpName %main "main" + OpDecorate %o SpecId 0 + %uint = OpTypeInt 32 0 + %o = OpSpecConstant %uint 0 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %main = OpFunction %void None %3 + %6 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/var/override/named/no_init/u32.wgsl.expected.wgsl b/test/var/override/named/no_init/u32.wgsl.expected.wgsl new file mode 100644 index 0000000000..ad44144470 --- /dev/null +++ b/test/var/override/named/no_init/u32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[override]] let o : u32; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/named/val_init/bool.wgsl b/test/var/override/named/val_init/bool.wgsl new file mode 100644 index 0000000000..0cb9566c06 --- /dev/null +++ b/test/var/override/named/val_init/bool.wgsl @@ -0,0 +1,6 @@ +[[override]] let o : bool = true; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/named/val_init/bool.wgsl.expected.hlsl b/test/var/override/named/val_init/bool.wgsl.expected.hlsl new file mode 100644 index 0000000000..e4ba79005c --- /dev/null +++ b/test/var/override/named/val_init/bool.wgsl.expected.hlsl @@ -0,0 +1,9 @@ +#ifndef WGSL_SPEC_CONSTANT_0 +#define WGSL_SPEC_CONSTANT_0 true +#endif +static const bool o = WGSL_SPEC_CONSTANT_0; + +[numthreads(1, 1, 1)] +void main() { + return; +} diff --git a/test/var/override/named/val_init/bool.wgsl.expected.msl b/test/var/override/named/val_init/bool.wgsl.expected.msl new file mode 100644 index 0000000000..95ce385236 --- /dev/null +++ b/test/var/override/named/val_init/bool.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +using namespace metal; +constant bool o [[function_constant(0)]]; +kernel void tint_symbol() { + return; +} + diff --git a/test/var/override/named/val_init/bool.wgsl.expected.spvasm b/test/var/override/named/val_init/bool.wgsl.expected.spvasm new file mode 100644 index 0000000000..277782e3ef --- /dev/null +++ b/test/var/override/named/val_init/bool.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 7 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %o "o" + OpName %main "main" + OpDecorate %o SpecId 0 + %bool = OpTypeBool + %o = OpSpecConstantTrue %bool + %void = OpTypeVoid + %3 = OpTypeFunction %void + %main = OpFunction %void None %3 + %6 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/var/override/named/val_init/bool.wgsl.expected.wgsl b/test/var/override/named/val_init/bool.wgsl.expected.wgsl new file mode 100644 index 0000000000..8e9edf50fe --- /dev/null +++ b/test/var/override/named/val_init/bool.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[override]] let o : bool = true; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/named/val_init/f32.wgsl b/test/var/override/named/val_init/f32.wgsl new file mode 100644 index 0000000000..904a8ad219 --- /dev/null +++ b/test/var/override/named/val_init/f32.wgsl @@ -0,0 +1,6 @@ +[[override]] let o : f32 = 1.0; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/named/val_init/f32.wgsl.expected.hlsl b/test/var/override/named/val_init/f32.wgsl.expected.hlsl new file mode 100644 index 0000000000..9d0794c1af --- /dev/null +++ b/test/var/override/named/val_init/f32.wgsl.expected.hlsl @@ -0,0 +1,9 @@ +#ifndef WGSL_SPEC_CONSTANT_0 +#define WGSL_SPEC_CONSTANT_0 1.0f +#endif +static const float o = WGSL_SPEC_CONSTANT_0; + +[numthreads(1, 1, 1)] +void main() { + return; +} diff --git a/test/var/override/named/val_init/f32.wgsl.expected.msl b/test/var/override/named/val_init/f32.wgsl.expected.msl new file mode 100644 index 0000000000..2ed922c18f --- /dev/null +++ b/test/var/override/named/val_init/f32.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +using namespace metal; +constant float o [[function_constant(0)]]; +kernel void tint_symbol() { + return; +} + diff --git a/test/var/override/named/val_init/f32.wgsl.expected.spvasm b/test/var/override/named/val_init/f32.wgsl.expected.spvasm new file mode 100644 index 0000000000..d5ce59ae6e --- /dev/null +++ b/test/var/override/named/val_init/f32.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 7 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %o "o" + OpName %main "main" + OpDecorate %o SpecId 0 + %float = OpTypeFloat 32 + %o = OpSpecConstant %float 1 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %main = OpFunction %void None %3 + %6 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/var/override/named/val_init/f32.wgsl.expected.wgsl b/test/var/override/named/val_init/f32.wgsl.expected.wgsl new file mode 100644 index 0000000000..32660b8598 --- /dev/null +++ b/test/var/override/named/val_init/f32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[override]] let o : f32 = 1.0; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/named/val_init/i32.wgsl b/test/var/override/named/val_init/i32.wgsl new file mode 100644 index 0000000000..c94ee67c20 --- /dev/null +++ b/test/var/override/named/val_init/i32.wgsl @@ -0,0 +1,6 @@ +[[override]] let o : i32 = 1; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/named/val_init/i32.wgsl.expected.hlsl b/test/var/override/named/val_init/i32.wgsl.expected.hlsl new file mode 100644 index 0000000000..fa2861c472 --- /dev/null +++ b/test/var/override/named/val_init/i32.wgsl.expected.hlsl @@ -0,0 +1,9 @@ +#ifndef WGSL_SPEC_CONSTANT_0 +#define WGSL_SPEC_CONSTANT_0 1 +#endif +static const int o = WGSL_SPEC_CONSTANT_0; + +[numthreads(1, 1, 1)] +void main() { + return; +} diff --git a/test/var/override/named/val_init/i32.wgsl.expected.msl b/test/var/override/named/val_init/i32.wgsl.expected.msl new file mode 100644 index 0000000000..b5b1fae715 --- /dev/null +++ b/test/var/override/named/val_init/i32.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +using namespace metal; +constant int o [[function_constant(0)]]; +kernel void tint_symbol() { + return; +} + diff --git a/test/var/override/named/val_init/i32.wgsl.expected.spvasm b/test/var/override/named/val_init/i32.wgsl.expected.spvasm new file mode 100644 index 0000000000..7d30d25b6e --- /dev/null +++ b/test/var/override/named/val_init/i32.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 7 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %o "o" + OpName %main "main" + OpDecorate %o SpecId 0 + %int = OpTypeInt 32 1 + %o = OpSpecConstant %int 1 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %main = OpFunction %void None %3 + %6 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/var/override/named/val_init/i32.wgsl.expected.wgsl b/test/var/override/named/val_init/i32.wgsl.expected.wgsl new file mode 100644 index 0000000000..b897fd53cf --- /dev/null +++ b/test/var/override/named/val_init/i32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[override]] let o : i32 = 1; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/named/val_init/u32.wgsl b/test/var/override/named/val_init/u32.wgsl new file mode 100644 index 0000000000..8c8df63c51 --- /dev/null +++ b/test/var/override/named/val_init/u32.wgsl @@ -0,0 +1,6 @@ +[[override]] let o : u32 = 1u; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/named/val_init/u32.wgsl.expected.hlsl b/test/var/override/named/val_init/u32.wgsl.expected.hlsl new file mode 100644 index 0000000000..9529784468 --- /dev/null +++ b/test/var/override/named/val_init/u32.wgsl.expected.hlsl @@ -0,0 +1,9 @@ +#ifndef WGSL_SPEC_CONSTANT_0 +#define WGSL_SPEC_CONSTANT_0 1u +#endif +static const uint o = WGSL_SPEC_CONSTANT_0; + +[numthreads(1, 1, 1)] +void main() { + return; +} diff --git a/test/var/override/named/val_init/u32.wgsl.expected.msl b/test/var/override/named/val_init/u32.wgsl.expected.msl new file mode 100644 index 0000000000..918348cdd0 --- /dev/null +++ b/test/var/override/named/val_init/u32.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +using namespace metal; +constant uint o [[function_constant(0)]]; +kernel void tint_symbol() { + return; +} + diff --git a/test/var/override/named/val_init/u32.wgsl.expected.spvasm b/test/var/override/named/val_init/u32.wgsl.expected.spvasm new file mode 100644 index 0000000000..93718f11d9 --- /dev/null +++ b/test/var/override/named/val_init/u32.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 7 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %o "o" + OpName %main "main" + OpDecorate %o SpecId 0 + %uint = OpTypeInt 32 0 + %o = OpSpecConstant %uint 1 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %main = OpFunction %void None %3 + %6 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/var/override/named/val_init/u32.wgsl.expected.wgsl b/test/var/override/named/val_init/u32.wgsl.expected.wgsl new file mode 100644 index 0000000000..3a62e22352 --- /dev/null +++ b/test/var/override/named/val_init/u32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[override]] let o : u32 = 1u; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/named/zero_init/bool.wgsl b/test/var/override/named/zero_init/bool.wgsl new file mode 100644 index 0000000000..1e03a45923 --- /dev/null +++ b/test/var/override/named/zero_init/bool.wgsl @@ -0,0 +1,6 @@ +[[override]] let o : bool = bool(); + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/named/zero_init/bool.wgsl.expected.hlsl b/test/var/override/named/zero_init/bool.wgsl.expected.hlsl new file mode 100644 index 0000000000..d14a339663 --- /dev/null +++ b/test/var/override/named/zero_init/bool.wgsl.expected.hlsl @@ -0,0 +1,9 @@ +#ifndef WGSL_SPEC_CONSTANT_0 +#define WGSL_SPEC_CONSTANT_0 false +#endif +static const bool o = WGSL_SPEC_CONSTANT_0; + +[numthreads(1, 1, 1)] +void main() { + return; +} diff --git a/test/var/override/named/zero_init/bool.wgsl.expected.msl b/test/var/override/named/zero_init/bool.wgsl.expected.msl new file mode 100644 index 0000000000..95ce385236 --- /dev/null +++ b/test/var/override/named/zero_init/bool.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +using namespace metal; +constant bool o [[function_constant(0)]]; +kernel void tint_symbol() { + return; +} + diff --git a/test/var/override/named/zero_init/bool.wgsl.expected.spvasm b/test/var/override/named/zero_init/bool.wgsl.expected.spvasm new file mode 100644 index 0000000000..ff12e004f7 --- /dev/null +++ b/test/var/override/named/zero_init/bool.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 7 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %o "o" + OpName %main "main" + OpDecorate %o SpecId 0 + %bool = OpTypeBool + %o = OpSpecConstantFalse %bool + %void = OpTypeVoid + %3 = OpTypeFunction %void + %main = OpFunction %void None %3 + %6 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/var/override/named/zero_init/bool.wgsl.expected.wgsl b/test/var/override/named/zero_init/bool.wgsl.expected.wgsl new file mode 100644 index 0000000000..511bd67843 --- /dev/null +++ b/test/var/override/named/zero_init/bool.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[override]] let o : bool = bool(); + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/named/zero_init/f32.wgsl b/test/var/override/named/zero_init/f32.wgsl new file mode 100644 index 0000000000..a47d934701 --- /dev/null +++ b/test/var/override/named/zero_init/f32.wgsl @@ -0,0 +1,6 @@ +[[override]] let o : f32 = f32(); + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/named/zero_init/f32.wgsl.expected.hlsl b/test/var/override/named/zero_init/f32.wgsl.expected.hlsl new file mode 100644 index 0000000000..5ae59a4ccf --- /dev/null +++ b/test/var/override/named/zero_init/f32.wgsl.expected.hlsl @@ -0,0 +1,9 @@ +#ifndef WGSL_SPEC_CONSTANT_0 +#define WGSL_SPEC_CONSTANT_0 0.0f +#endif +static const float o = WGSL_SPEC_CONSTANT_0; + +[numthreads(1, 1, 1)] +void main() { + return; +} diff --git a/test/var/override/named/zero_init/f32.wgsl.expected.msl b/test/var/override/named/zero_init/f32.wgsl.expected.msl new file mode 100644 index 0000000000..2ed922c18f --- /dev/null +++ b/test/var/override/named/zero_init/f32.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +using namespace metal; +constant float o [[function_constant(0)]]; +kernel void tint_symbol() { + return; +} + diff --git a/test/var/override/named/zero_init/f32.wgsl.expected.spvasm b/test/var/override/named/zero_init/f32.wgsl.expected.spvasm new file mode 100644 index 0000000000..7ae9060492 --- /dev/null +++ b/test/var/override/named/zero_init/f32.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 7 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %o "o" + OpName %main "main" + OpDecorate %o SpecId 0 + %float = OpTypeFloat 32 + %o = OpSpecConstant %float 0 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %main = OpFunction %void None %3 + %6 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/var/override/named/zero_init/f32.wgsl.expected.wgsl b/test/var/override/named/zero_init/f32.wgsl.expected.wgsl new file mode 100644 index 0000000000..fe1468f1b5 --- /dev/null +++ b/test/var/override/named/zero_init/f32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[override]] let o : f32 = f32(); + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/named/zero_init/i32.wgsl b/test/var/override/named/zero_init/i32.wgsl new file mode 100644 index 0000000000..64552272a3 --- /dev/null +++ b/test/var/override/named/zero_init/i32.wgsl @@ -0,0 +1,6 @@ +[[override]] let o : i32 = i32(); + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/named/zero_init/i32.wgsl.expected.hlsl b/test/var/override/named/zero_init/i32.wgsl.expected.hlsl new file mode 100644 index 0000000000..6af794b4cb --- /dev/null +++ b/test/var/override/named/zero_init/i32.wgsl.expected.hlsl @@ -0,0 +1,9 @@ +#ifndef WGSL_SPEC_CONSTANT_0 +#define WGSL_SPEC_CONSTANT_0 0 +#endif +static const int o = WGSL_SPEC_CONSTANT_0; + +[numthreads(1, 1, 1)] +void main() { + return; +} diff --git a/test/var/override/named/zero_init/i32.wgsl.expected.msl b/test/var/override/named/zero_init/i32.wgsl.expected.msl new file mode 100644 index 0000000000..b5b1fae715 --- /dev/null +++ b/test/var/override/named/zero_init/i32.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +using namespace metal; +constant int o [[function_constant(0)]]; +kernel void tint_symbol() { + return; +} + diff --git a/test/var/override/named/zero_init/i32.wgsl.expected.spvasm b/test/var/override/named/zero_init/i32.wgsl.expected.spvasm new file mode 100644 index 0000000000..18c8840440 --- /dev/null +++ b/test/var/override/named/zero_init/i32.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 7 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %o "o" + OpName %main "main" + OpDecorate %o SpecId 0 + %int = OpTypeInt 32 1 + %o = OpSpecConstant %int 0 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %main = OpFunction %void None %3 + %6 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/var/override/named/zero_init/i32.wgsl.expected.wgsl b/test/var/override/named/zero_init/i32.wgsl.expected.wgsl new file mode 100644 index 0000000000..5cde67078e --- /dev/null +++ b/test/var/override/named/zero_init/i32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[override]] let o : i32 = i32(); + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/named/zero_init/u32.wgsl b/test/var/override/named/zero_init/u32.wgsl new file mode 100644 index 0000000000..9d1205e72b --- /dev/null +++ b/test/var/override/named/zero_init/u32.wgsl @@ -0,0 +1,6 @@ +[[override]] let o : u32 = u32(); + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/named/zero_init/u32.wgsl.expected.hlsl b/test/var/override/named/zero_init/u32.wgsl.expected.hlsl new file mode 100644 index 0000000000..17f5ddf612 --- /dev/null +++ b/test/var/override/named/zero_init/u32.wgsl.expected.hlsl @@ -0,0 +1,9 @@ +#ifndef WGSL_SPEC_CONSTANT_0 +#define WGSL_SPEC_CONSTANT_0 0u +#endif +static const uint o = WGSL_SPEC_CONSTANT_0; + +[numthreads(1, 1, 1)] +void main() { + return; +} diff --git a/test/var/override/named/zero_init/u32.wgsl.expected.msl b/test/var/override/named/zero_init/u32.wgsl.expected.msl new file mode 100644 index 0000000000..918348cdd0 --- /dev/null +++ b/test/var/override/named/zero_init/u32.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +using namespace metal; +constant uint o [[function_constant(0)]]; +kernel void tint_symbol() { + return; +} + diff --git a/test/var/override/named/zero_init/u32.wgsl.expected.spvasm b/test/var/override/named/zero_init/u32.wgsl.expected.spvasm new file mode 100644 index 0000000000..e5486b0dfb --- /dev/null +++ b/test/var/override/named/zero_init/u32.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 7 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %o "o" + OpName %main "main" + OpDecorate %o SpecId 0 + %uint = OpTypeInt 32 0 + %o = OpSpecConstant %uint 0 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %main = OpFunction %void None %3 + %6 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/var/override/named/zero_init/u32.wgsl.expected.wgsl b/test/var/override/named/zero_init/u32.wgsl.expected.wgsl new file mode 100644 index 0000000000..f9d5644b9d --- /dev/null +++ b/test/var/override/named/zero_init/u32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[override]] let o : u32 = u32(); + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/numbered/no_init/bool.wgsl b/test/var/override/numbered/no_init/bool.wgsl new file mode 100644 index 0000000000..a8b81c7090 --- /dev/null +++ b/test/var/override/numbered/no_init/bool.wgsl @@ -0,0 +1,6 @@ +[[override(1234)]] let o : bool; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/numbered/no_init/bool.wgsl.expected.hlsl b/test/var/override/numbered/no_init/bool.wgsl.expected.hlsl new file mode 100644 index 0000000000..eca597705d --- /dev/null +++ b/test/var/override/numbered/no_init/bool.wgsl.expected.hlsl @@ -0,0 +1,16 @@ +SKIP: FAILED + +#ifndef WGSL_SPEC_CONSTANT_1234 +#error spec constant required for constant id 1234 +#endif +static const bool o = WGSL_SPEC_CONSTANT_1234; + +[numthreads(1, 1, 1)] +void main() { + return; +} +/tmp/tint_4uqZfq:2:2: error: spec constant required for constant id 1234 +#error spec constant required for constant id 1234 + ^ + + diff --git a/test/var/override/numbered/no_init/bool.wgsl.expected.msl b/test/var/override/numbered/no_init/bool.wgsl.expected.msl new file mode 100644 index 0000000000..d6cbfe9df6 --- /dev/null +++ b/test/var/override/numbered/no_init/bool.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +using namespace metal; +constant bool o [[function_constant(1234)]]; +kernel void tint_symbol() { + return; +} + diff --git a/test/var/override/numbered/no_init/bool.wgsl.expected.spvasm b/test/var/override/numbered/no_init/bool.wgsl.expected.spvasm new file mode 100644 index 0000000000..b1169e6faa --- /dev/null +++ b/test/var/override/numbered/no_init/bool.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 7 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %o "o" + OpName %main "main" + OpDecorate %o SpecId 1234 + %bool = OpTypeBool + %o = OpSpecConstantFalse %bool + %void = OpTypeVoid + %3 = OpTypeFunction %void + %main = OpFunction %void None %3 + %6 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/var/override/numbered/no_init/bool.wgsl.expected.wgsl b/test/var/override/numbered/no_init/bool.wgsl.expected.wgsl new file mode 100644 index 0000000000..e8d6dafc2d --- /dev/null +++ b/test/var/override/numbered/no_init/bool.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[override(1234)]] let o : bool; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/numbered/no_init/f32.wgsl b/test/var/override/numbered/no_init/f32.wgsl new file mode 100644 index 0000000000..95e9dfda24 --- /dev/null +++ b/test/var/override/numbered/no_init/f32.wgsl @@ -0,0 +1,6 @@ +[[override(1234)]] let o : f32; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/numbered/no_init/f32.wgsl.expected.hlsl b/test/var/override/numbered/no_init/f32.wgsl.expected.hlsl new file mode 100644 index 0000000000..f5c5a75fd6 --- /dev/null +++ b/test/var/override/numbered/no_init/f32.wgsl.expected.hlsl @@ -0,0 +1,16 @@ +SKIP: FAILED + +#ifndef WGSL_SPEC_CONSTANT_1234 +#error spec constant required for constant id 1234 +#endif +static const float o = WGSL_SPEC_CONSTANT_1234; + +[numthreads(1, 1, 1)] +void main() { + return; +} +/tmp/tint_UepNPT:2:2: error: spec constant required for constant id 1234 +#error spec constant required for constant id 1234 + ^ + + diff --git a/test/var/override/numbered/no_init/f32.wgsl.expected.msl b/test/var/override/numbered/no_init/f32.wgsl.expected.msl new file mode 100644 index 0000000000..41134e2213 --- /dev/null +++ b/test/var/override/numbered/no_init/f32.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +using namespace metal; +constant float o [[function_constant(1234)]]; +kernel void tint_symbol() { + return; +} + diff --git a/test/var/override/numbered/no_init/f32.wgsl.expected.spvasm b/test/var/override/numbered/no_init/f32.wgsl.expected.spvasm new file mode 100644 index 0000000000..a253f64b1c --- /dev/null +++ b/test/var/override/numbered/no_init/f32.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 7 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %o "o" + OpName %main "main" + OpDecorate %o SpecId 1234 + %float = OpTypeFloat 32 + %o = OpSpecConstant %float 0 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %main = OpFunction %void None %3 + %6 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/var/override/numbered/no_init/f32.wgsl.expected.wgsl b/test/var/override/numbered/no_init/f32.wgsl.expected.wgsl new file mode 100644 index 0000000000..cabed6a160 --- /dev/null +++ b/test/var/override/numbered/no_init/f32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[override(1234)]] let o : f32; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/numbered/no_init/i32.wgsl b/test/var/override/numbered/no_init/i32.wgsl new file mode 100644 index 0000000000..74846ef2f9 --- /dev/null +++ b/test/var/override/numbered/no_init/i32.wgsl @@ -0,0 +1,6 @@ +[[override(1234)]] let o : i32; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/numbered/no_init/i32.wgsl.expected.hlsl b/test/var/override/numbered/no_init/i32.wgsl.expected.hlsl new file mode 100644 index 0000000000..0b62c77b78 --- /dev/null +++ b/test/var/override/numbered/no_init/i32.wgsl.expected.hlsl @@ -0,0 +1,16 @@ +SKIP: FAILED + +#ifndef WGSL_SPEC_CONSTANT_1234 +#error spec constant required for constant id 1234 +#endif +static const int o = WGSL_SPEC_CONSTANT_1234; + +[numthreads(1, 1, 1)] +void main() { + return; +} +/tmp/tint_mwzbqN:2:2: error: spec constant required for constant id 1234 +#error spec constant required for constant id 1234 + ^ + + diff --git a/test/var/override/numbered/no_init/i32.wgsl.expected.msl b/test/var/override/numbered/no_init/i32.wgsl.expected.msl new file mode 100644 index 0000000000..812c466766 --- /dev/null +++ b/test/var/override/numbered/no_init/i32.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +using namespace metal; +constant int o [[function_constant(1234)]]; +kernel void tint_symbol() { + return; +} + diff --git a/test/var/override/numbered/no_init/i32.wgsl.expected.spvasm b/test/var/override/numbered/no_init/i32.wgsl.expected.spvasm new file mode 100644 index 0000000000..aca0b07388 --- /dev/null +++ b/test/var/override/numbered/no_init/i32.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 7 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %o "o" + OpName %main "main" + OpDecorate %o SpecId 1234 + %int = OpTypeInt 32 1 + %o = OpSpecConstant %int 0 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %main = OpFunction %void None %3 + %6 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/var/override/numbered/no_init/i32.wgsl.expected.wgsl b/test/var/override/numbered/no_init/i32.wgsl.expected.wgsl new file mode 100644 index 0000000000..68fb678baf --- /dev/null +++ b/test/var/override/numbered/no_init/i32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[override(1234)]] let o : i32; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/numbered/no_init/u32.wgsl b/test/var/override/numbered/no_init/u32.wgsl new file mode 100644 index 0000000000..152c709d42 --- /dev/null +++ b/test/var/override/numbered/no_init/u32.wgsl @@ -0,0 +1,6 @@ +[[override(1234)]] let o : u32; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/numbered/no_init/u32.wgsl.expected.hlsl b/test/var/override/numbered/no_init/u32.wgsl.expected.hlsl new file mode 100644 index 0000000000..443a2b155e --- /dev/null +++ b/test/var/override/numbered/no_init/u32.wgsl.expected.hlsl @@ -0,0 +1,16 @@ +SKIP: FAILED + +#ifndef WGSL_SPEC_CONSTANT_1234 +#error spec constant required for constant id 1234 +#endif +static const uint o = WGSL_SPEC_CONSTANT_1234; + +[numthreads(1, 1, 1)] +void main() { + return; +} +/tmp/tint_U0EZFZ:2:2: error: spec constant required for constant id 1234 +#error spec constant required for constant id 1234 + ^ + + diff --git a/test/var/override/numbered/no_init/u32.wgsl.expected.msl b/test/var/override/numbered/no_init/u32.wgsl.expected.msl new file mode 100644 index 0000000000..799a682390 --- /dev/null +++ b/test/var/override/numbered/no_init/u32.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +using namespace metal; +constant uint o [[function_constant(1234)]]; +kernel void tint_symbol() { + return; +} + diff --git a/test/var/override/numbered/no_init/u32.wgsl.expected.spvasm b/test/var/override/numbered/no_init/u32.wgsl.expected.spvasm new file mode 100644 index 0000000000..ca38a9a94c --- /dev/null +++ b/test/var/override/numbered/no_init/u32.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 7 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %o "o" + OpName %main "main" + OpDecorate %o SpecId 1234 + %uint = OpTypeInt 32 0 + %o = OpSpecConstant %uint 0 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %main = OpFunction %void None %3 + %6 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/var/override/numbered/no_init/u32.wgsl.expected.wgsl b/test/var/override/numbered/no_init/u32.wgsl.expected.wgsl new file mode 100644 index 0000000000..40cb6c0d4a --- /dev/null +++ b/test/var/override/numbered/no_init/u32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[override(1234)]] let o : u32; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/numbered/val_init/bool.wgsl b/test/var/override/numbered/val_init/bool.wgsl new file mode 100644 index 0000000000..c29ec30a0d --- /dev/null +++ b/test/var/override/numbered/val_init/bool.wgsl @@ -0,0 +1,6 @@ +[[override(1234)]] let o : bool = true; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/numbered/val_init/bool.wgsl.expected.hlsl b/test/var/override/numbered/val_init/bool.wgsl.expected.hlsl new file mode 100644 index 0000000000..6e189ede17 --- /dev/null +++ b/test/var/override/numbered/val_init/bool.wgsl.expected.hlsl @@ -0,0 +1,9 @@ +#ifndef WGSL_SPEC_CONSTANT_1234 +#define WGSL_SPEC_CONSTANT_1234 true +#endif +static const bool o = WGSL_SPEC_CONSTANT_1234; + +[numthreads(1, 1, 1)] +void main() { + return; +} diff --git a/test/var/override/numbered/val_init/bool.wgsl.expected.msl b/test/var/override/numbered/val_init/bool.wgsl.expected.msl new file mode 100644 index 0000000000..d6cbfe9df6 --- /dev/null +++ b/test/var/override/numbered/val_init/bool.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +using namespace metal; +constant bool o [[function_constant(1234)]]; +kernel void tint_symbol() { + return; +} + diff --git a/test/var/override/numbered/val_init/bool.wgsl.expected.spvasm b/test/var/override/numbered/val_init/bool.wgsl.expected.spvasm new file mode 100644 index 0000000000..060f8603f6 --- /dev/null +++ b/test/var/override/numbered/val_init/bool.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 7 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %o "o" + OpName %main "main" + OpDecorate %o SpecId 1234 + %bool = OpTypeBool + %o = OpSpecConstantTrue %bool + %void = OpTypeVoid + %3 = OpTypeFunction %void + %main = OpFunction %void None %3 + %6 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/var/override/numbered/val_init/bool.wgsl.expected.wgsl b/test/var/override/numbered/val_init/bool.wgsl.expected.wgsl new file mode 100644 index 0000000000..2ab9051487 --- /dev/null +++ b/test/var/override/numbered/val_init/bool.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[override(1234)]] let o : bool = true; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/numbered/val_init/f32.wgsl b/test/var/override/numbered/val_init/f32.wgsl new file mode 100644 index 0000000000..649d4e8658 --- /dev/null +++ b/test/var/override/numbered/val_init/f32.wgsl @@ -0,0 +1,6 @@ +[[override(1234)]] let o : f32 = 1.0; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/numbered/val_init/f32.wgsl.expected.hlsl b/test/var/override/numbered/val_init/f32.wgsl.expected.hlsl new file mode 100644 index 0000000000..4671b12b3b --- /dev/null +++ b/test/var/override/numbered/val_init/f32.wgsl.expected.hlsl @@ -0,0 +1,9 @@ +#ifndef WGSL_SPEC_CONSTANT_1234 +#define WGSL_SPEC_CONSTANT_1234 1.0f +#endif +static const float o = WGSL_SPEC_CONSTANT_1234; + +[numthreads(1, 1, 1)] +void main() { + return; +} diff --git a/test/var/override/numbered/val_init/f32.wgsl.expected.msl b/test/var/override/numbered/val_init/f32.wgsl.expected.msl new file mode 100644 index 0000000000..41134e2213 --- /dev/null +++ b/test/var/override/numbered/val_init/f32.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +using namespace metal; +constant float o [[function_constant(1234)]]; +kernel void tint_symbol() { + return; +} + diff --git a/test/var/override/numbered/val_init/f32.wgsl.expected.spvasm b/test/var/override/numbered/val_init/f32.wgsl.expected.spvasm new file mode 100644 index 0000000000..c75755700c --- /dev/null +++ b/test/var/override/numbered/val_init/f32.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 7 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %o "o" + OpName %main "main" + OpDecorate %o SpecId 1234 + %float = OpTypeFloat 32 + %o = OpSpecConstant %float 1 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %main = OpFunction %void None %3 + %6 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/var/override/numbered/val_init/f32.wgsl.expected.wgsl b/test/var/override/numbered/val_init/f32.wgsl.expected.wgsl new file mode 100644 index 0000000000..bb12693e58 --- /dev/null +++ b/test/var/override/numbered/val_init/f32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[override(1234)]] let o : f32 = 1.0; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/numbered/val_init/i32.wgsl b/test/var/override/numbered/val_init/i32.wgsl new file mode 100644 index 0000000000..d867f6a8c8 --- /dev/null +++ b/test/var/override/numbered/val_init/i32.wgsl @@ -0,0 +1,6 @@ +[[override(1234)]] let o : i32 = 1; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/numbered/val_init/i32.wgsl.expected.hlsl b/test/var/override/numbered/val_init/i32.wgsl.expected.hlsl new file mode 100644 index 0000000000..cfb2fd9371 --- /dev/null +++ b/test/var/override/numbered/val_init/i32.wgsl.expected.hlsl @@ -0,0 +1,9 @@ +#ifndef WGSL_SPEC_CONSTANT_1234 +#define WGSL_SPEC_CONSTANT_1234 1 +#endif +static const int o = WGSL_SPEC_CONSTANT_1234; + +[numthreads(1, 1, 1)] +void main() { + return; +} diff --git a/test/var/override/numbered/val_init/i32.wgsl.expected.msl b/test/var/override/numbered/val_init/i32.wgsl.expected.msl new file mode 100644 index 0000000000..812c466766 --- /dev/null +++ b/test/var/override/numbered/val_init/i32.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +using namespace metal; +constant int o [[function_constant(1234)]]; +kernel void tint_symbol() { + return; +} + diff --git a/test/var/override/numbered/val_init/i32.wgsl.expected.spvasm b/test/var/override/numbered/val_init/i32.wgsl.expected.spvasm new file mode 100644 index 0000000000..f66889a42e --- /dev/null +++ b/test/var/override/numbered/val_init/i32.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 7 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %o "o" + OpName %main "main" + OpDecorate %o SpecId 1234 + %int = OpTypeInt 32 1 + %o = OpSpecConstant %int 1 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %main = OpFunction %void None %3 + %6 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/var/override/numbered/val_init/i32.wgsl.expected.wgsl b/test/var/override/numbered/val_init/i32.wgsl.expected.wgsl new file mode 100644 index 0000000000..592dd92695 --- /dev/null +++ b/test/var/override/numbered/val_init/i32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[override(1234)]] let o : i32 = 1; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/numbered/val_init/u32.wgsl b/test/var/override/numbered/val_init/u32.wgsl new file mode 100644 index 0000000000..22c55f27d5 --- /dev/null +++ b/test/var/override/numbered/val_init/u32.wgsl @@ -0,0 +1,6 @@ +[[override(1234)]] let o : u32 = 1u; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/numbered/val_init/u32.wgsl.expected.hlsl b/test/var/override/numbered/val_init/u32.wgsl.expected.hlsl new file mode 100644 index 0000000000..36437dd9e2 --- /dev/null +++ b/test/var/override/numbered/val_init/u32.wgsl.expected.hlsl @@ -0,0 +1,9 @@ +#ifndef WGSL_SPEC_CONSTANT_1234 +#define WGSL_SPEC_CONSTANT_1234 1u +#endif +static const uint o = WGSL_SPEC_CONSTANT_1234; + +[numthreads(1, 1, 1)] +void main() { + return; +} diff --git a/test/var/override/numbered/val_init/u32.wgsl.expected.msl b/test/var/override/numbered/val_init/u32.wgsl.expected.msl new file mode 100644 index 0000000000..799a682390 --- /dev/null +++ b/test/var/override/numbered/val_init/u32.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +using namespace metal; +constant uint o [[function_constant(1234)]]; +kernel void tint_symbol() { + return; +} + diff --git a/test/var/override/numbered/val_init/u32.wgsl.expected.spvasm b/test/var/override/numbered/val_init/u32.wgsl.expected.spvasm new file mode 100644 index 0000000000..2b51d1ce62 --- /dev/null +++ b/test/var/override/numbered/val_init/u32.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 7 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %o "o" + OpName %main "main" + OpDecorate %o SpecId 1234 + %uint = OpTypeInt 32 0 + %o = OpSpecConstant %uint 1 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %main = OpFunction %void None %3 + %6 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/var/override/numbered/val_init/u32.wgsl.expected.wgsl b/test/var/override/numbered/val_init/u32.wgsl.expected.wgsl new file mode 100644 index 0000000000..18bb63b255 --- /dev/null +++ b/test/var/override/numbered/val_init/u32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[override(1234)]] let o : u32 = 1u; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/numbered/zero_init/bool.wgsl b/test/var/override/numbered/zero_init/bool.wgsl new file mode 100644 index 0000000000..894b5847fb --- /dev/null +++ b/test/var/override/numbered/zero_init/bool.wgsl @@ -0,0 +1,6 @@ +[[override(1234)]] let o : bool = bool(); + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/numbered/zero_init/bool.wgsl.expected.hlsl b/test/var/override/numbered/zero_init/bool.wgsl.expected.hlsl new file mode 100644 index 0000000000..86e9b98f87 --- /dev/null +++ b/test/var/override/numbered/zero_init/bool.wgsl.expected.hlsl @@ -0,0 +1,9 @@ +#ifndef WGSL_SPEC_CONSTANT_1234 +#define WGSL_SPEC_CONSTANT_1234 false +#endif +static const bool o = WGSL_SPEC_CONSTANT_1234; + +[numthreads(1, 1, 1)] +void main() { + return; +} diff --git a/test/var/override/numbered/zero_init/bool.wgsl.expected.msl b/test/var/override/numbered/zero_init/bool.wgsl.expected.msl new file mode 100644 index 0000000000..d6cbfe9df6 --- /dev/null +++ b/test/var/override/numbered/zero_init/bool.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +using namespace metal; +constant bool o [[function_constant(1234)]]; +kernel void tint_symbol() { + return; +} + diff --git a/test/var/override/numbered/zero_init/bool.wgsl.expected.spvasm b/test/var/override/numbered/zero_init/bool.wgsl.expected.spvasm new file mode 100644 index 0000000000..b1169e6faa --- /dev/null +++ b/test/var/override/numbered/zero_init/bool.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 7 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %o "o" + OpName %main "main" + OpDecorate %o SpecId 1234 + %bool = OpTypeBool + %o = OpSpecConstantFalse %bool + %void = OpTypeVoid + %3 = OpTypeFunction %void + %main = OpFunction %void None %3 + %6 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/var/override/numbered/zero_init/bool.wgsl.expected.wgsl b/test/var/override/numbered/zero_init/bool.wgsl.expected.wgsl new file mode 100644 index 0000000000..653b1da727 --- /dev/null +++ b/test/var/override/numbered/zero_init/bool.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[override(1234)]] let o : bool = bool(); + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/numbered/zero_init/f32.wgsl b/test/var/override/numbered/zero_init/f32.wgsl new file mode 100644 index 0000000000..0a3908cc67 --- /dev/null +++ b/test/var/override/numbered/zero_init/f32.wgsl @@ -0,0 +1,6 @@ +[[override(1234)]] let o : f32 = f32(); + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/numbered/zero_init/f32.wgsl.expected.hlsl b/test/var/override/numbered/zero_init/f32.wgsl.expected.hlsl new file mode 100644 index 0000000000..70f868fce9 --- /dev/null +++ b/test/var/override/numbered/zero_init/f32.wgsl.expected.hlsl @@ -0,0 +1,9 @@ +#ifndef WGSL_SPEC_CONSTANT_1234 +#define WGSL_SPEC_CONSTANT_1234 0.0f +#endif +static const float o = WGSL_SPEC_CONSTANT_1234; + +[numthreads(1, 1, 1)] +void main() { + return; +} diff --git a/test/var/override/numbered/zero_init/f32.wgsl.expected.msl b/test/var/override/numbered/zero_init/f32.wgsl.expected.msl new file mode 100644 index 0000000000..41134e2213 --- /dev/null +++ b/test/var/override/numbered/zero_init/f32.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +using namespace metal; +constant float o [[function_constant(1234)]]; +kernel void tint_symbol() { + return; +} + diff --git a/test/var/override/numbered/zero_init/f32.wgsl.expected.spvasm b/test/var/override/numbered/zero_init/f32.wgsl.expected.spvasm new file mode 100644 index 0000000000..a253f64b1c --- /dev/null +++ b/test/var/override/numbered/zero_init/f32.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 7 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %o "o" + OpName %main "main" + OpDecorate %o SpecId 1234 + %float = OpTypeFloat 32 + %o = OpSpecConstant %float 0 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %main = OpFunction %void None %3 + %6 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/var/override/numbered/zero_init/f32.wgsl.expected.wgsl b/test/var/override/numbered/zero_init/f32.wgsl.expected.wgsl new file mode 100644 index 0000000000..589751bd77 --- /dev/null +++ b/test/var/override/numbered/zero_init/f32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[override(1234)]] let o : f32 = f32(); + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/numbered/zero_init/i32.wgsl b/test/var/override/numbered/zero_init/i32.wgsl new file mode 100644 index 0000000000..4dd63fd44a --- /dev/null +++ b/test/var/override/numbered/zero_init/i32.wgsl @@ -0,0 +1,6 @@ +[[override(1234)]] let o : i32 = i32(); + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/numbered/zero_init/i32.wgsl.expected.hlsl b/test/var/override/numbered/zero_init/i32.wgsl.expected.hlsl new file mode 100644 index 0000000000..654ffd6295 --- /dev/null +++ b/test/var/override/numbered/zero_init/i32.wgsl.expected.hlsl @@ -0,0 +1,9 @@ +#ifndef WGSL_SPEC_CONSTANT_1234 +#define WGSL_SPEC_CONSTANT_1234 0 +#endif +static const int o = WGSL_SPEC_CONSTANT_1234; + +[numthreads(1, 1, 1)] +void main() { + return; +} diff --git a/test/var/override/numbered/zero_init/i32.wgsl.expected.msl b/test/var/override/numbered/zero_init/i32.wgsl.expected.msl new file mode 100644 index 0000000000..812c466766 --- /dev/null +++ b/test/var/override/numbered/zero_init/i32.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +using namespace metal; +constant int o [[function_constant(1234)]]; +kernel void tint_symbol() { + return; +} + diff --git a/test/var/override/numbered/zero_init/i32.wgsl.expected.spvasm b/test/var/override/numbered/zero_init/i32.wgsl.expected.spvasm new file mode 100644 index 0000000000..aca0b07388 --- /dev/null +++ b/test/var/override/numbered/zero_init/i32.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 7 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %o "o" + OpName %main "main" + OpDecorate %o SpecId 1234 + %int = OpTypeInt 32 1 + %o = OpSpecConstant %int 0 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %main = OpFunction %void None %3 + %6 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/var/override/numbered/zero_init/i32.wgsl.expected.wgsl b/test/var/override/numbered/zero_init/i32.wgsl.expected.wgsl new file mode 100644 index 0000000000..8c0adcaf38 --- /dev/null +++ b/test/var/override/numbered/zero_init/i32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[override(1234)]] let o : i32 = i32(); + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/numbered/zero_init/u32.wgsl b/test/var/override/numbered/zero_init/u32.wgsl new file mode 100644 index 0000000000..d525204378 --- /dev/null +++ b/test/var/override/numbered/zero_init/u32.wgsl @@ -0,0 +1,6 @@ +[[override(1234)]] let o : u32 = u32(); + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +} diff --git a/test/var/override/numbered/zero_init/u32.wgsl.expected.hlsl b/test/var/override/numbered/zero_init/u32.wgsl.expected.hlsl new file mode 100644 index 0000000000..2fe25d834a --- /dev/null +++ b/test/var/override/numbered/zero_init/u32.wgsl.expected.hlsl @@ -0,0 +1,9 @@ +#ifndef WGSL_SPEC_CONSTANT_1234 +#define WGSL_SPEC_CONSTANT_1234 0u +#endif +static const uint o = WGSL_SPEC_CONSTANT_1234; + +[numthreads(1, 1, 1)] +void main() { + return; +} diff --git a/test/var/override/numbered/zero_init/u32.wgsl.expected.msl b/test/var/override/numbered/zero_init/u32.wgsl.expected.msl new file mode 100644 index 0000000000..799a682390 --- /dev/null +++ b/test/var/override/numbered/zero_init/u32.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +using namespace metal; +constant uint o [[function_constant(1234)]]; +kernel void tint_symbol() { + return; +} + diff --git a/test/var/override/numbered/zero_init/u32.wgsl.expected.spvasm b/test/var/override/numbered/zero_init/u32.wgsl.expected.spvasm new file mode 100644 index 0000000000..ca38a9a94c --- /dev/null +++ b/test/var/override/numbered/zero_init/u32.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 7 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %o "o" + OpName %main "main" + OpDecorate %o SpecId 1234 + %uint = OpTypeInt 32 0 + %o = OpSpecConstant %uint 0 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %main = OpFunction %void None %3 + %6 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/var/override/numbered/zero_init/u32.wgsl.expected.wgsl b/test/var/override/numbered/zero_init/u32.wgsl.expected.wgsl new file mode 100644 index 0000000000..d45213b780 --- /dev/null +++ b/test/var/override/numbered/zero_init/u32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[override(1234)]] let o : u32 = u32(); + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = o; +}