encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

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 <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
This commit is contained in:
Ben Clayton 2021-10-26 10:05:02 +00:00 committed by Tint LUCI CQ
parent 0436d045ab
commit 6c4effe654
125 changed files with 1356 additions and 52 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

50
src/writer/spirv/builder.cc
View File

@ -577,7 +577,7 @@ uint32_t Builder::GenerateExpression(const ast::Expression* expr) {
return GenerateCallExpression(c);
}
if (auto* c = expr->As<ast::ConstructorExpression>()) {
return GenerateConstructorExpression(nullptr, c, false);
return GenerateConstructorExpression(nullptr, c);
}
if (auto* i = expr->As<ast::IdentifierExpression>()) {
return GenerateIdentifierExpression(i);
@ -763,7 +763,7 @@ bool Builder::GenerateGlobalVariable(const ast::Variable* var) {
}
init_id = GenerateConstructorExpression(
var, var->constructor->As<ast::ConstructorExpression>(), true);
var, var->constructor->As<ast::ConstructorExpression>());
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<ast::ScalarConstructorExpression>()) {
return GenerateLiteralIfNeeded(var, scalar->literal);
}
if (auto* type = expr->As<ast::TypeConstructorExpression>()) {
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<sem::GlobalVariable>(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<sem::I32>()) {
return GenerateConstantIfNeeded(
ScalarConstant::I32(0).AsSpecOp(constant_id));
}
if (result_type->Is<sem::U32>()) {
return GenerateConstantIfNeeded(
ScalarConstant::U32(0).AsSpecOp(constant_id));
}
if (result_type->Is<sem::F32>()) {
return GenerateConstantIfNeeded(
ScalarConstant::F32(0).AsSpecOp(constant_id));
}
if (result_type->Is<sem::Bool>()) {
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<ast::ConstructorExpression>(), is_global_init);
id = GenerateConstructorExpression(nullptr,
e->As<ast::ConstructorExpression>());
} 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;
}

13
src/writer/spirv/builder.h
View File

@ -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

58
src/writer/spirv/builder_constructor_expression_test.cc
View File

