blob: 94cbe057dd91259af47adaf1d4c755e5679a98aa [file] [log] [blame]
// Copyright 2023 The Dawn & Tint Authors
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are met:
//
// 1. Redistributions of source code must retain the above copyright notice, this
// list of conditions and the following disclaimer.
//
// 2. Redistributions in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// 3. Neither the name of the copyright holder nor the names of its
// contributors may be used to endorse or promote products derived from
// this software without specific prior written permission.
//
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
#include "src/tint/lang/spirv/reader/parser/helper_test.h"
namespace tint::spirv::reader {
namespace {
TEST_F(SpirvParserTest, FunctionVar) {
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
%void = OpTypeVoid
%u32 = OpTypeInt 32 0
%u32_ptr = OpTypePointer Function %u32
%ep_type = OpTypeFunction %void
%main = OpFunction %void None %ep_type
%main_start = OpLabel
%var = OpVariable %u32_ptr Function
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:ptr<function, u32, read_write> = var
ret
}
}
)");
}
TEST_F(SpirvParserTest, FunctionVar_Initializer) {
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
%void = OpTypeVoid
%u32 = OpTypeInt 32 0
%u32_42 = OpConstant %u32 42
%u32_ptr = OpTypePointer Function %u32
%ep_type = OpTypeFunction %void
%main = OpFunction %void None %ep_type
%main_start = OpLabel
%var = OpVariable %u32_ptr Function %u32_42
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:ptr<function, u32, read_write> = var, 42u
ret
}
}
)");
}
TEST_F(SpirvParserTest, PrivateVar) {
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
%void = OpTypeVoid
%u32 = OpTypeInt 32 0
%u32_ptr = OpTypePointer Private %u32
%ep_type = OpTypeFunction %void
%var = OpVariable %u32_ptr Private
%main = OpFunction %void None %ep_type
%main_start = OpLabel
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%1:ptr<private, u32, read_write> = var
}
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B2: {
ret
}
}
)");
}
TEST_F(SpirvParserTest, PrivateVar_Initializer) {
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
%void = OpTypeVoid
%u32 = OpTypeInt 32 0
%u32_42 = OpConstant %u32 42
%u32_ptr = OpTypePointer Private %u32
%ep_type = OpTypeFunction %void
%var = OpVariable %u32_ptr Private %u32_42
%main = OpFunction %void None %ep_type
%main_start = OpLabel
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%1:ptr<private, u32, read_write> = var, 42u
}
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B2: {
ret
}
}
)");
}
TEST_F(SpirvParserTest, StorageVar_ReadOnly) {
EXPECT_IR(R"(
OpCapability Shader
OpExtension "SPV_KHR_storage_buffer_storage_class"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %1 "main"
OpExecutionMode %1 LocalSize 1 1 1
OpDecorate %str Block
OpMemberDecorate %str 0 Offset 0
OpDecorate %6 NonWritable
OpDecorate %6 DescriptorSet 1
OpDecorate %6 Binding 2
%void = OpTypeVoid
%uint = OpTypeInt 32 0
%str = OpTypeStruct %uint
%_ptr_StorageBuffer_str = OpTypePointer StorageBuffer %str
%5 = OpTypeFunction %void
%6 = OpVariable %_ptr_StorageBuffer_str StorageBuffer
%1 = OpFunction %void None %5
%7 = OpLabel
OpReturn
OpFunctionEnd
)",
R"(
tint_symbol_1 = struct @align(4) {
tint_symbol:u32 @offset(0)
}
$B1: { # root
%1:ptr<storage, tint_symbol_1, read> = var @binding_point(1, 2)
}
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B2: {
ret
}
}
)");
}
TEST_F(SpirvParserTest, StorageVar_ReadWrite) {
EXPECT_IR(R"(
OpCapability Shader
OpExtension "SPV_KHR_storage_buffer_storage_class"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %1 "main"
OpExecutionMode %1 LocalSize 1 1 1
OpDecorate %str Block
OpMemberDecorate %str 0 Offset 0
OpDecorate %6 DescriptorSet 1
OpDecorate %6 Binding 2
%void = OpTypeVoid
%uint = OpTypeInt 32 0
%str = OpTypeStruct %uint
%_ptr_StorageBuffer_str = OpTypePointer StorageBuffer %str
%5 = OpTypeFunction %void
%6 = OpVariable %_ptr_StorageBuffer_str StorageBuffer
%1 = OpFunction %void None %5
%7 = OpLabel
OpReturn
OpFunctionEnd
)",
R"(
tint_symbol_1 = struct @align(4) {
tint_symbol:u32 @offset(0)
}
$B1: { # root
%1:ptr<storage, tint_symbol_1, read_write> = var @binding_point(1, 2)
}
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B2: {
ret
}
}
)");
}
TEST_F(SpirvParserTest, StorageVar_ReadOnly_And_ReadWrite) {
EXPECT_IR(R"(
OpCapability Shader
OpExtension "SPV_KHR_storage_buffer_storage_class"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %1 "main"
OpExecutionMode %1 LocalSize 1 1 1
OpDecorate %str Block
OpMemberDecorate %str 0 Offset 0
OpDecorate %6 NonWritable
OpDecorate %6 DescriptorSet 1
OpDecorate %6 Binding 2
OpDecorate %7 DescriptorSet 1
OpDecorate %7 Binding 3
%void = OpTypeVoid
%uint = OpTypeInt 32 0
%str = OpTypeStruct %uint
%_ptr_StorageBuffer_str = OpTypePointer StorageBuffer %str
%5 = OpTypeFunction %void
%6 = OpVariable %_ptr_StorageBuffer_str StorageBuffer
%7 = OpVariable %_ptr_StorageBuffer_str StorageBuffer
%1 = OpFunction %void None %5
%8 = OpLabel
OpReturn
OpFunctionEnd
)",
R"(
tint_symbol_1 = struct @align(4) {
tint_symbol:u32 @offset(0)
}
$B1: { # root
%1:ptr<storage, tint_symbol_1, read> = var @binding_point(1, 2)
%2:ptr<storage, tint_symbol_1, read_write> = var @binding_point(1, 3)
}
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B2: {
ret
}
}
)");
}
TEST_F(SpirvParserTest, UniformVar) {
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %1 "main"
OpExecutionMode %1 LocalSize 1 1 1
OpDecorate %str Block
OpMemberDecorate %str 0 Offset 0
OpDecorate %6 NonWritable
OpDecorate %6 DescriptorSet 1
OpDecorate %6 Binding 2
%void = OpTypeVoid
%uint = OpTypeInt 32 0
%str = OpTypeStruct %uint
%_ptr_Uniform_str = OpTypePointer Uniform %str
%5 = OpTypeFunction %void
%6 = OpVariable %_ptr_Uniform_str Uniform
%1 = OpFunction %void None %5
%7 = OpLabel
OpReturn
OpFunctionEnd
)",
R"(
tint_symbol_1 = struct @align(4) {
tint_symbol:u32 @offset(0)
}
$B1: { # root
%1:ptr<uniform, tint_symbol_1, read> = var @binding_point(1, 2)
}
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B2: {
ret
}
}
)");
}
struct BuiltinCase {
std::string spirv_type;
std::string spirv_builtin;
std::string ir;
};
std::string PrintBuiltinCase(testing::TestParamInfo<BuiltinCase> bc) {
return bc.param.spirv_builtin;
}
using BuiltinInputTest = SpirvParserTestWithParam<BuiltinCase>;
TEST_P(BuiltinInputTest, Enum) {
auto params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpCapability SampleRateShading
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
OpDecorate %var BuiltIn )" +
params.spirv_builtin + R"(
%void = OpTypeVoid
%bool = OpTypeBool
%u32 = OpTypeInt 32 0
%vec3u = OpTypeVector %u32 3
%f32 = OpTypeFloat 32
%vec4f = OpTypeVector %f32 4
%u32_1 = OpConstant %u32 1
%arr_u32_1 = OpTypeArray %u32 %u32_1
%fn_type = OpTypeFunction %void
%_ptr_Input = OpTypePointer Input %)" +
params.spirv_type + R"(
%var = OpVariable %_ptr_Input Input
%main = OpFunction %void None %fn_type
%main_start = OpLabel
OpReturn
OpFunctionEnd
)",
params.ir);
}
INSTANTIATE_TEST_SUITE_P(
SpirvParser,
BuiltinInputTest,
testing::Values(
BuiltinCase{
"vec4f",
"FragCoord",
"%1:ptr<__in, vec4<f32>, read> = var @builtin(position)",
},
BuiltinCase{
"bool",
"FrontFacing",
"%1:ptr<__in, bool, read> = var @builtin(front_facing)",
},
BuiltinCase{
"vec3u",
"GlobalInvocationId",
"%1:ptr<__in, vec3<u32>, read> = var @builtin(global_invocation_id)",
},
BuiltinCase{
"u32",
"InstanceIndex",
"%1:ptr<__in, u32, read> = var @builtin(instance_index)",
},
BuiltinCase{
"vec3u",
"LocalInvocationId",
"%1:ptr<__in, vec3<u32>, read> = var @builtin(local_invocation_id)",
},
BuiltinCase{
"u32",
"LocalInvocationIndex",
"%1:ptr<__in, u32, read> = var @builtin(local_invocation_index)",
},
BuiltinCase{
"vec3u",
"NumWorkgroups",
"%1:ptr<__in, vec3<u32>, read> = var @builtin(num_workgroups)",
},
BuiltinCase{
"u32",
"SampleId",
"%1:ptr<__in, u32, read> = var @builtin(sample_index)",
},
BuiltinCase{
"arr_u32_1",
"SampleMask",
"%1:ptr<__in, array<u32, 1>, read> = var @builtin(sample_mask)",
},
BuiltinCase{
"u32",
"VertexIndex",
"%1:ptr<__in, u32, read> = var @builtin(vertex_index)",
},
BuiltinCase{
"vec3u",
"WorkgroupId",
"%1:ptr<__in, vec3<u32>, read> = var @builtin(workgroup_id)",
}),
PrintBuiltinCase);
using BuiltinOutputTest = SpirvParserTestWithParam<BuiltinCase>;
TEST_P(BuiltinOutputTest, Enum) {
auto params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
OpDecorate %var BuiltIn )" +
params.spirv_builtin + R"(
%void = OpTypeVoid
%u32 = OpTypeInt 32 0
%f32 = OpTypeFloat 32
%vec4f = OpTypeVector %f32 4
%u32_1 = OpConstant %u32 1
%arr_u32_1 = OpTypeArray %u32 %u32_1
%fn_type = OpTypeFunction %void
%_ptr_Output = OpTypePointer Output %)" +
params.spirv_type + R"(
%var = OpVariable %_ptr_Output Output
%main = OpFunction %void None %fn_type
%main_start = OpLabel
OpReturn
OpFunctionEnd
)",
params.ir);
}
INSTANTIATE_TEST_SUITE_P(
SpirvParser,
BuiltinOutputTest,
testing::Values(
BuiltinCase{
"f32",
"PointSize",
"%1:ptr<__out, f32, read_write> = var @builtin(__point_size)",
},
BuiltinCase{
"vec4f",
"Position",
"%1:ptr<__out, vec4<f32>, read_write> = var @builtin(position)",
},
BuiltinCase{
"arr_u32_1",
"SampleMask",
"%1:ptr<__out, array<u32, 1>, read_write> = var @builtin(sample_mask)",
}),
PrintBuiltinCase);
TEST_F(SpirvParserTest, Invariant_OnVariable) {
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
OpDecorate %var BuiltIn Position
OpDecorate %var Invariant
%void = OpTypeVoid
%f32 = OpTypeFloat 32
%vec4f = OpTypeVector %f32 4
%fn_type = OpTypeFunction %void
%_ptr_Output = OpTypePointer Output %vec4f
%var = OpVariable %_ptr_Output Output
%main = OpFunction %void None %fn_type
%main_start = OpLabel
OpReturn
OpFunctionEnd
)",
"%1:ptr<__out, vec4<f32>, read_write> = var @invariant @builtin(position)");
}
struct LocationCase {
std::string spirv_decorations;
std::string ir;
};
using LocationVarTest = SpirvParserTestWithParam<LocationCase>;
TEST_P(LocationVarTest, Input) {
auto params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpCapability SampleRateShading
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
)" + params.spirv_decorations +
R"(
%void = OpTypeVoid
%f32 = OpTypeFloat 32
%vec4f = OpTypeVector %f32 4
%fn_type = OpTypeFunction %void
%_ptr_Input = OpTypePointer Input %vec4f
%var = OpVariable %_ptr_Input Input
%main = OpFunction %void None %fn_type
%main_start = OpLabel
OpReturn
OpFunctionEnd
)",
"%1:ptr<__in, vec4<f32>, read> = " + params.ir);
}
TEST_P(LocationVarTest, Output) {
auto params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpCapability SampleRateShading
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
)" + params.spirv_decorations +
R"(
%void = OpTypeVoid
%f32 = OpTypeFloat 32
%vec4f = OpTypeVector %f32 4
%fn_type = OpTypeFunction %void
%_ptr_Output = OpTypePointer Output %vec4f
%var = OpVariable %_ptr_Output Output
%main = OpFunction %void None %fn_type
%main_start = OpLabel
OpReturn
OpFunctionEnd
)",
"%1:ptr<__out, vec4<f32>, read_write> = " + params.ir);
}
INSTANTIATE_TEST_SUITE_P(SpirvParser,
LocationVarTest,
testing::Values(
LocationCase{
"OpDecorate %var Location 1 ",
"var @location(1)",
},
LocationCase{
"OpDecorate %var Location 2 "
"OpDecorate %var NoPerspective ",
"var @location(2) @interpolate(linear, center)",
},
LocationCase{
"OpDecorate %var Location 3 "
"OpDecorate %var Flat ",
"var @location(3) @interpolate(flat, center)",
},
LocationCase{
"OpDecorate %var Location 4 "
"OpDecorate %var Centroid ",
"var @location(4) @interpolate(perspective, centroid)",
},
LocationCase{
"OpDecorate %var Location 5 "
"OpDecorate %var Sample ",
"var @location(5) @interpolate(perspective, sample)",
},
LocationCase{
"OpDecorate %var Location 6 "
"OpDecorate %var NoPerspective "
"OpDecorate %var Centroid ",
"var @location(6) @interpolate(linear, centroid)",
}));
} // namespace
} // namespace tint::spirv::reader