@ -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<f32>(1.0f, rel);
WrapInFunction(t);
auto* g = Global("g", ty.vec2<f32>(), 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<f32>(2.0f);
WrapInFunction(cast);
auto* g = Global("g", ty.vec2<f32>(), 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<f32>(2.0f);
WrapInFunction(cast);
auto* g = Global("g", ty.vec3<f32>(), 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<f32>(2.0f, vec2<f32>(2.0f, 2.0f));
WrapInFunction(cast);
auto* g = Global("g", ty.vec3<f32>(), 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<f32>(vec2<f32>(2.0f, 2.0f), 2.0f);
WrapInFunction(cast);
auto* g = Global("g", ty.vec3<f32>(), 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<f32>(2.0f);
WrapInFunction(cast);
auto* g = Global("g", ty.vec4<f32>(), 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<f32>(2.0f, 2.0f, vec2<f32>(2.0f, 2.0f));
WrapInFunction(cast);
auto* g = Global("g", ty.vec4<f32>(), 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<f32>(2.0f, vec2<f32>(2.0f, 2.0f), 2.0f);
WrapInFunction(cast);
auto* g = Global("g", ty.vec4<f32>(), 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<f32>(vec2<f32>(2.0f, 2.0f), 2.0f, 2.0f);
WrapInFunction(cast);
auto* g = Global("g", ty.vec4<f32>(), 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<f32>(vec2<f32>(2.0f, 2.0f), vec2<f32>(2.0f, 2.0f));
WrapInFunction(cast);
auto* g = Global("g", ty.vec4<f32>(), 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<f32>(2.0f, vec3<f32>(2.0f, 2.0f, 2.0f));
WrapInFunction(cast);
auto* g = Global("g", ty.vec4<f32>(), 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<f32>(vec3<f32>(2.0f, 2.0f, 2.0f), 2.0f);
WrapInFunction(cast);
auto* g = Global("g", ty.vec4<f32>(), 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<f32>();
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<bool>();
WrapInFunction(t);

36
src/writer/spirv/builder_global_variable_test.cc
View File

@ -167,6 +167,24 @@ TEST_F(BuilderTest, GlobalVar_Override_Bool) {
)");
}
TEST_F(BuilderTest, GlobalVar_Override_Bool_ZeroValue) {
auto* v = GlobalConst("var", ty.bool_(), Construct<bool>(),
ast::DecorationList{
create<ast::OverrideDecoration>(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<f32>(),
ast::DecorationList{
create<ast::OverrideDecoration>(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{

19
src/writer/spirv/scalar_constant.h
View File

@ -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

6
test/var/override/named/no_init/bool.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override]] let o : bool;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

16
test/var/override/named/no_init/bool.wgsl.expected.hlsl Normal file
View File

@ -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
^

8
test/var/override/named/no_init/bool.wgsl.expected.msl Normal file
View File

@ -0,0 +1,8 @@
#include <metal_stdlib>
using namespace metal;
constant bool o [[function_constant(0)]];
kernel void tint_symbol() {
return;
}

20
test/var/override/named/no_init/bool.wgsl.expected.spvasm Normal file
View File

@ -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

6
test/var/override/named/no_init/bool.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override]] let o : bool;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

6
test/var/override/named/no_init/f32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override]] let o : f32;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

16
test/var/override/named/no_init/f32.wgsl.expected.hlsl Normal file
View File

@ -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
^

8
test/var/override/named/no_init/f32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,8 @@
#include <metal_stdlib>
using namespace metal;
constant float o [[function_constant(0)]];
kernel void tint_symbol() {
return;
}

20
test/var/override/named/no_init/f32.wgsl.expected.spvasm Normal file
View File

@ -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

6
test/var/override/named/no_init/f32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override]] let o : f32;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

6
test/var/override/named/no_init/i32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override]] let o : i32;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

16
test/var/override/named/no_init/i32.wgsl.expected.hlsl Normal file
View File

@ -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
^

8
test/var/override/named/no_init/i32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,8 @@
#include <metal_stdlib>
using namespace metal;
constant int o [[function_constant(0)]];
kernel void tint_symbol() {
return;
}

20
test/var/override/named/no_init/i32.wgsl.expected.spvasm Normal file
View File

@ -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

6
test/var/override/named/no_init/i32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override]] let o : i32;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

6
test/var/override/named/no_init/u32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override]] let o : u32;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

16
test/var/override/named/no_init/u32.wgsl.expected.hlsl Normal file
View File

@ -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
^

8
test/var/override/named/no_init/u32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,8 @@
#include <metal_stdlib>
using namespace metal;
constant uint o [[function_constant(0)]];
kernel void tint_symbol() {
return;
}

20
test/var/override/named/no_init/u32.wgsl.expected.spvasm Normal file
View File

@ -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

6
test/var/override/named/no_init/u32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override]] let o : u32;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

6
test/var/override/named/val_init/bool.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override]] let o : bool = true;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

9
test/var/override/named/val_init/bool.wgsl.expected.hlsl Normal file
View File

@ -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;
}

8
test/var/override/named/val_init/bool.wgsl.expected.msl Normal file
View File

@ -0,0 +1,8 @@
#include <metal_stdlib>
using namespace metal;
constant bool o [[function_constant(0)]];
kernel void tint_symbol() {
return;
}

20
test/var/override/named/val_init/bool.wgsl.expected.spvasm Normal file
View File

@ -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

6
test/var/override/named/val_init/bool.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override]] let o : bool = true;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

6
test/var/override/named/val_init/f32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override]] let o : f32 = 1.0;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

9
test/var/override/named/val_init/f32.wgsl.expected.hlsl Normal file
View File

@ -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;
}

8
test/var/override/named/val_init/f32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,8 @@
#include <metal_stdlib>
using namespace metal;
constant float o [[function_constant(0)]];
kernel void tint_symbol() {
return;
}

20
test/var/override/named/val_init/f32.wgsl.expected.spvasm Normal file
View File

@ -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

6
test/var/override/named/val_init/f32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override]] let o : f32 = 1.0;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

6
test/var/override/named/val_init/i32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override]] let o : i32 = 1;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

9
test/var/override/named/val_init/i32.wgsl.expected.hlsl Normal file
View File

@ -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;
}

8
test/var/override/named/val_init/i32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,8 @@
#include <metal_stdlib>
using namespace metal;
constant int o [[function_constant(0)]];
kernel void tint_symbol() {
return;
}

20
test/var/override/named/val_init/i32.wgsl.expected.spvasm Normal file
View File

@ -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

6
test/var/override/named/val_init/i32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override]] let o : i32 = 1;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

6
test/var/override/named/val_init/u32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override]] let o : u32 = 1u;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

9
test/var/override/named/val_init/u32.wgsl.expected.hlsl Normal file
View File

@ -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;
}

8
test/var/override/named/val_init/u32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,8 @@
#include <metal_stdlib>
using namespace metal;
constant uint o [[function_constant(0)]];
kernel void tint_symbol() {
return;
}

20
test/var/override/named/val_init/u32.wgsl.expected.spvasm Normal file
View File

@ -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

6
test/var/override/named/val_init/u32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override]] let o : u32 = 1u;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

6
test/var/override/named/zero_init/bool.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override]] let o : bool = bool();
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

9
test/var/override/named/zero_init/bool.wgsl.expected.hlsl Normal file
View File

@ -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;
}

8
test/var/override/named/zero_init/bool.wgsl.expected.msl Normal file
View File

@ -0,0 +1,8 @@
#include <metal_stdlib>
using namespace metal;
constant bool o [[function_constant(0)]];
kernel void tint_symbol() {
return;
}

20
test/var/override/named/zero_init/bool.wgsl.expected.spvasm Normal file
View File

@ -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

6
test/var/override/named/zero_init/bool.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override]] let o : bool = bool();
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

6
test/var/override/named/zero_init/f32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override]] let o : f32 = f32();
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

9
test/var/override/named/zero_init/f32.wgsl.expected.hlsl Normal file
View File

@ -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;
}

8
test/var/override/named/zero_init/f32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,8 @@
#include <metal_stdlib>
using namespace metal;
constant float o [[function_constant(0)]];
kernel void tint_symbol() {
return;
}

20
test/var/override/named/zero_init/f32.wgsl.expected.spvasm Normal file
View File

@ -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

6
test/var/override/named/zero_init/f32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override]] let o : f32 = f32();
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

6
test/var/override/named/zero_init/i32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override]] let o : i32 = i32();
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

9
test/var/override/named/zero_init/i32.wgsl.expected.hlsl Normal file
View File

@ -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;
}

8
test/var/override/named/zero_init/i32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,8 @@
#include <metal_stdlib>
using namespace metal;
constant int o [[function_constant(0)]];
kernel void tint_symbol() {
return;
}

20
test/var/override/named/zero_init/i32.wgsl.expected.spvasm Normal file
View File

@ -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

6
test/var/override/named/zero_init/i32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override]] let o : i32 = i32();
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

6
test/var/override/named/zero_init/u32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override]] let o : u32 = u32();
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

9
test/var/override/named/zero_init/u32.wgsl.expected.hlsl Normal file
View File

@ -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;
}

8
test/var/override/named/zero_init/u32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,8 @@
#include <metal_stdlib>
using namespace metal;
constant uint o [[function_constant(0)]];
kernel void tint_symbol() {
return;
}

20
test/var/override/named/zero_init/u32.wgsl.expected.spvasm Normal file
View File

@ -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

6
test/var/override/named/zero_init/u32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override]] let o : u32 = u32();
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

6
test/var/override/numbered/no_init/bool.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override(1234)]] let o : bool;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

16
test/var/override/numbered/no_init/bool.wgsl.expected.hlsl Normal file
View File

@ -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
^

8
test/var/override/numbered/no_init/bool.wgsl.expected.msl Normal file
View File

@ -0,0 +1,8 @@
#include <metal_stdlib>
using namespace metal;
constant bool o [[function_constant(1234)]];
kernel void tint_symbol() {
return;
}

20
test/var/override/numbered/no_init/bool.wgsl.expected.spvasm Normal file
View File

@ -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

6
test/var/override/numbered/no_init/bool.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override(1234)]] let o : bool;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

6
test/var/override/numbered/no_init/f32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override(1234)]] let o : f32;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

16
test/var/override/numbered/no_init/f32.wgsl.expected.hlsl Normal file
View File

@ -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
^

8
test/var/override/numbered/no_init/f32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,8 @@
#include <metal_stdlib>
using namespace metal;
constant float o [[function_constant(1234)]];
kernel void tint_symbol() {
return;
}

20
test/var/override/numbered/no_init/f32.wgsl.expected.spvasm Normal file
View File

@ -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

6
test/var/override/numbered/no_init/f32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override(1234)]] let o : f32;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

6
test/var/override/numbered/no_init/i32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override(1234)]] let o : i32;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

16
test/var/override/numbered/no_init/i32.wgsl.expected.hlsl Normal file
View File

@ -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
^

8
test/var/override/numbered/no_init/i32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,8 @@
#include <metal_stdlib>
using namespace metal;
constant int o [[function_constant(1234)]];
kernel void tint_symbol() {
return;
}

20
test/var/override/numbered/no_init/i32.wgsl.expected.spvasm Normal file
View File

@ -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

6
test/var/override/numbered/no_init/i32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override(1234)]] let o : i32;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

6
test/var/override/numbered/no_init/u32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override(1234)]] let o : u32;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

16
test/var/override/numbered/no_init/u32.wgsl.expected.hlsl Normal file
View File

@ -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
^

8
test/var/override/numbered/no_init/u32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,8 @@
#include <metal_stdlib>
using namespace metal;
constant uint o [[function_constant(1234)]];
kernel void tint_symbol() {
return;
}

20
test/var/override/numbered/no_init/u32.wgsl.expected.spvasm Normal file
View File

@ -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

6
test/var/override/numbered/no_init/u32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override(1234)]] let o : u32;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

6
test/var/override/numbered/val_init/bool.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override(1234)]] let o : bool = true;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

9
test/var/override/numbered/val_init/bool.wgsl.expected.hlsl Normal file
View File

@ -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;
}

8
test/var/override/numbered/val_init/bool.wgsl.expected.msl Normal file
View File

@ -0,0 +1,8 @@
#include <metal_stdlib>
using namespace metal;
constant bool o [[function_constant(1234)]];
kernel void tint_symbol() {
return;
}

20
test/var/override/numbered/val_init/bool.wgsl.expected.spvasm Normal file
View File

@ -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

6
test/var/override/numbered/val_init/bool.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override(1234)]] let o : bool = true;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

6
test/var/override/numbered/val_init/f32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override(1234)]] let o : f32 = 1.0;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

9
test/var/override/numbered/val_init/f32.wgsl.expected.hlsl Normal file
View File

@ -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;
}

8
test/var/override/numbered/val_init/f32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,8 @@
#include <metal_stdlib>
using namespace metal;
constant float o [[function_constant(1234)]];
kernel void tint_symbol() {
return;
}

20
test/var/override/numbered/val_init/f32.wgsl.expected.spvasm Normal file
View File

@ -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

6
test/var/override/numbered/val_init/f32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override(1234)]] let o : f32 = 1.0;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

6
test/var/override/numbered/val_init/i32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override(1234)]] let o : i32 = 1;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

9
test/var/override/numbered/val_init/i32.wgsl.expected.hlsl Normal file
View File

@ -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;
}

8
test/var/override/numbered/val_init/i32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,8 @@
#include <metal_stdlib>
using namespace metal;
constant int o [[function_constant(1234)]];
kernel void tint_symbol() {
return;
}

20
test/var/override/numbered/val_init/i32.wgsl.expected.spvasm Normal file
View File

@ -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

6
test/var/override/numbered/val_init/i32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[override(1234)]] let o : i32 = 1;
[[stage(compute), workgroup_size(1)]]
fn main() {
_ = o;
}

Some files were not shown because too many files have changed in this diff Show More