Google Git
Sign in
dawn / dawn / 9cf619c021907af643d7ab9cf128e9fe02362ed7 / . / src / tint / lang / spirv / reader / texture_test.cc
blob: 24d69ccc7ab91d06ee07dbe8bf1ca1337e1c404d [file] [log] [blame]
// Copyright 2025 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/helper_test.h"
namespace tint::spirv::reader {
namespace {
TEST_F(SpirvReaderTest, Handle_Sampler) {
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical Simple
OpEntryPoint Fragment %main "main"
OpExecutionMode %main OriginUpperLeft
OpName %10 "wg"
OpDecorate %10 DescriptorSet 0
OpDecorate %10 Binding 0
%sampler = OpTypeSampler
%ptr = OpTypePointer UniformConstant %sampler
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%10 = OpVariable %ptr UniformConstant
%main = OpFunction %void None %voidfn
%entry = OpLabel
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%wg:ptr<handle, sampler, read> = var undef @binding_point(0, 0)
}
%main = @fragment func():void {
$B2: {
ret
}
}
)");
}
TEST_F(SpirvReaderTest, Handle_Image_Texture) {
EXPECT_IR(R"(
OpCapability Shader
OpCapability Sampled1D
OpMemoryModel Logical Simple
OpEntryPoint Fragment %main "main"
OpExecutionMode %main OriginUpperLeft
OpName %10 "wg"
OpDecorate %10 DescriptorSet 0
OpDecorate %10 Binding 0
%float = OpTypeFloat 32
%tex = OpTypeImage %float 1D 0 0 0 1 Unknown
%ptr_tex = OpTypePointer UniformConstant %tex
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%10 = OpVariable %ptr_tex UniformConstant
%main = OpFunction %void None %voidfn
%entry = OpLabel
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%wg:ptr<handle, texture_1d<f32>, read> = var undef @binding_point(0, 0)
}
%main = @fragment func():void {
$B2: {
ret
}
}
)");
}
TEST_F(SpirvReaderTest, Handle_Image_NonWritable) {
EXPECT_IR(R"(
OpCapability Shader
OpCapability Sampled1D
OpCapability StorageImageExtendedFormats
OpMemoryModel Logical Simple
OpEntryPoint Fragment %main "main"
OpExecutionMode %main OriginUpperLeft
OpName %10 "wg"
OpDecorate %10 DescriptorSet 0
OpDecorate %10 Binding 0
OpDecorate %10 NonWritable
%float = OpTypeFloat 32
%storage = OpTypeImage %float 1D 0 0 0 2 Rg32f
%ptr_storage = OpTypePointer UniformConstant %storage
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%10 = OpVariable %ptr_storage UniformConstant
%main = OpFunction %void None %voidfn
%entry = OpLabel
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%wg:ptr<handle, texture_storage_1d<rg32float, read>, read> = var undef @binding_point(0, 0)
}
%main = @fragment func():void {
$B2: {
ret
}
}
)");
}
TEST_F(SpirvReaderTest, Handle_Image_NonReadable) {
EXPECT_IR(R"(
OpCapability Shader
OpCapability Sampled1D
OpCapability StorageImageExtendedFormats
OpMemoryModel Logical Simple
OpEntryPoint Fragment %main "main"
OpExecutionMode %main OriginUpperLeft
OpName %10 "wg"
OpDecorate %10 DescriptorSet 0
OpDecorate %10 Binding 0
OpDecorate %10 NonReadable
%float = OpTypeFloat 32
%storage = OpTypeImage %float 1D 0 0 0 2 Rg32f
%ptr_storage = OpTypePointer UniformConstant %storage
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%10 = OpVariable %ptr_storage UniformConstant
%main = OpFunction %void None %voidfn
%entry = OpLabel
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%wg:ptr<handle, texture_storage_1d<rg32float, write>, read> = var undef @binding_point(0, 0)
}
%main = @fragment func():void {
$B2: {
ret
}
}
)");
}
TEST_F(SpirvReaderTest, Handle_MS_2D) {
EXPECT_IR(R"(
OpCapability Shader
OpCapability ImageQuery
OpMemoryModel Logical Simple
OpEntryPoint Fragment %main "main"
OpExecutionMode %main OriginUpperLeft
OpName %10 "wg"
OpDecorate %10 DescriptorSet 0
OpDecorate %10 Binding 0
%float = OpTypeFloat 32
%uint = OpTypeInt 32 0
%tex = OpTypeImage %float 2D 0 0 1 1 Unknown
%ptr_tex = OpTypePointer UniformConstant %tex
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%10 = OpVariable %ptr_tex UniformConstant
%main = OpFunction %void None %voidfn
%entry = OpLabel
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%wg:ptr<handle, texture_multisampled_2d<f32>, read> = var undef @binding_point(0, 0)
}
%main = @fragment func():void {
$B2: {
ret
}
}
)");
}
TEST_F(SpirvReaderTest, Handle_Depth_MS_2D) {
EXPECT_IR(R"(
OpCapability Shader
OpCapability ImageQuery
OpMemoryModel Logical Simple
OpEntryPoint Fragment %main "main"
OpExecutionMode %main OriginUpperLeft
OpName %10 "wg"
OpDecorate %10 DescriptorSet 0
OpDecorate %10 Binding 0
%float = OpTypeFloat 32
%uint = OpTypeInt 32 0
%tex = OpTypeImage %float 2D 1 0 1 1 Unknown
%ptr_tex = OpTypePointer UniformConstant %tex
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%10 = OpVariable %ptr_tex UniformConstant
%main = OpFunction %void None %voidfn
%entry = OpLabel
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%wg:ptr<handle, texture_depth_multisampled_2d, read> = var undef @binding_point(0, 0)
}
%main = @fragment func():void {
$B2: {
ret
}
}
)");
}
TEST_F(SpirvReaderTest, ImageGather) {
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical Simple
OpEntryPoint Fragment %main "main"
OpExecutionMode %main OriginUpperLeft
OpName %10 "wg"
OpDecorate %var_sampler DescriptorSet 0
OpDecorate %var_sampler Binding 0
OpDecorate %10 DescriptorSet 2
OpDecorate %10 Binding 0
%int = OpTypeInt 32 1
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%v4float = OpTypeVector %float 4
%tex = OpTypeImage %float 2D 0 0 0 1 Unknown
%ptr_tex = OpTypePointer UniformConstant %tex
%sampled_img = OpTypeSampledImage %tex
%sampler = OpTypeSampler
%ptr_sampler = OpTypePointer UniformConstant %sampler
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%int_1 = OpConstant %int 1
%float_1 = OpConstant %float 1
%float_2 = OpConstant %float 2
%coords2 = OpConstantComposite %v2float %float_1 %float_2
%var_sampler = OpVariable %ptr_sampler UniformConstant
%10 = OpVariable %ptr_tex UniformConstant
%main = OpFunction %void None %voidfn
%entry = OpLabel
%sam = OpLoad %sampler %var_sampler
%im = OpLoad %tex %10
%sampled_image = OpSampledImage %sampled_img %im %sam
%result = OpImageGather %v4float %sampled_image %coords2 %int_1
%r2 = OpFAdd %v4float %result %result
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%1:ptr<handle, sampler, read> = var undef @binding_point(0, 0)
%wg:ptr<handle, texture_2d<f32>, read> = var undef @binding_point(2, 0)
}
%main = @fragment func():void {
$B2: {
%4:sampler = load %1
%5:texture_2d<f32> = load %wg
%6:vec4<f32> = textureGather 1i, %5, %4, vec2<f32>(1.0f, 2.0f)
%7:vec4<f32> = add %6, %6
ret
}
}
)");
}
TEST_F(SpirvReaderTest, ImageSampleImplicitLod) {
EXPECT_IR(R"(
OpCapability Shader
OpCapability Sampled1D
OpMemoryModel Logical Simple
OpEntryPoint Fragment %main "main"
OpExecutionMode %main OriginUpperLeft
OpName %10 "wg"
OpDecorate %var_sampler DescriptorSet 0
OpDecorate %var_sampler Binding 0
OpDecorate %10 DescriptorSet 2
OpDecorate %10 Binding 0
%int = OpTypeInt 32 1
%float = OpTypeFloat 32
%v2int = OpTypeVector %int 2
%v2float = OpTypeVector %float 2
%v4float = OpTypeVector %float 4
%tex = OpTypeImage %float 2D 0 0 0 1 Unknown
%ptr_tex = OpTypePointer UniformConstant %tex
%sampled_img = OpTypeSampledImage %tex
%sampler = OpTypeSampler
%ptr_sampler = OpTypePointer UniformConstant %sampler
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%int_1 = OpConstant %int 1
%float_1 = OpConstant %float 1
%float_2 = OpConstant %float 2
%coords2 = OpConstantComposite %v2float %float_1 %float_2
%var_sampler = OpVariable %ptr_sampler UniformConstant
%10 = OpVariable %ptr_tex UniformConstant
%int_10 = OpConstant %int 10
%int_11 = OpConstant %int 11
%offset2i = OpConstantComposite %v2int %int_10 %int_11
%main = OpFunction %void None %voidfn
%entry = OpLabel
%sam = OpLoad %sampler %var_sampler
%im = OpLoad %tex %10
%sampled_image = OpSampledImage %sampled_img %im %sam
%result = OpImageSampleImplicitLod %v4float %sampled_image %coords2 ConstOffset %offset2i
%r2 = OpFAdd %v4float %result %result
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%1:ptr<handle, sampler, read> = var undef @binding_point(0, 0)
%wg:ptr<handle, texture_2d<f32>, read> = var undef @binding_point(2, 0)
}
%main = @fragment func():void {
$B2: {
%4:sampler = load %1
%5:texture_2d<f32> = load %wg
%6:vec4<f32> = textureSample %5, %4, vec2<f32>(1.0f, 2.0f), vec2<i32>(10i, 11i)
%7:vec4<f32> = add %6, %6
ret
}
}
)");
}
TEST_F(SpirvReaderTest, ImageSampleExplicitLod) {
EXPECT_IR(R"(
OpCapability Shader
OpCapability Sampled1D
OpMemoryModel Logical Simple
OpEntryPoint Fragment %main "main"
OpExecutionMode %main OriginUpperLeft
OpName %10 "wg"
OpDecorate %var_sampler DescriptorSet 0
OpDecorate %var_sampler Binding 0
OpDecorate %10 DescriptorSet 2
OpDecorate %10 Binding 0
%int = OpTypeInt 32 1
%float = OpTypeFloat 32
%v2int = OpTypeVector %int 2
%v2float = OpTypeVector %float 2
%v4float = OpTypeVector %float 4
%tex = OpTypeImage %float 2D 0 0 0 1 Unknown
%ptr_tex = OpTypePointer UniformConstant %tex
%sampled_img = OpTypeSampledImage %tex
%sampler = OpTypeSampler
%ptr_sampler = OpTypePointer UniformConstant %sampler
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%int_1 = OpConstant %int 1
%float_1 = OpConstant %float 1
%float_2 = OpConstant %float 2
%coords2 = OpConstantComposite %v2float %float_1 %float_2
%var_sampler = OpVariable %ptr_sampler UniformConstant
%10 = OpVariable %ptr_tex UniformConstant
%int_10 = OpConstant %int 10
%int_11 = OpConstant %int 11
%offset2i = OpConstantComposite %v2int %int_10 %int_11
%main = OpFunction %void None %voidfn
%entry = OpLabel
%sam = OpLoad %sampler %var_sampler
%im = OpLoad %tex %10
%sampled_image = OpSampledImage %sampled_img %im %sam
%result = OpImageSampleExplicitLod %v4float %sampled_image %coords2 Lod|ConstOffset %float_1 %offset2i
%r2 = OpFAdd %v4float %result %result
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%1:ptr<handle, sampler, read> = var undef @binding_point(0, 0)
%wg:ptr<handle, texture_2d<f32>, read> = var undef @binding_point(2, 0)
}
%main = @fragment func():void {
$B2: {
%4:sampler = load %1
%5:texture_2d<f32> = load %wg
%6:vec4<f32> = textureSampleLevel %5, %4, vec2<f32>(1.0f, 2.0f), 1.0f, vec2<i32>(10i, 11i)
%7:vec4<f32> = add %6, %6
ret
}
}
)");
}
TEST_F(SpirvReaderTest, ImageQuerySize) {
EXPECT_IR(R"(
OpCapability Shader
OpCapability Sampled1D
OpCapability ImageQuery
OpMemoryModel Logical Simple
OpEntryPoint Fragment %main "main"
OpExecutionMode %main OriginUpperLeft
OpName %wg "wg"
OpDecorate %wg DescriptorSet 2
OpDecorate %wg Binding 0
%int = OpTypeInt 32 1
%float = OpTypeFloat 32
%tex = OpTypeImage %float 1D 0 0 0 2 R32f
%ptr_tex = OpTypePointer UniformConstant %tex
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%wg = OpVariable %ptr_tex UniformConstant
%main = OpFunction %void None %voidfn
%entry = OpLabel
%im = OpLoad %tex %wg
%result = OpImageQuerySize %int %im
%r2 = OpIAdd %int %result %result
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%wg:ptr<handle, texture_storage_1d<r32float, read_write>, read> = var undef @binding_point(2, 0)
}
%main = @fragment func():void {
$B2: {
%3:texture_storage_1d<r32float, read_write> = load %wg
%4:u32 = textureDimensions %3
%5:i32 = convert %4
%6:i32 = add %5, %5
ret
}
}
)");
}
TEST_F(SpirvReaderTest, ImageQueryLevels) {
EXPECT_IR(R"(
OpCapability Shader
OpCapability Sampled1D
OpCapability ImageQuery
OpMemoryModel Logical Simple
OpEntryPoint Fragment %main "main"
OpExecutionMode %main OriginUpperLeft
OpName %wg "wg"
OpDecorate %wg DescriptorSet 2
OpDecorate %wg Binding 0
%int = OpTypeInt 32 1
%float = OpTypeFloat 32
%tex = OpTypeImage %float 1D 0 0 0 1 R32f
%ptr_tex = OpTypePointer UniformConstant %tex
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%wg = OpVariable %ptr_tex UniformConstant
%main = OpFunction %void None %voidfn
%entry = OpLabel
%im = OpLoad %tex %wg
%result = OpImageQueryLevels %int %im
%r2 = OpIAdd %int %result %result
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%wg:ptr<handle, texture_1d<f32>, read> = var undef @binding_point(2, 0)
}
%main = @fragment func():void {
$B2: {
%3:texture_1d<f32> = load %wg
%4:u32 = textureNumLevels %3
%5:i32 = convert %4
%6:i32 = add %5, %5
ret
}
}
)");
}
TEST_F(SpirvReaderTest, ImageQuerySamples) {
EXPECT_IR(R"(
OpCapability Shader
OpCapability Sampled1D
OpCapability ImageQuery
OpMemoryModel Logical Simple
OpEntryPoint Fragment %main "main"
OpExecutionMode %main OriginUpperLeft
OpName %wg "wg"
OpDecorate %wg DescriptorSet 2
OpDecorate %wg Binding 0
%int = OpTypeInt 32 1
%float = OpTypeFloat 32
%tex = OpTypeImage %float 2D 0 0 1 1 R32f
%ptr_tex = OpTypePointer UniformConstant %tex
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%wg = OpVariable %ptr_tex UniformConstant
%main = OpFunction %void None %voidfn
%entry = OpLabel
%im = OpLoad %tex %wg
%result = OpImageQuerySamples %int %im
%r2 = OpIAdd %int %result %result
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%wg:ptr<handle, texture_multisampled_2d<f32>, read> = var undef @binding_point(2, 0)
}
%main = @fragment func():void {
$B2: {
%3:texture_multisampled_2d<f32> = load %wg
%4:u32 = textureNumSamples %3
%5:i32 = convert %4
%6:i32 = add %5, %5
ret
}
}
)");
}
TEST_F(SpirvReaderTest, ImageQuerySizeLod) {
EXPECT_IR(R"(
OpCapability Shader
OpCapability Sampled1D
OpCapability ImageQuery
OpMemoryModel Logical Simple
OpEntryPoint Fragment %main "main"
OpExecutionMode %main OriginUpperLeft
OpName %wg "wg"
OpDecorate %wg DescriptorSet 2
OpDecorate %wg Binding 0
%int = OpTypeInt 32 1
%float = OpTypeFloat 32
%tex = OpTypeImage %float 1D 0 0 0 1 R32f
%ptr_tex = OpTypePointer UniformConstant %tex
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%wg = OpVariable %ptr_tex UniformConstant
%int_1 = OpConstant %int 1
%main = OpFunction %void None %voidfn
%entry = OpLabel
%im = OpLoad %tex %wg
%result = OpImageQuerySizeLod %int %im %int_1
%r2 = OpIAdd %int %result %result
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%wg:ptr<handle, texture_1d<f32>, read> = var undef @binding_point(2, 0)
}
%main = @fragment func():void {
$B2: {
%3:texture_1d<f32> = load %wg
%4:u32 = textureDimensions %3, 1i
%5:i32 = convert %4
%6:i32 = add %5, %5
ret
}
}
)");
}
struct ImgData {
std::string name;
std::string spirv_type;
std::string spirv_fn;
std::string wgsl_type;
std::string wgsl_fn;
};
[[maybe_unused]] std::ostream& operator<<(std::ostream& out, const ImgData& d) {
out << d.name;
return out;
}
using SamplerTest = SpirvReaderTestWithParam<ImgData>;
TEST_P(SamplerTest, Handle) {
auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpCapability Sampled1D
OpMemoryModel Logical Simple
OpEntryPoint Fragment %main "main"
OpExecutionMode %main OriginUpperLeft
OpName %10 "wg"
OpDecorate %var_sampler DescriptorSet 0
OpDecorate %var_sampler Binding 0
OpDecorate %10 DescriptorSet 2
OpDecorate %10 Binding 0
%int = OpTypeInt 32 1
%uint = OpTypeInt 32 0
%float = OpTypeFloat 32
%v2int = OpTypeVector %int 2
%v2uint = OpTypeVector %uint 2
%v2float = OpTypeVector %float 2
%v3float = OpTypeVector %float 3
%v4float = OpTypeVector %float 4
%tex = OpTypeImage )" +
params.spirv_type + R"(
%ptr_tex = OpTypePointer UniformConstant %tex
%sampled_img = OpTypeSampledImage %tex
%sampler = OpTypeSampler
%ptr_sampler = OpTypePointer UniformConstant %sampler
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%int_1 = OpConstant %int 1
%float_null = OpConstantNull %float
%float_1 = OpConstant %float 1
%float_2 = OpConstant %float 2
%float_3 = OpConstant %float 3
%float_4 = OpConstant %float 4
%float_5 = OpConstant %float 5
%coords2 = OpConstantComposite %v2float %float_1 %float_2
%coords3 = OpConstantComposite %v3float %float_1 %float_2 %float_3
%coords4 = OpConstantComposite %v4float %float_1 %float_2 %float_3 %float_4
%vf12 = OpConstantComposite %v2float %float_1 %float_2
%vf21 = OpConstantComposite %v2float %float_2 %float_1
%int_10 = OpConstant %int 10
%int_11 = OpConstant %int 11
%offset2i = OpConstantComposite %v2int %int_10 %int_11
%uint_20 = OpConstant %uint 20
%uint_21 = OpConstant %uint 21
%offset2u = OpConstantComposite %v2uint %uint_20 %uint_21
%depth = OpConstant %float 1
%var_sampler = OpVariable %ptr_sampler UniformConstant
%10 = OpVariable %ptr_tex UniformConstant
%main = OpFunction %void None %voidfn
%entry = OpLabel
%sam = OpLoad %sampler %var_sampler
%im = OpLoad %tex %10
%sampled_image = OpSampledImage %sampled_img %im %sam
%result = )" + params.spirv_fn +
R"(
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%1:ptr<handle, sampler, read> = var undef @binding_point(0, 0)
%wg:ptr<handle, )" +
params.wgsl_type + R"(, read> = var undef @binding_point(2, 0)
}
%main = @fragment func():void {
$B2: {
%4:sampler = load %1
%5:)" + params.wgsl_type +
R"( = load %wg)" + params.wgsl_fn +
R"(
ret
}
}
)");
}
INSTANTIATE_TEST_SUITE_P(
SpirvReaderTest_ImageGather,
SamplerTest,
::testing::Values(
ImgData{
.name = "2D",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "OpImageGather %v4float %sampled_image %coords2 %int_1",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%6:vec4<f32> = textureGather 1i, %5, %4, vec2<f32>(1.0f, 2.0f))",
},
ImgData{
.name = "2D ConstOffset Signed",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn =
"OpImageGather %v4float %sampled_image %coords2 %int_1 ConstOffset %offset2i",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%6:vec4<f32> = textureGather 1i, %5, %4, vec2<f32>(1.0f, 2.0f), vec2<i32>(10i, 11i))",
},
ImgData{
.name = "2D ConstOffset Unsigned",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn =
"OpImageGather %v4float %sampled_image %coords2 %int_1 ConstOffset %offset2u",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%6:vec2<i32> = convert vec2<u32>(20u, 21u)
%7:vec4<f32> = textureGather 1i, %5, %4, vec2<f32>(1.0f, 2.0f), %6)",
},
ImgData{
.name = "2D Array",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn = "OpImageGather %v4float %sampled_image %coords3 %int_1",
.wgsl_type = "texture_2d_array<f32>",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:i32 = convert %7
%9:vec4<f32> = textureGather 1i, %5, %4, %6, %8)",
},
ImgData{
.name = "2D Array ConstOffset Signed",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn =
"OpImageGather %v4float %sampled_image %coords3 %int_1 ConstOffset %offset2i",
.wgsl_type = "texture_2d_array<f32>",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:i32 = convert %7
%9:vec4<f32> = textureGather 1i, %5, %4, %6, %8, vec2<i32>(10i, 11i))",
},
ImgData{
.name = "2D Array ConstOffset Unsigned",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn =
"OpImageGather %v4float %sampled_image %coords3 %int_1 ConstOffset %offset2u",
.wgsl_type = "texture_2d_array<f32>",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:i32 = convert %7
%9:vec2<i32> = convert vec2<u32>(20u, 21u)
%10:vec4<f32> = textureGather 1i, %5, %4, %6, %8, %9)",
},
ImgData{
.name = "2D Depth",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "OpImageGather %v4float %sampled_image %coords2 %int_1",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = R"(
%6:vec4<f32> = textureGather %5, %4, vec2<f32>(1.0f, 2.0f))",
},
ImgData{
.name = "2D Depth ConstOffset Signed",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn =
"OpImageGather %v4float %sampled_image %coords2 %int_1 ConstOffset %offset2i",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = R"(
%6:vec4<f32> = textureGather %5, %4, vec2<f32>(1.0f, 2.0f), vec2<i32>(10i, 11i))",
},
ImgData{
.name = "2D Depth ConstOffset Unsigned",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn =
"OpImageGather %v4float %sampled_image %coords2 %int_1 ConstOffset %offset2u",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = R"(
%6:vec2<i32> = convert vec2<u32>(20u, 21u)
%7:vec4<f32> = textureGather %5, %4, vec2<f32>(1.0f, 2.0f), %6)",
},
ImgData{
.name = "2D Depth Array",
.spirv_type = "%float 2D 1 1 0 1 Unknown",
.spirv_fn = "OpImageGather %v4float %sampled_image %coords3 %int_1",
.wgsl_type = "texture_depth_2d_array",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:i32 = convert %7
%9:vec4<f32> = textureGather %5, %4, %6, %8)",
},
ImgData{
.name = "2D Depth Array ConstOffset Signed",
.spirv_type = "%float 2D 1 1 0 1 Unknown",
.spirv_fn =
"OpImageGather %v4float %sampled_image %coords3 %int_1 ConstOffset %offset2i",
.wgsl_type = "texture_depth_2d_array",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:i32 = convert %7
%9:vec4<f32> = textureGather %5, %4, %6, %8, vec2<i32>(10i, 11i))",
},
ImgData{
.name = "2D Depth Array ConstOffset Unsigned",
.spirv_type = "%float 2D 1 1 0 1 Unknown",
.spirv_fn =
"OpImageGather %v4float %sampled_image %coords3 %int_1 ConstOffset %offset2u",
.wgsl_type = "texture_depth_2d_array",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:i32 = convert %7
%9:vec2<i32> = convert vec2<u32>(20u, 21u)
%10:vec4<f32> = textureGather %5, %4, %6, %8, %9)",
},
ImgData{
.name = "Cube",
.spirv_type = "%float Cube 0 0 0 1 Unknown",
.spirv_fn = "OpImageGather %v4float %sampled_image %coords3 %int_1",
.wgsl_type = "texture_cube<f32>",
.wgsl_fn = R"(
%6:vec4<f32> = textureGather 1i, %5, %4, vec3<f32>(1.0f, 2.0f, 3.0f))",
},
ImgData{
.name = "Cube Depth",
.spirv_type = "%float Cube 1 0 0 1 Unknown",
.spirv_fn = "OpImageGather %v4float %sampled_image %coords3 %int_1",
.wgsl_type = "texture_depth_cube",
.wgsl_fn = R"(
%6:vec4<f32> = textureGather %5, %4, vec3<f32>(1.0f, 2.0f, 3.0f))",
},
ImgData{
.name = "Cube Array",
.spirv_type = "%float Cube 0 1 0 1 Unknown",
.spirv_fn = "OpImageGather %v4float %sampled_image %coords4 %int_1",
.wgsl_type = "texture_cube_array<f32>",
.wgsl_fn = R"(
%6:vec3<f32> = swizzle vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f), xyz
%7:f32 = swizzle vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f), w
%8:i32 = convert %7
%9:vec4<f32> = textureGather 1i, %5, %4, %6, %8)",
},
ImgData{
.name = "Cube Depth Array",
.spirv_type = "%float Cube 1 1 0 1 Unknown",
.spirv_fn = "OpImageGather %v4float %sampled_image %coords4 %int_1",
.wgsl_type = "texture_depth_cube_array",
.wgsl_fn = R"(
%6:vec3<f32> = swizzle vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f), xyz
%7:f32 = swizzle vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f), w
%8:i32 = convert %7
%9:vec4<f32> = textureGather %5, %4, %6, %8)",
}));
INSTANTIATE_TEST_SUITE_P(
DISABLED_SpirvReaderTest_ImageDrefGather,
SamplerTest,
::testing::Values(
ImgData{
.name = "2d Depth",
.spirv_type = "%float 2d 1 0 0 1 Unknown",
.spirv_fn = "OpImageDrefGather %v4float %sampled_image %coords2 %depth",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = "textureGatherCompare %4, %5, vec2<f32>(1.0f, 2.0f), 1.0f",
},
ImgData{
.name = "2d Depth ConstOffset Signed",
.spirv_type = "%float 2d 1 0 0 1 Unknown",
.spirv_fn =
"OpImageDrefGather %v4float %sampled_image %coords2 %depth ConstOffset %offset2i",
.wgsl_type = "texture_depth_2d",
.wgsl_fn =
"textureGatherCompare %4, %5, vec2<f32>(1.0f, 2.0f), 1.0f, vec2<i32>(10i, 11i)",
},
ImgData{
.name = "2d Depth ConstOffset Unsigned",
.spirv_type = "%float 2d 1 0 0 1 Unknown",
.spirv_fn =
"OpImageDrefGather %v4float %sampled_image %coords2 %depth ConstOfset %offset2u",
.wgsl_type = "texture_depth_2d",
.wgsl_fn =
"textureGatherCompare %4, %5, vec2<f32>(1.0f, 2.0f), 1.0f, vec2<u32>(20u, 21u)",
},
ImgData{
.name = "2d Depth Array",
.spirv_type = "%float 2d 1 1 0 1 Unknown",
.spirv_fn = "OpImageDrefGather %v4float %sampled_image %coords3 %depth",
.wgsl_type = "texture_depth_2d_array",
.wgsl_fn = "textureGatherCompare %4, %5, vec2<f32>(1.0f, 2.0f), 3i, 1.0f",
},
ImgData{
.name = "2d Depth Array ConstOffset Signed",
.spirv_type = "%float 2d 1 1 0 1 Unknown",
.spirv_fn =
"OpImageDrefGather %v4float %sampled_image %coords3 %depth ConstOffset %offset2i",
.wgsl_type = "texture_depth_2_arrayd",
.wgsl_fn =
"textureGatherCompare %4, %5, vec2<f32>(1.0f, 2.0f), 3i, 1.0f, vec2<i32>(10i, 11i)",
},
ImgData{
.name = "2d Depth Array ConstOffset Unsigned",
.spirv_type = "%float 2d 1 1 0 1 Unknown",
.spirv_fn =
"OpImageDrefGather %v4float %sampled_image %coords3 %depth ConstOfset %offset2u",
.wgsl_type = "texture_depth_2d_array",
.wgsl_fn =
"textureGatherCompare %4, %5, vec2<f32>(1.0f, 2.0f), 3i, 1.0f, vec2<u32>(20u, 21u)",
},
ImgData{
.name = "Cube Depth",
.spirv_type = "%float Cube 1 0 0 1 Unknown",
.spirv_fn = "OpImageDrefGather %v4float %sampled_image %coords3 %depth",
.wgsl_type = "texture_depth_cube",
.wgsl_fn = "textureGatherCompare %4, %5, vec3<f32>(1.0f, 2.0f, 3.0f), 1.0f",
},
ImgData{
.name = "Cube Depth Array",
.spirv_type = "%float Cube 1 1 0 1 Unknown",
.spirv_fn = "OpImageDrefGather %v4float %sampled_image %coords4 %depth",
.wgsl_type = "texture_depth_cube_array",
.wgsl_fn = "textureGatherCompare %4, %5, vec3<f32>(1.0f, 2.0f, 3.0f), 4i, 1.0f",
}));
INSTANTIATE_TEST_SUITE_P(
SpirvReaderTest_ImageSampleImplicitLod,
SamplerTest,
::testing::Values(
ImgData{
.name = "2D",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "OpImageSampleImplicitLod %v4float %sampled_image %coords2",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%6:vec4<f32> = textureSample %5, %4, vec2<f32>(1.0f, 2.0f))",
},
ImgData{
.name = "2D ConstOffset",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn =
"OpImageSampleImplicitLod %v4float %sampled_image %coords2 ConstOffset %offset2i",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%6:vec4<f32> = textureSample %5, %4, vec2<f32>(1.0f, 2.0f), vec2<i32>(10i, 11i))",
},
ImgData{
.name = "2D Bias",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "OpImageSampleImplicitLod %v4float %sampled_image %coords2 Bias %float_5",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%6:vec4<f32> = textureSampleBias %5, %4, vec2<f32>(1.0f, 2.0f), 5.0f)",
},
ImgData{
.name = "2D Bias ConstOffset Signed",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "OpImageSampleImplicitLod %v4float %sampled_image %coords2 "
"Bias|ConstOffset %float_5 %offset2i",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%6:vec4<f32> = textureSampleBias %5, %4, vec2<f32>(1.0f, 2.0f), 5.0f, vec2<i32>(10i, 11i))",
},
ImgData{
.name = "2D Bias ConstOffset Unsigned",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "OpImageSampleImplicitLod %v4float %sampled_image %coords2 "
"Bias|ConstOffset %float_5 %offset2u",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%6:vec2<i32> = convert vec2<u32>(20u, 21u)
%7:vec4<f32> = textureSampleBias %5, %4, vec2<f32>(1.0f, 2.0f), 5.0f, %6)",
},
ImgData{
.name = "2D Array",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn = "OpImageSampleImplicitLod %v4float %sampled_image %coords3",
.wgsl_type = "texture_2d_array<f32>",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:i32 = convert %7
%9:vec4<f32> = textureSample %5, %4, %6, %8)",
},
ImgData{
.name = "2D Array ConstOffset",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn =
"OpImageSampleImplicitLod %v4float %sampled_image %coords3 ConstOffset %offset2i",
.wgsl_type = "texture_2d_array<f32>",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:i32 = convert %7
%9:vec4<f32> = textureSample %5, %4, %6, %8, vec2<i32>(10i, 11i))",
},
ImgData{
.name = "2D Array Bias",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn = "OpImageSampleImplicitLod %v4float %sampled_image %coords3 Bias %float_5",
.wgsl_type = "texture_2d_array<f32>",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:i32 = convert %7
%9:vec4<f32> = textureSampleBias %5, %4, %6, %8, 5.0f)",
},
ImgData{
.name = "2D Array Bias ConstOffset",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn = "OpImageSampleImplicitLod %v4float %sampled_image %coords3 "
"Bias|ConstOffset %float_5 %offset2i",
.wgsl_type = "texture_2d_array<f32>",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:i32 = convert %7
%9:vec4<f32> = textureSampleBias %5, %4, %6, %8, 5.0f, vec2<i32>(10i, 11i))",
},
ImgData{
.name = "2D Array Bias ConstOffset",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn = "OpImageSampleImplicitLod %v4float %sampled_image %coords3 "
"Bias|ConstOffset %float_5 %offset2i",
.wgsl_type = "texture_2d_array<f32>",
.wgsl_fn =
R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:i32 = convert %7
%9:vec4<f32> = textureSampleBias %5, %4, %6, %8, 5.0f, vec2<i32>(10i, 11i))",
}));
INSTANTIATE_TEST_SUITE_P(
SpirvReaderTest_ImageSampleProjImplicitLod,
SamplerTest,
::testing::Values(
ImgData{
.name = "1D",
.spirv_type = "%float 1D 0 0 0 1 Unknown",
.spirv_fn = "OpImageSampleProjImplicitLod %v4float %sampled_image %coords2",
.wgsl_type = "texture_1d<f32>",
.wgsl_fn = R"(
%6:f32 = swizzle vec2<f32>(1.0f, 2.0f), x
%7:f32 = swizzle vec2<f32>(1.0f, 2.0f), y
%8:f32 = div %6, %7
%9:vec4<f32> = textureSample %5, %4, %8)",
},
ImgData{
.name = "2D",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "OpImageSampleProjImplicitLod %v4float %sampled_image %coords3",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:vec2<f32> = div %6, %7
%9:vec4<f32> = textureSample %5, %4, %8)",
},
ImgData{
.name = "3D",
.spirv_type = "%float 3D 0 0 0 1 Unknown",
.spirv_fn = "OpImageSampleProjImplicitLod %v4float %sampled_image %coords4",
.wgsl_type = "texture_3d<f32>",
.wgsl_fn = R"(
%6:vec3<f32> = swizzle vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f), xyz
%7:f32 = swizzle vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f), w
%8:vec3<f32> = div %6, %7
%9:vec4<f32> = textureSample %5, %4, %8)",
},
ImgData{
.name = "2D ConstOffset",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "OpImageSampleProjImplicitLod %v4float %sampled_image %coords3 ConstOffset "
"%offset2i",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:vec2<f32> = div %6, %7
%9:vec4<f32> = textureSample %5, %4, %8, vec2<i32>(10i, 11i))",
},
ImgData{
.name = "2D Bias",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn =
"OpImageSampleProjImplicitLod %v4float %sampled_image %coords3 Bias %float_5",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:vec2<f32> = div %6, %7
%9:vec4<f32> = textureSampleBias %5, %4, %8, 5.0f)",
},
ImgData{
.name = "2D Bias ConstOffset Signed",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "OpImageSampleProjImplicitLod %v4float %sampled_image %coords3 "
"Bias|ConstOffset %float_5 %offset2i",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn =
R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:vec2<f32> = div %6, %7
%9:vec4<f32> = textureSampleBias %5, %4, %8, 5.0f, vec2<i32>(10i, 11i))",
},
ImgData{
.name = "2D Bias ConstOffset Unsigned",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "OpImageSampleProjImplicitLod %v4float %sampled_image %coords3 "
"Bias|ConstOffset %float_5 %offset2u",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn =
R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:vec2<f32> = div %6, %7
%9:vec2<i32> = convert vec2<u32>(20u, 21u)
%10:vec4<f32> = textureSampleBias %5, %4, %8, 5.0f, %9)",
},
ImgData{
.name = "2D Depth",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "OpImageSampleProjImplicitLod %v4float %sampled_image %coords3",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:vec2<f32> = div %6, %7
%9:f32 = textureSample %5, %4, %8
%10:vec4<f32> = construct %9, 0.0f, 0.0f, 0.0f)",
}));
INSTANTIATE_TEST_SUITE_P(
SpirvReaderTest_ImageSampleExplicitLod,
SamplerTest,
::testing::Values(
ImgData{
.name = "2D Lod",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "OpImageSampleExplicitLod %v4float %sampled_image %coords2 Lod %float_null",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%6:vec4<f32> = textureSampleLevel %5, %4, vec2<f32>(1.0f, 2.0f), 0.0f)",
},
ImgData{
.name = "2D Array Lod",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn = "OpImageSampleExplicitLod %v4float %sampled_image %coords3 Lod %float_null",
.wgsl_type = "texture_2d_array<f32>",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:i32 = convert %7
%9:vec4<f32> = textureSampleLevel %5, %4, %6, %8, 0.0f)",
},
ImgData{
.name = "2D Lod ConstOffset signed",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "OpImageSampleExplicitLod %v4float %sampled_image %coords2 Lod|ConstOffset "
"%float_null %offset2i",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%6:vec4<f32> = textureSampleLevel %5, %4, vec2<f32>(1.0f, 2.0f), 0.0f, vec2<i32>(10i, 11i))",
},
ImgData{
.name = "2D lod ConstOffset unsigned",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "OpImageSampleExplicitLod %v4float %sampled_image %coords2 Lod|ConstOffset "
"%float_null %offset2u",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%6:vec2<i32> = convert vec2<u32>(20u, 21u)
%7:vec4<f32> = textureSampleLevel %5, %4, vec2<f32>(1.0f, 2.0f), 0.0f, %6)",
},
ImgData{
.name = "2D Array lod ConstOffset",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn = "OpImageSampleExplicitLod %v4float %sampled_image %coords3 Lod|ConstOffset "
"%float_null %offset2i",
.wgsl_type = "texture_2d_array<f32>",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:i32 = convert %7
%9:vec4<f32> = textureSampleLevel %5, %4, %6, %8, 0.0f, vec2<i32>(10i, 11i))",
},
ImgData{
.name = "2D Grad",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn =
"OpImageSampleExplicitLod %v4float %sampled_image %coords2 Grad %vf12 %vf21",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%6:vec4<f32> = textureSampleGrad %5, %4, vec2<f32>(1.0f, 2.0f), vec2<f32>(1.0f, 2.0f), vec2<f32>(2.0f, 1.0f))",
},
ImgData{
.name = "2D Array Grad",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn =
"OpImageSampleExplicitLod %v4float %sampled_image %coords3 Grad %vf12 %vf21",
.wgsl_type = "texture_2d_array<f32>",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:i32 = convert %7
%9:vec4<f32> = textureSampleGrad %5, %4, %6, %8, vec2<f32>(1.0f, 2.0f), vec2<f32>(2.0f, 1.0f))",
},
ImgData{
.name = "2D Grad ConstOffset signed",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "OpImageSampleExplicitLod %v4float %sampled_image %coords2 "
"Grad|ConstOffset %vf12 %vf21 %offset2i",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%6:vec4<f32> = textureSampleGrad %5, %4, vec2<f32>(1.0f, 2.0f), vec2<f32>(1.0f, 2.0f), vec2<f32>(2.0f, 1.0f), vec2<i32>(10i, 11i))",
},
ImgData{
.name = "2D Grad ConstOffset Unsigned",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "OpImageSampleExplicitLod %v4float %sampled_image %coords2 "
"Grad|ConstOffset %vf12 %vf21 %offset2u",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%6:vec2<i32> = convert vec2<u32>(20u, 21u)
%7:vec4<f32> = textureSampleGrad %5, %4, vec2<f32>(1.0f, 2.0f), vec2<f32>(1.0f, 2.0f), vec2<f32>(2.0f, 1.0f), %6)",
},
ImgData{
.name = "2D Array Grad ConstOffset",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn = "OpImageSampleExplicitLod %v4float %sampled_image %coords3 "
"Grad|ConstOffset %vf12 %vf21 %offset2i",
.wgsl_type = "texture_2d_array<f32>",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:i32 = convert %7
%9:vec4<f32> = textureSampleGrad %5, %4, %6, %8, vec2<f32>(1.0f, 2.0f), vec2<f32>(2.0f, 1.0f), vec2<i32>(10i, 11i))",
},
ImgData{
.name = "2D Array Grad ConstOffset Unsigned",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn = "OpImageSampleExplicitLod %v4float %sampled_image %coords3 "
"Grad|ConstOffset %vf12 %vf21 %offset2u",
.wgsl_type = "texture_2d_array<f32>",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:i32 = convert %7
%9:vec2<i32> = convert vec2<u32>(20u, 21u)
%10:vec4<f32> = textureSampleGrad %5, %4, %6, %8, vec2<f32>(1.0f, 2.0f), vec2<f32>(2.0f, 1.0f), %9)",
},
ImgData{
.name = "2D Depth",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "OpImageSampleExplicitLod %v4float %sampled_image %vf12 Lod %float_1",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = R"(
%6:i32 = convert 1.0f
%7:f32 = textureSampleLevel %5, %4, vec2<f32>(1.0f, 2.0f), %6
%8:vec4<f32> = construct %7, 0.0f, 0.0f, 0.0f)",
}));
INSTANTIATE_TEST_SUITE_P(
SpirvReaderTest_ImageSampleProjExplicitLod,
SamplerTest,
::testing::Values(
ImgData{
.name = "2D",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn =
"OpImageSampleProjExplicitLod %v4float %sampled_image %coords3 Lod %float_1",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:vec2<f32> = div %6, %7
%9:vec4<f32> = textureSampleLevel %5, %4, %8, 1.0f)",
},
ImgData{
.name = "2D Lod ConstOffset",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "OpImageSampleProjExplicitLod %v4float %sampled_image %coords3 "
"Lod|ConstOffset %float_1 %offset2i",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:vec2<f32> = div %6, %7
%9:vec4<f32> = textureSampleLevel %5, %4, %8, 1.0f, vec2<i32>(10i, 11i))",
},
ImgData{
.name = "2D Grad",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn =
"OpImageSampleProjExplicitLod %v4float %sampled_image %coords3 Grad %vf12 %vf21",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:vec2<f32> = div %6, %7
%9:vec4<f32> = textureSampleGrad %5, %4, %8, vec2<f32>(1.0f, 2.0f), vec2<f32>(2.0f, 1.0f))",
},
ImgData{
.name = "2D Lod Grad ConstOffset",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "OpImageSampleProjExplicitLod %v4float %sampled_image %coords3 "
"Grad|ConstOffset %vf12 %vf21 %offset2i",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:vec2<f32> = div %6, %7
%9:vec4<f32> = textureSampleGrad %5, %4, %8, vec2<f32>(1.0f, 2.0f), vec2<f32>(2.0f, 1.0f), vec2<i32>(10i, 11i))",
}));
using SamplerComparisonTest = SpirvReaderTestWithParam<ImgData>;
TEST_P(SamplerComparisonTest, Handle) {
auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpCapability Sampled1D
OpMemoryModel Logical Simple
OpEntryPoint Fragment %main "main"
OpExecutionMode %main OriginUpperLeft
OpName %10 "wg"
OpDecorate %var_sampler DescriptorSet 0
OpDecorate %var_sampler Binding 0
OpDecorate %10 DescriptorSet 2
OpDecorate %10 Binding 0
%int = OpTypeInt 32 1
%uint = OpTypeInt 32 0
%float = OpTypeFloat 32
%v2int = OpTypeVector %int 2
%v2uint = OpTypeVector %uint 2
%v2float = OpTypeVector %float 2
%v3float = OpTypeVector %float 3
%v4float = OpTypeVector %float 4
%tex = OpTypeImage )" +
params.spirv_type +
R"(
%ptr_tex = OpTypePointer UniformConstant %tex
%sampled_img = OpTypeSampledImage %tex
%sampler = OpTypeSampler
%ptr_sampler = OpTypePointer UniformConstant %sampler
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%int_1 = OpConstant %int 1
%float_0 = OpConstant %float 0
%float_1 = OpConstant %float 1
%float_2 = OpConstant %float 2
%float_3 = OpConstant %float 3
%float_4 = OpConstant %float 4
%float_5 = OpConstant %float 5
%coords2 = OpConstantComposite %v2float %float_1 %float_2
%coords3 = OpConstantComposite %v3float %float_1 %float_2 %float_3
%coords4 = OpConstantComposite %v4float %float_1 %float_2 %float_3 %float_4
%int_10 = OpConstant %int 10
%int_11 = OpConstant %int 11
%offset2i = OpConstantComposite %v2int %int_10 %int_11
%uint_20 = OpConstant %uint 20
%uint_21 = OpConstant %uint 21
%offset2u = OpConstantComposite %v2uint %uint_20 %uint_21
%depth = OpConstant %float 1
%var_sampler = OpVariable %ptr_sampler UniformConstant
%10 = OpVariable %ptr_tex UniformConstant
%main = OpFunction %void None %voidfn
%entry = OpLabel
%sam = OpLoad %sampler %var_sampler
%im = OpLoad %tex %10
%sampled_image = OpSampledImage %sampled_img %im %sam
%result = )" + params.spirv_fn +
R"(
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%1:ptr<handle, sampler_comparison, read> = var undef @binding_point(0, 0)
%wg:ptr<handle, )" +
params.wgsl_type +
R"(, read> = var undef @binding_point(2, 0)
}
%main = @fragment func():void {
$B2: {
%4:sampler_comparison = load %1
%5:)" + params.wgsl_type +
R"( = load %wg)" + params.wgsl_fn +
R"(
ret
}
}
)");
}
INSTANTIATE_TEST_SUITE_P(
SpirvReaderTest_ImageSampleDrefImplicitLod,
SamplerComparisonTest,
::testing::Values(
ImgData{
.name = "2D",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "OpImageSampleDrefImplicitLod %float %sampled_image %coords2 %depth",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = R"(
%6:f32 = textureSampleCompare %5, %4, vec2<f32>(1.0f, 2.0f), 1.0f)",
},
ImgData{
.name = "2D ConstOffset",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "OpImageSampleDrefImplicitLod %float %sampled_image %coords2 %depth "
"ConstOffset %offset2i",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = R"(
%6:f32 = textureSampleCompare %5, %4, vec2<f32>(1.0f, 2.0f), 1.0f, vec2<i32>(10i, 11i))",
},
ImgData{
.name = "2D Array",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn = "OpImageSampleDrefImplicitLod %float %sampled_image %coords3 %depth",
.wgsl_type = "texture_depth_2d_array",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:i32 = convert %7
%9:f32 = textureSampleCompare %5, %4, %6, %8, 1.0f)",
},
ImgData{
.name = "2D Array ConstOffset",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn = "OpImageSampleDrefImplicitLod %float %sampled_image %coords3 %depth "
"ConstOffset %offset2i",
.wgsl_type = "texture_depth_2d_array",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:i32 = convert %7
%9:f32 = textureSampleCompare %5, %4, %6, %8, 1.0f, vec2<i32>(10i, 11i))",
}));
INSTANTIATE_TEST_SUITE_P(
DISABLED_SpirvReaderTest_ImageSampleDrefExplicitLod,
SamplerComparisonTest,
::testing::Values(
ImgData{
.name = "2D",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "OpImageSampleDrefExplicitLod %float %sampled_image %coords2 "
"%depth Lod %float_0",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = "textureSampleCompareLevel %4, %5, vec2<f32>(1.0f, 2.0f), 1.0f",
},
ImgData{
.name = "2D array",
.spirv_type = "%float 2D 1 1 0 1 Unknown",
.spirv_fn = "OpImageSampleDrefExplicitLod %float %sampled_image %coords3 "
"%depth Lod %float_0",
.wgsl_type = "texture_depth_2d_array",
.wgsl_fn = "textureSampleCompareLevel %4, %5, vec2<f32>(1.0f, 2.0f), 3i, 1.0f",
},
ImgData{
.name = "2D ConstOffset",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "OpImageSampleDrefExplicitLod %float %sampled_image %coords2 %depth "
"Lod|ConstOffset %float_0 %offset2i",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = "textureSampleCompareLevel %4, %5, vec2<f32>(1.0f, 2.0f), 1.0f, "
"vec2<i32>(10i, 11i)",
},
ImgData{
.name = "2D array ConstOffset",
.spirv_type = "%float 2D 1 1 0 1 Unknown",
.spirv_fn = "OpImageSampleDrefExplicitLod %float %sampled_image %coords3 %depth "
"Lod|ConstOffset %float_0 %offset2i",
.wgsl_type = "texture_depth_2d_array",
.wgsl_fn = "textureSampleCompareLevel %4, %5, vec2<f32>(1.0f, 2.0f), 3i, 1.0f, "
"vec2<i32>(10i, 11i)",
},
ImgData{
.name = "Cube",
.spirv_type = "%float Cube 1 0 0 1 Unknown",
.spirv_fn = "OpImageSampleDrefExplicitLod %float %sampled_image %coords3 "
"%depth Lod %float_0",
.wgsl_type = "texture_depth_cube",
.wgsl_fn = "textureSampleCompareLevel %4, %5, vec3<f32>(1.0f, 2.0f, 3.0f), 1.0f",
},
ImgData{
.name = "Cube array",
.spirv_type = "%float Cube 1 1 0 1 Unknown",
.spirv_fn = "OpImageSampleDrefExplicitLod %float %sampled_image %coords4 "
"%depth Lod %float_0",
.wgsl_type = "texture_depth_cube_array",
.wgsl_fn = "textureSampleCompareLevel %4, %5, vec3<f32>(1.0f, 2.0f, 3.0f), "
"4i, 1.0f",
}));
INSTANTIATE_TEST_SUITE_P(
DISABLED_SpirvReaderTest_ImageSampleProjDrefImplicitLod,
SamplerComparisonTest,
::testing::Values(
ImgData{
.name = "2D Depth",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "OpImageSampleProjDrefImplicitLod %float %sampled_image %coords3 %float_1",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = "textureSampleCompare %4, %5, vec3<f32>(1.0f, 2.0f, 3.0f), 1.0f",
},
ImgData{
.name = "2D Depth ConstOffset",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "OpImageSampleProjDrefImplicitLod %float %sampled_image %coords3 %float_1 "
"ConstOffset %offset2i",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = "textureSampleCompare %4, %5, vec3<f32>(1.0f, 2.0f, 3.0f), 1.0f, "
"vec2<i32>(10i, 11i)",
}));
INSTANTIATE_TEST_SUITE_P(
DISABLED_SpirvReaderTest_ImageSampleProjDrefExplicitLod,
SamplerComparisonTest,
::testing::Values(
ImgData{
.name = "2D Depth Lod",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "OpImageSampleProjDrefExplicitLod %float %sampled_image %coords3 %float_1 "
"Lod %float_0",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = "textureSampleCompare %4, %5, vec3<f32>(1.0f, 2.0f, 3.0f), 1.0f, 0.0f",
},
ImgData{
.name = "2D Depth Lod ConstOffset",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "OpImageSampleProjDrefExplicitLod %float %sampled_image %coords3 %float_1 "
"Lod|ConstOffset %float_0 %offset2i",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = "textureSampleCompareLevel %4, %5, vec3<f32>(1.0f, 2.0f, 3.0f), 1.0f, 0.0f, "
"vec2<i32>(10i, 11i)",
}));
// This test shows the use of a sampled image used with both regular
// sampling and depth-reference sampling. The texture is a depth-texture,
// and we use builtins textureSample and textureSampleCompare
TEST_F(SpirvReaderTest, ImageSampleImplicitLod_BothDrefAndNonDref) {
EXPECT_IR(R"(
OpCapability Shader
OpCapability Sampled1D
OpMemoryModel Logical Simple
OpEntryPoint Fragment %main "main"
OpExecutionMode %main OriginUpperLeft
OpName %10 "wg"
OpDecorate %var_sampler1 DescriptorSet 0
OpDecorate %var_sampler1 Binding 0
OpDecorate %var_sampler2 DescriptorSet 0
OpDecorate %var_sampler2 Binding 1
OpDecorate %10 DescriptorSet 2
OpDecorate %10 Binding 0
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%v4float = OpTypeVector %float 4
%tex = OpTypeImage %float 2D 0 0 0 1 Unknown
%ptr_tex = OpTypePointer UniformConstant %tex
%sampled_img = OpTypeSampledImage %tex
%sampler = OpTypeSampler
%ptr_sampler = OpTypePointer UniformConstant %sampler
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%float_0 = OpConstant %float 0
%float_1 = OpConstant %float 1
%float_2 = OpConstant %float 2
%coords2 = OpConstantComposite %v2float %float_1 %float_2
%depth = OpConstant %float 1
%var_sampler1 = OpVariable %ptr_sampler UniformConstant
%var_sampler2 = OpVariable %ptr_sampler UniformConstant
%10 = OpVariable %ptr_tex UniformConstant
%main = OpFunction %void None %voidfn
%entry = OpLabel
%sam = OpLoad %sampler %var_sampler1
%im = OpLoad %tex %10
%sampled_image = OpSampledImage %sampled_img %im %sam
%sam_dref = OpLoad %sampler %var_sampler2
%sampled_dref_image = OpSampledImage %sampled_img %im %sam_dref
%200 = OpImageSampleImplicitLod %v4float %sampled_image %coords2
%210 = OpImageSampleDrefImplicitLod %float %sampled_dref_image %coords2 %depth
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%1:ptr<handle, sampler, read> = var undef @binding_point(0, 0)
%2:ptr<handle, sampler_comparison, read> = var undef @binding_point(0, 1)
%wg:ptr<handle, texture_depth_2d, read> = var undef @binding_point(2, 0)
}
%main = @fragment func():void {
$B2: {
%5:sampler = load %1
%6:texture_depth_2d = load %wg
%7:sampler_comparison = load %2
%8:f32 = textureSample %6, %5, vec2<f32>(1.0f, 2.0f)
%9:vec4<f32> = construct %8, 0.0f, 0.0f, 0.0f
%10:f32 = textureSampleCompare %6, %7, vec2<f32>(1.0f, 2.0f), 1.0f
ret
}
}
)");
}
using ImageAccessTest = SpirvReaderTestWithParam<ImgData>;
TEST_P(ImageAccessTest, Variable) {
auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpCapability StorageImageExtendedFormats
OpMemoryModel Logical Simple
OpEntryPoint Fragment %main "main"
OpExecutionMode %main OriginUpperLeft
OpName %main "main"
OpName %wg "wg"
OpName %offset2i "offset2i"
OpDecorate %wg DescriptorSet 2
OpDecorate %wg Binding 1
%void = OpTypeVoid
%float = OpTypeFloat 32
%int = OpTypeInt 32 1
%uint = OpTypeInt 32 0
%v2int = OpTypeVector %int 2
%v3int = OpTypeVector %int 3
%v4int = OpTypeVector %int 4
%v2uint = OpTypeVector %uint 2
%v3uint = OpTypeVector %uint 3
%v4uint = OpTypeVector %uint 4
%v2float = OpTypeVector %float 2
%v3float = OpTypeVector %float 3
%v4float = OpTypeVector %float 4
%im_ty = OpTypeImage )" +
params.spirv_type +
R"(
%ptr_im_ty = OpTypePointer UniformConstant %im_ty
%voidfn = OpTypeFunction %void
%int_1 = OpConstant %int 1
%int_2 = OpConstant %int 2
%int_3 = OpConstant %int 3
%int_4 = OpConstant %int 4
%int_10 = OpConstant %int 10
%int_11 = OpConstant %int 11
%uint_1 = OpConstant %uint 1
%uint_2 = OpConstant %uint 2
%uint_3 = OpConstant %uint 3
%uint_4 = OpConstant %uint 4
%float_1 = OpConstant %float 1
%float_2 = OpConstant %float 2
%float_3 = OpConstant %float 3
%float_4 = OpConstant %float 4
%wg = OpVariable %ptr_im_ty UniformConstant
%offset2i = OpConstantComposite %v2int %int_10 %int_11
%vi12 = OpConstantComposite %v2int %int_1 %int_2
%vi123 = OpConstantComposite %v3int %int_1 %int_2 %int_3
%vi1234 = OpConstantComposite %v4int %int_1 %int_2 %int_3 %int_4
%vf12 = OpConstantComposite %v2float %float_1 %float_2
%vf123 = OpConstantComposite %v3float %float_1 %float_2 %float_3
%vf1234 = OpConstantComposite %v4float %float_1 %float_2 %float_3 %float_4
%vu12 = OpConstantComposite %v2uint %uint_1 %uint_2
%vu123 = OpConstantComposite %v3uint %uint_1 %uint_2 %uint_3
%vu1234 = OpConstantComposite %v4uint %uint_1 %uint_2 %uint_3 %uint_4
%main = OpFunction %void None %voidfn
%entry = OpLabel
%im = OpLoad %im_ty %wg
)" + params.spirv_fn +
R"(
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%wg:ptr<handle, )" +
params.wgsl_type +
R"(, read> = var undef @binding_point(2, 1)
}
%main = @fragment func():void {
$B2: {
%3:)" + params.wgsl_type +
R"( = load %wg)" + params.wgsl_fn +
R"(
ret
}
}
)");
}
INSTANTIATE_TEST_SUITE_P(SpirvReaderTest_ImageWrite_OptionalParams,
ImageAccessTest,
::testing::Values(ImgData{
.name = "No extra params",
.spirv_type = "%float 2D 0 0 0 2 Rgba32f",
.spirv_fn = "OpImageWrite %im %vi12 %vf1234",
.wgsl_type = "texture_storage_2d<rgba32float, read_write>",
.wgsl_fn = R"(
%4:void = textureStore %3, vec2<i32>(1i, 2i), vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f))",
}));
// SPIR-V's texel parameter is a scalar or vector with at least as many
// components as there are channels in the underlying format, and the
// component type matches the sampled type (modulo signed/unsigned integer).
// WGSL's texel parameter is a 4-element vector scalar or vector, with
// component type equal to the 32-bit form of the channel type.
INSTANTIATE_TEST_SUITE_P(
SpirvReaderTest_ImageWrite_ConvertTexelOperand_Arity_Float,
ImageAccessTest,
::testing::Values(ImgData{.name = "Source 1 component",
.spirv_type = "%float 2D 0 0 0 2 R32f",
.spirv_fn = "OpImageWrite %im %vi12 %float_1",
.wgsl_type = "texture_storage_2d<r32float, read_write>",
.wgsl_fn = R"(
%4:vec4<f32> = construct 1.0f
%5:void = textureStore %3, vec2<i32>(1i, 2i), %4)"},
ImgData{
.name = "Source 2 component, dest 1 component",
.spirv_type = "%float 2D 0 0 0 2 R32f",
.spirv_fn = "OpImageWrite %im %vi12 %vf12",
.wgsl_type = "texture_storage_2d<r32float, read_write>",
.wgsl_fn = R"(
%4:vec4<f32> = construct vec2<f32>(1.0f, 2.0f), vec2<f32>(0.0f)
%5:void = textureStore %3, vec2<i32>(1i, 2i), %4)",
},
ImgData{
.name = "Source 3 component, dest 1 component",
.spirv_type = "%float 2D 0 0 0 2 R32f",
.spirv_fn = "OpImageWrite %im %vi12 %vf123",
.wgsl_type = "texture_storage_2d<r32float, read_write>",
.wgsl_fn = R"(
%4:vec4<f32> = construct vec3<f32>(1.0f, 2.0f, 3.0f), 0.0f
%5:void = textureStore %3, vec2<i32>(1i, 2i), %4)",
},
ImgData{
.name = "Source 4 component, dest 1 component",
.spirv_type = "%float 2D 0 0 0 2 R32f",
.spirv_fn = "OpImageWrite %im %vi12 %vf1234",
.wgsl_type = "texture_storage_2d<r32float, read_write>",
.wgsl_fn = R"(
%4:void = textureStore %3, vec2<i32>(1i, 2i), vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f))",
},
ImgData{
.name = "Source 2 component, dest 2 component",
.spirv_type = "%float 2D 0 0 0 2 Rg32f",
.spirv_fn = "OpImageWrite %im %vi12 %vf12",
.wgsl_type = "texture_storage_2d<rg32float, read_write>",
.wgsl_fn = R"(
%4:vec4<f32> = construct vec2<f32>(1.0f, 2.0f), vec2<f32>(0.0f)
%5:void = textureStore %3, vec2<i32>(1i, 2i), %4)",
},
ImgData{
.name = "Source 3 component, dest 2 component",
.spirv_type = "%float 2D 0 0 0 2 Rg32f",
.spirv_fn = "OpImageWrite %im %vi12 %vf123",
.wgsl_type = "texture_storage_2d<rg32float, read_write>",
.wgsl_fn = R"(
%4:vec4<f32> = construct vec3<f32>(1.0f, 2.0f, 3.0f), 0.0f
%5:void = textureStore %3, vec2<i32>(1i, 2i), %4)",
},
ImgData{
.name = "Source 4 component, dest 2 component",
.spirv_type = "%float 2D 0 0 0 2 Rg32f",
.spirv_fn = "OpImageWrite %im %vi12 %vf1234",
.wgsl_type = "texture_storage_2d<rg32float, read_write>",
.wgsl_fn = R"(
%4:void = textureStore %3, vec2<i32>(1i, 2i), vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f))",
},
ImgData{
.name = "Source 4 component, dest 4 component",
.spirv_type = "%float 2D 0 0 0 2 Rgba32f",
.spirv_fn = "OpImageWrite %im %vi12 %vf1234",
.wgsl_type = "texture_storage_2d<rgba32float, read_write>",
.wgsl_fn = R"(
%4:void = textureStore %3, vec2<i32>(1i, 2i), vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f))",
}));
// As above, but unsigned integer.
INSTANTIATE_TEST_SUITE_P(SpirvReaderTest_ImageWrite_ConvertTexelOperand_Arity_Uint,
ImageAccessTest,
::testing::Values(
ImgData{
.name = "Source 1 component",
.spirv_type = "%uint 2D 0 0 0 2 R32ui",
.spirv_fn = "OpImageWrite %im %vi12 %uint_1",
.wgsl_type = "texture_storage_2d<r32uint, read_write>",
.wgsl_fn = R"(
%4:vec4<u32> = construct 1u
%5:void = textureStore %3, vec2<i32>(1i, 2i), %4)",
},
ImgData{
.name = "Source 2 component, dest 1 component",
.spirv_type = "%uint 2D 0 0 0 2 R32ui",
.spirv_fn = "OpImageWrite %im %vi12 %vu12",
.wgsl_type = "texture_storage_2d<r32uint, read_write>",
.wgsl_fn = R"(
%4:vec4<u32> = construct vec2<u32>(1u, 2u), vec2<u32>(0u)
%5:void = textureStore %3, vec2<i32>(1i, 2i), %4)",
},
ImgData{
.name = "Source 3 component, dest 1 component",
.spirv_type = "%uint 2D 0 0 0 2 R32ui",
.spirv_fn = "OpImageWrite %im %vi12 %vu123",
.wgsl_type = "texture_storage_2d<r32uint, read_write>",
.wgsl_fn = R"(
%4:vec4<u32> = construct vec3<u32>(1u, 2u, 3u), 0u
%5:void = textureStore %3, vec2<i32>(1i, 2i), %4)",
},
ImgData{
.name = "Source 4 component, dest 1 component",
.spirv_type = "%uint 2D 0 0 0 2 R32ui",
.spirv_fn = "OpImageWrite %im %vi12 %vu1234",
.wgsl_type = "texture_storage_2d<r32uint, read_write>",
.wgsl_fn = R"(
%4:void = textureStore %3, vec2<i32>(1i, 2i), vec4<u32>(1u, 2u, 3u, 4u))",
},
ImgData{
.name = "Source 2 component, dest 2 component",
.spirv_type = "%uint 2D 0 0 0 2 Rg32ui",
.spirv_fn = "OpImageWrite %im %vi12 %vu12",
.wgsl_type = "texture_storage_2d<rg32uint, read_write>",
.wgsl_fn = R"(
%4:vec4<u32> = construct vec2<u32>(1u, 2u), vec2<u32>(0u)
%5:void = textureStore %3, vec2<i32>(1i, 2i), %4)",
},
ImgData{
.name = "Source 3 component, dest 2 component",
.spirv_type = "%uint 2D 0 0 0 2 Rg32ui",
.spirv_fn = "OpImageWrite %im %vi12 %vu123",
.wgsl_type = "texture_storage_2d<rg32uint, read_write>",
.wgsl_fn = R"(
%4:vec4<u32> = construct vec3<u32>(1u, 2u, 3u), 0u
%5:void = textureStore %3, vec2<i32>(1i, 2i), %4)",
},
ImgData{
.name = "Source 4 component, dest 2 component",
.spirv_type = "%uint 2D 0 0 0 2 Rg32ui",
.spirv_fn = "OpImageWrite %im %vi12 %vu1234",
.wgsl_type = "texture_storage_2d<rg32uint, read_write>",
.wgsl_fn = R"(
%4:void = textureStore %3, vec2<i32>(1i, 2i), vec4<u32>(1u, 2u, 3u, 4u))",
},
ImgData{
.name = "Source 4 component, dest 4 component",
.spirv_type = "%uint 2D 0 0 0 2 Rgba32ui",
.spirv_fn = "OpImageWrite %im %vi12 %vu1234",
.wgsl_type = "texture_storage_2d<rgba32uint, read_write>",
.wgsl_fn = R"(
%4:void = textureStore %3, vec2<i32>(1i, 2i), vec4<u32>(1u, 2u, 3u, 4u))",
}));
// As above, but signed integer.
INSTANTIATE_TEST_SUITE_P(SpirvReaderTest_ImageWrite_ConvertTexelOperand_Arity_Sint,
ImageAccessTest,
::testing::Values(
ImgData{
.name = "Source 1 component",
.spirv_type = "%int 2D 0 0 0 2 R32i",
.spirv_fn = "OpImageWrite %im %vi12 %int_1",
.wgsl_type = "texture_storage_2d<r32sint, read_write>",
.wgsl_fn = R"(
%4:vec4<i32> = construct 1i
%5:void = textureStore %3, vec2<i32>(1i, 2i), %4)",
},
ImgData{
.name = "Source 2 component, dest 1 component",
.spirv_type = "%int 2D 0 0 0 2 R32i",
.spirv_fn = "OpImageWrite %im %vi12 %vi12",
.wgsl_type = "texture_storage_2d<r32sint, read_write>",
.wgsl_fn = R"(
%4:vec4<i32> = construct vec2<i32>(1i, 2i), vec2<i32>(0i)
%5:void = textureStore %3, vec2<i32>(1i, 2i), %4)",
},
ImgData{
.name = "Source 3 component, dest 1 component",
.spirv_type = "%int 2D 0 0 0 2 R32i",
.spirv_fn = "OpImageWrite %im %vi12 %vi123",
.wgsl_type = "texture_storage_2d<r32sint, read_write>",
.wgsl_fn = R"(
%4:vec4<i32> = construct vec3<i32>(1i, 2i, 3i), 0i
%5:void = textureStore %3, vec2<i32>(1i, 2i), %4)",
},
ImgData{
.name = "Source 4 component, dest 1 component",
.spirv_type = "%int 2D 0 0 0 2 R32i",
.spirv_fn = "OpImageWrite %im %vi12 %vi1234",
.wgsl_type = "texture_storage_2d<r32sint, read_write>",
.wgsl_fn = R"(
%4:void = textureStore %3, vec2<i32>(1i, 2i), vec4<i32>(1i, 2i, 3i, 4i))",
},
ImgData{
.name = "Source 2 component, dest 2 component",
.spirv_type = "%int 2D 0 0 0 2 Rg32i",
.spirv_fn = "OpImageWrite %im %vi12 %vi12",
.wgsl_type = "texture_storage_2d<rg32sint, read_write>",
.wgsl_fn = R"(
%4:vec4<i32> = construct vec2<i32>(1i, 2i), vec2<i32>(0i)
%5:void = textureStore %3, vec2<i32>(1i, 2i), %4)",
},
ImgData{
.name = "Source 3 component, dest 2 component",
.spirv_type = "%int 2D 0 0 0 2 Rg32i",
.spirv_fn = "OpImageWrite %im %vi12 %vi123",
.wgsl_type = "texture_storage_2d<rg32sint, read_write>",
.wgsl_fn = R"(
%4:vec4<i32> = construct vec3<i32>(1i, 2i, 3i), 0i
%5:void = textureStore %3, vec2<i32>(1i, 2i), %4)",
},
ImgData{
.name = "Source 4 component, dest 2 component",
.spirv_type = "%int 2D 0 0 0 2 Rg32i",
.spirv_fn = "OpImageWrite %im %vi12 %vi1234",
.wgsl_type = "texture_storage_2d<rg32sint, read_write>",
.wgsl_fn = R"(
%4:void = textureStore %3, vec2<i32>(1i, 2i), vec4<i32>(1i, 2i, 3i, 4i))",
},
ImgData{
.name = "Source 4 component, dest 4 component",
.spirv_type = "%int 2D 0 0 0 2 Rgba32i",
.spirv_fn = "OpImageWrite %im %vi12 %vi1234",
.wgsl_type = "texture_storage_2d<rgba32sint, read_write>",
.wgsl_fn = R"(
%4:void = textureStore %3, vec2<i32>(1i, 2i), vec4<i32>(1i, 2i, 3i, 4i))",
}));
// The texel operand signedness must match the channel type signedness.
// SPIR-V validation checks that.
// This suite is for the cases where they are integral and the same
// signedness.
INSTANTIATE_TEST_SUITE_P(SpirvReaderTest_ImageWrite_ConvertTexelOperand_SameSignedness,
ImageAccessTest,
::testing::Values(
ImgData{
.name = "Sampled type is unsigned int, texel is unsigned int",
.spirv_type = "%uint 2D 0 0 0 2 Rgba32ui",
.spirv_fn = "OpImageWrite %im %vi12 %vu1234",
.wgsl_type = "texture_storage_2d<rgba32uint, read_write>",
.wgsl_fn = R"(
%4:void = textureStore %3, vec2<i32>(1i, 2i), vec4<u32>(1u, 2u, 3u, 4u))",
},
ImgData{
.name = "Sampled type is signed int, texel is signed int",
.spirv_type = "%int 2D 0 0 0 2 Rgba32i",
.spirv_fn = "OpImageWrite %im %vi12 %vi1234",
.wgsl_type = "texture_storage_2d<rgba32sint, read_write>",
.wgsl_fn = R"(
%4:void = textureStore %3, vec2<i32>(1i, 2i), vec4<i32>(1i, 2i, 3i, 4i))",
}));
// Show that zeros of the correct integer signedness are
// created when expanding an integer vector.
INSTANTIATE_TEST_SUITE_P(SpirvReaderTest_ImageWrite_ConvertTexelOperand_Signedness_AndWidening,
ImageAccessTest,
::testing::Values(
ImgData{
.name = "Source unsigned, dest unsigned",
.spirv_type = "%uint 2D 0 0 0 2 R32ui",
.spirv_fn = "OpImageWrite %im %vi12 %vu12",
.wgsl_type = "texture_storage_2d<r32uint, read_write>",
.wgsl_fn = R"(
%4:vec4<u32> = construct vec2<u32>(1u, 2u), vec2<u32>(0u)
%5:void = textureStore %3, vec2<i32>(1i, 2i), %4)",
},
ImgData{
.name = "Source signed, dest signed",
.spirv_type = "%int 2D 0 0 0 2 R32i",
.spirv_fn = "OpImageWrite %im %vi12 %vi12",
.wgsl_type = "texture_storage_2d<r32sint, read_write>",
.wgsl_fn = R"(
%4:vec4<i32> = construct vec2<i32>(1i, 2i), vec2<i32>(0i)
%5:void = textureStore %3, vec2<i32>(1i, 2i), %4)",
}));
INSTANTIATE_TEST_SUITE_P(
SpirvReaderTest_ImageFetch,
ImageAccessTest,
::testing::Values(
ImgData{
.name = "OpImageFetch with no extra params, on sampled texture",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageFetch %v4float %im %vi12",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%4:vec4<f32> = textureLoad %3, vec2<i32>(1i, 2i), 0i)",
},
ImgData{
.name = "OpImageFetch with explicit level, on sampled texture",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageFetch %v4float %im %vi12 Lod %int_3",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%4:vec4<f32> = textureLoad %3, vec2<i32>(1i, 2i), 3i)",
},
ImgData{
.name = "OpImageFetch with no extra params, on depth texture",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageFetch %v4float %im %vi12",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = R"(
%4:f32 = textureLoad %3, vec2<i32>(1i, 2i), 0i
%5:vec4<f32> = construct %4, 0.0f, 0.0f, 0.0f)",
},
ImgData{
.name = "OpImageFetch with extra params, on depth texture",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageFetch %v4float %im %vi12 Lod %int_3",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = R"(
%4:f32 = textureLoad %3, vec2<i32>(1i, 2i), 3i
%5:vec4<f32> = construct %4, 0.0f, 0.0f, 0.0f)",
},
// In SPIR-V OpImageFetch always yields a vector of 4 elements, even for depth images. But
// in WGSL, textureLoad on a depth image yields f32.
// https://crbug.com/tint/439
ImgData{
.name = "ImageFetch on depth image.",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageFetch %v4float %im %vi12",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = R"(
%4:f32 = textureLoad %3, vec2<i32>(1i, 2i), 0i
%5:vec4<f32> = construct %4, 0.0f, 0.0f, 0.0f)",
},
// In SPIR-V OpImageFetch always yields a vector of 4 elements, even for depth images. But
// in WGSL, textureLoad on a depth image yields f32.
// https://crbug.com/tint/439
ImgData{
.name = "ImageFetch on multisampled depth image.",
.spirv_type = "%float 2D 1 0 1 1 Unknown",
.spirv_fn = "%99 = OpImageFetch %v4float %im %vi12 Sample %int_1",
.wgsl_type = "texture_depth_multisampled_2d",
.wgsl_fn = R"(
%4:f32 = textureLoad %3, vec2<i32>(1i, 2i), 1i
%5:vec4<f32> = construct %4, 0.0f, 0.0f, 0.0f)",
},
// SPIR-V requires a Sample image operand when operating on a multisampled image.
ImgData{
.name = "ImageFetch non-arrayed",
.spirv_type = "%float 2D 0 0 1 1 Unknown",
.spirv_fn = "%99 = OpImageFetch %v4float %im %vi12 Sample %int_1",
.wgsl_type = "texture_multisampled_2d<f32>",
.wgsl_fn = R"(
%4:vec4<f32> = textureLoad %3, vec2<i32>(1i, 2i), 1i)",
},
ImgData{
.name = "2d",
.spirv_type = "%float 2D 0 0 1 1 Unknown",
.spirv_fn = "%99 = OpImageFetch %v4float %im %vi12 Sample %uint_1",
.wgsl_type = "texture_multisampled_2d<f32>",
.wgsl_fn = R"(
%4:i32 = convert 1u
%5:vec4<f32> = textureLoad %3, vec2<i32>(1i, 2i), %4)",
}));
using SampledImageAccessTest = SpirvReaderTestWithParam<ImgData>;
TEST_P(SampledImageAccessTest, Variable) {
auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpCapability ImageQuery
OpCapability Sampled1D
OpMemoryModel Logical Simple
OpEntryPoint Fragment %main "main"
OpExecutionMode %main OriginUpperLeft
OpName %20 "wg"
OpDecorate %20 DescriptorSet 2
OpDecorate %20 Binding 1
%float = OpTypeFloat 32
%int = OpTypeInt 32 1
%uint = OpTypeInt 32 0
%v2int = OpTypeVector %int 2
%v3int = OpTypeVector %int 3
%v4int = OpTypeVector %int 4
%v2uint = OpTypeVector %uint 2
%v4uint = OpTypeVector %uint 4
%v2float = OpTypeVector %float 2
%v4float = OpTypeVector %float 4
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%im_ty = OpTypeImage )" +
params.spirv_type + R"(
%ptr_im_ty = OpTypePointer UniformConstant %im_ty
%20 = OpVariable %ptr_im_ty UniformConstant
%int_1 = OpConstant %int 1
%int_2 = OpConstant %int 2
%int_3 = OpConstant %int 3
%int_4 = OpConstant %int 4
%uint_1 = OpConstant %uint 1
%uint_3 = OpConstant %uint 3
%uint_4 = OpConstant %uint 4
%vi12 = OpConstantComposite %v2int %int_1 %int_2
%float_1 = OpConstant %float 1
%float_2 = OpConstant %float 2
%vf12 = OpConstantComposite %v2float %float_1 %float_2
%main = OpFunction %void None %voidfn
%entry = OpLabel
%im = OpLoad %im_ty %20
)" + params.spirv_fn +
R"(
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%wg:ptr<handle, )" +
params.wgsl_type +
R"(, read> = var undef @binding_point(2, 1)
}
%main = @fragment func():void {
$B2: {
%3:)" + params.wgsl_type +
R"( = load %wg)" + params.wgsl_fn +
R"(
ret
}
}
)");
}
INSTANTIATE_TEST_SUITE_P(SpirvReaderTest_ConvertResultSignedness,
SampledImageAccessTest,
::testing::Values(
ImgData{
.name = "2d no conversion float",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageFetch %v4float %im %vi12",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%4:vec4<f32> = textureLoad %3, vec2<i32>(1i, 2i), 0i)",
},
ImgData{
.name = "2d no conversion uint",
.spirv_type = "%uint 2D 0 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageFetch %v4uint %im %vi12",
.wgsl_type = "texture_2d<u32>",
.wgsl_fn = R"(
%4:vec4<u32> = textureLoad %3, vec2<i32>(1i, 2i), 0i)",
},
ImgData{
.name = "2d no conversion int",
.spirv_type = "%int 2D 0 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageFetch %v4int %im %vi12",
.wgsl_type = "texture_2d<i32>",
.wgsl_fn = R"(
%4:vec4<i32> = textureLoad %3, vec2<i32>(1i, 2i), 0i)",
}));
// From VUID-StandaloneSpirv-OpImageQuerySizeLod-04659:
// ImageQuerySizeLod requires Sampled=1
INSTANTIATE_TEST_SUITE_P(SpirvReaderTest_ImageQuerySizeLod_NonArrayed_SignedResult_SignedLevel,
SampledImageAccessTest,
::testing::Values(
ImgData{
.name = "1D",
.spirv_type = "%float 1D 0 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageQuerySizeLod %int %im %int_1\n",
.wgsl_type = "texture_1d<f32>",
.wgsl_fn = R"(
%4:u32 = textureDimensions %3, 1i
%5:i32 = convert %4)",
},
ImgData{
.name = "2D",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageQuerySizeLod %v2int %im %int_1\n",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%4:vec2<u32> = textureDimensions %3, 1i
%5:vec2<i32> = convert %4)",
},
ImgData{
.name = "3D",
.spirv_type = "%float 3D 0 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageQuerySizeLod %v3int %im %int_1\n",
.wgsl_type = "texture_3d<f32>",
.wgsl_fn = R"(
%4:vec3<u32> = textureDimensions %3, 1i
%5:vec3<i32> = convert %4)",
},
ImgData{
.name = "Cube",
.spirv_type = "%float Cube 0 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageQuerySizeLod %v2int %im %int_1\n",
.wgsl_type = "texture_cube<f32>",
.wgsl_fn = R"(
%4:vec2<u32> = textureDimensions %3, 1i
%5:vec2<i32> = convert %4)",
},
ImgData{
.name = "Depth 2D",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageQuerySizeLod %v2int %im %int_1\n",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = R"(
%4:vec2<u32> = textureDimensions %3, 1i
%5:vec2<i32> = convert %4)",
},
ImgData{
.name = "Depth Cube",
.spirv_type = "%float Cube 1 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageQuerySizeLod %v2int %im %int_1\n",
.wgsl_type = "texture_depth_cube",
.wgsl_fn = R"(
%4:vec2<u32> = textureDimensions %3, 1i
%5:vec2<i32> = convert %4)",
}));
// ImageQuerySize requires storage image or multisampled
// For storage image, use another instruction to indicate whether it
// is readonly or writeonly.
INSTANTIATE_TEST_SUITE_P(SpirvReaderTest_ImageQuerySizeLod_Arrayed_SignedResult_SignedLevel,
SampledImageAccessTest,
::testing::Values(
ImgData{
.name = "2D array",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn = "%99 = OpImageQuerySizeLod %v3int %im %int_1\n",
.wgsl_type = "texture_2d_array<f32>",
.wgsl_fn =
R"(
%4:vec2<u32> = textureDimensions %3, 1i
%5:u32 = textureNumLayers %3
%6:vec3<u32> = construct %4, %5
%7:vec3<i32> = convert %6)",
},
ImgData{
.name = "Cube array",
.spirv_type = "%float Cube 0 1 0 1 Unknown",
.spirv_fn = "%99 = OpImageQuerySizeLod %v3int %im %int_1\n",
.wgsl_type = "texture_cube_array<f32>",
.wgsl_fn =
R"(
%4:vec2<u32> = textureDimensions %3, 1i
%5:u32 = textureNumLayers %3
%6:vec3<u32> = construct %4, %5
%7:vec3<i32> = convert %6)",
},
ImgData{
.name = "Depth 2D array",
.spirv_type = "%float 2D 1 1 0 1 Unknown",
.spirv_fn = "%99 = OpImageQuerySizeLod %v3int %im %int_1\n",
.wgsl_type = "texture_depth_2d_array",
.wgsl_fn =
R"(
%4:vec2<u32> = textureDimensions %3, 1i
%5:u32 = textureNumLayers %3
%6:vec3<u32> = construct %4, %5
%7:vec3<i32> = convert %6)",
},
ImgData{
.name = "Depth Cube Array",
.spirv_type = "%float Cube 1 1 0 1 Unknown",
.spirv_fn = "%99 = OpImageQuerySizeLod %v3int %im %int_1\n",
.wgsl_type = "texture_depth_cube_array",
.wgsl_fn =
R"(
%4:vec2<u32> = textureDimensions %3, 1i
%5:u32 = textureNumLayers %3
%6:vec3<u32> = construct %4, %5
%7:vec3<i32> = convert %6)",
}));
// textureDimensions accepts both signed and unsigned the level-of-detail values.
INSTANTIATE_TEST_SUITE_P(SpirvReaderTest_ImageQuerySizeLod_NonArrayed_SignedResult_UnsignedLevel,
SampledImageAccessTest,
::testing::Values(ImgData{
.name = "1D",
.spirv_type = "%float 1D 0 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageQuerySizeLod %int %im %uint_1\n",
.wgsl_type = "texture_1d<f32>",
.wgsl_fn = R"(
%4:u32 = textureDimensions %3, 1u
%5:i32 = convert %4)",
}));
// When SPIR-V wants the result type to be unsigned, we have to insert a value constructor
// for WGSL to do the type coercion.
INSTANTIATE_TEST_SUITE_P(SpirvReaderTest_ImageQuerySizeLod_NonArrayed_UnsignedResult_SignedLevel,
SampledImageAccessTest,
::testing::Values(ImgData{
.name = "1D",
.spirv_type = "%float 1D 0 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageQuerySizeLod %uint %im %int_1",
.wgsl_type = "texture_1d<f32>",
.wgsl_fn = R"(
%4:u32 = textureDimensions %3, 1i)",
}));
INSTANTIATE_TEST_SUITE_P(SpirvReaderTest_ImageQueryLevels_SignedResult,
SampledImageAccessTest,
::testing::Values(
ImgData{
.name = "2D",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageQueryLevels %int %im",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%4:u32 = textureNumLevels %3
%5:i32 = convert %4)",
},
ImgData{
.name = "2D array",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn = "%99 = OpImageQueryLevels %int %im",
.wgsl_type = "texture_2d_array<f32>",
.wgsl_fn = R"(
%4:u32 = textureNumLevels %3
%5:i32 = convert %4)",
},
ImgData{
.name = "3D",
.spirv_type = "%float 3D 0 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageQueryLevels %int %im",
.wgsl_type = "texture_3d<f32>",
.wgsl_fn = R"(
%4:u32 = textureNumLevels %3
%5:i32 = convert %4)",
},
ImgData{
.name = "Cube",
.spirv_type = "%float Cube 0 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageQueryLevels %int %im",
.wgsl_type = "texture_cube<f32>",
.wgsl_fn = R"(
%4:u32 = textureNumLevels %3
%5:i32 = convert %4)",
},
ImgData{
.name = "Cube array",
.spirv_type = "%float Cube 0 1 0 1 Unknown",
.spirv_fn = "%99 = OpImageQueryLevels %int %im",
.wgsl_type = "texture_cube_array<f32>",
.wgsl_fn = R"(
%4:u32 = textureNumLevels %3
%5:i32 = convert %4)",
},
ImgData{
.name = "depth 2d",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageQueryLevels %int %im",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = R"(
%4:u32 = textureNumLevels %3
%5:i32 = convert %4)",
},
ImgData{
.name = "depth 2d array",
.spirv_type = "%float 2D 1 1 0 1 Unknown",
.spirv_fn = "%99 = OpImageQueryLevels %int %im",
.wgsl_type = "texture_depth_2d_array",
.wgsl_fn = R"(
%4:u32 = textureNumLevels %3
%5:i32 = convert %4)",
},
ImgData{
.name = "depth cube",
.spirv_type = "%float Cube 1 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageQueryLevels %int %im",
.wgsl_type = "texture_depth_cube",
.wgsl_fn = R"(
%4:u32 = textureNumLevels %3
%5:i32 = convert %4)",
},
ImgData{
.name = "depth cube array",
.spirv_type = "%float Cube 1 1 0 1 Unknown",
.spirv_fn = "%99 = OpImageQueryLevels %int %im",
.wgsl_type = "texture_depth_cube_array",
.wgsl_fn = R"(
%4:u32 = textureNumLevels %3
%5:i32 = convert %4)",
}));
// Spot check that a value conversion is inserted when SPIR-V asks for an unsigned int result.
INSTANTIATE_TEST_SUITE_P(SpirvReaderTest_ImageQueryLevels_UnsignedResult,
SampledImageAccessTest,
::testing::Values(ImgData{
.name = "2D",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageQueryLevels %uint %im\n",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%4:u32 = textureNumLevels %3)",
}));
using NoSamplerImageAccessTest = SpirvReaderTestWithParam<ImgData>;
TEST_P(NoSamplerImageAccessTest, Variable) {
auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpCapability ImageQuery
OpCapability Image1D
OpCapability Sampled1D
OpMemoryModel Logical Simple
OpEntryPoint Fragment %main "main"
OpExecutionMode %main OriginUpperLeft
OpName %wg "wg"
OpDecorate %wg DescriptorSet 2
OpDecorate %wg Binding 1
%float = OpTypeFloat 32
%int = OpTypeInt 32 1
%uint = OpTypeInt 32 0
%v2int = OpTypeVector %int 2
%v3int = OpTypeVector %int 3
%v4int = OpTypeVector %int 4
%v2uint = OpTypeVector %uint 2
%v3uint = OpTypeVector %uint 3
%v4uint = OpTypeVector %uint 4
%v4float = OpTypeVector %float 4
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%im_ty = OpTypeImage )" +
params.spirv_type + R"(
%ptr_im_ty = OpTypePointer UniformConstant %im_ty
%int_1 = OpConstant %int 1
%int_2 = OpConstant %int 2
%int_3 = OpConstant %int 3
%vi12 = OpConstantComposite %v2int %int_1 %int_2
%vi123 = OpConstantComposite %v3int %int_1 %int_2 %int_3
%wg = OpVariable %ptr_im_ty UniformConstant
%main = OpFunction %void None %voidfn
%entry = OpLabel
%im = OpLoad %im_ty %wg
)" + params.spirv_fn +
R"(
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%wg:ptr<handle, )" +
params.wgsl_type +
R"(, read> = var undef @binding_point(2, 1)
}
%main = @fragment func():void {
$B2: {
%3:)" + params.wgsl_type +
R"( = load %wg)" + params.wgsl_fn +
R"(
ret
}
}
)");
}
INSTANTIATE_TEST_SUITE_P(SpirvReaderTest_CheckResultSignedness,
NoSamplerImageAccessTest,
::testing::Values(
ImgData{
.name = "no conversion, float -> v4float",
.spirv_type = "%float 2D 0 0 0 2 Rgba32f",
.spirv_fn = "%99 = OpImageRead %v4float %im %vi12",
.wgsl_type = "texture_storage_2d<rgba32float, read_write>",
.wgsl_fn = R"(
%4:vec4<f32> = textureLoad %3, vec2<i32>(1i, 2i))",
},
ImgData{
.name = "no conversion, uint -> v4uint",
.spirv_type = "%uint 2D 0 0 0 2 Rgba32ui",
.spirv_fn = "%99 = OpImageRead %v4uint %im %vi12",
.wgsl_type = "texture_storage_2d<rgba32uint, read_write>",
.wgsl_fn = R"(
%4:vec4<u32> = textureLoad %3, vec2<i32>(1i, 2i))",
},
ImgData{
.name = "no conversion, int -> v4int",
.spirv_type = "%int 2D 0 0 0 2 Rgba32i",
.spirv_fn = "%99 = OpImageRead %v4int %im %vi12",
.wgsl_type = "texture_storage_2d<rgba32sint, read_write>",
.wgsl_fn = R"(
%4:vec4<i32> = textureLoad %3, vec2<i32>(1i, 2i))",
}));
// ImageQuerySize requires storage image or multisampled
// For storage image, use another instruction to indicate whether it is readonly or writeonly.
INSTANTIATE_TEST_SUITE_P(SpirvReaderTest_ImageQuerySize_NonArrayed_SignedResult,
NoSamplerImageAccessTest,
::testing::Values(
ImgData{
.name = "1D storage image",
.spirv_type = "%float 1D 0 0 0 2 Rgba32f",
.spirv_fn = "%99 = OpImageQuerySize %int %im",
.wgsl_type = "texture_storage_1d<rgba32float, read_write>",
.wgsl_fn = R"(
%4:u32 = textureDimensions %3
%5:i32 = convert %4)",
},
ImgData{
.name = "2D storage image",
.spirv_type = "%float 2D 0 0 0 2 Rgba32f",
.spirv_fn = "%99 = OpImageQuerySize %v2int %im",
.wgsl_type = "texture_storage_2d<rgba32float, read_write>",
.wgsl_fn = R"(
%4:vec2<u32> = textureDimensions %3
%5:vec2<i32> = convert %4)",
},
ImgData{
.name = "3D storage image",
.spirv_type = "%float 3D 0 0 0 2 Rgba32f",
.spirv_fn = "%99 = OpImageQuerySize %v3int %im",
.wgsl_type = "texture_storage_3d<rgba32float, read_write>",
.wgsl_fn = R"(
%4:vec3<u32> = textureDimensions %3
%5:vec3<i32> = convert %4)",
},
ImgData{
.name = "Multisampled",
.spirv_type = "%float 2D 0 0 1 1 Unknown",
.spirv_fn = "%99 = OpImageQuerySize %v2int %im",
.wgsl_type = "texture_multisampled_2d<f32>",
.wgsl_fn = R"(
%4:vec2<u32> = textureDimensions %3
%5:vec2<i32> = convert %4)",
}));
// ImageQuerySize requires storage image or multisampled
// For storage image, use another instruction to indicate whether it is readonly or writeonly.
INSTANTIATE_TEST_SUITE_P(SpirvReaderTest_ImageQuerySize_Arrayed_SignedResult,
NoSamplerImageAccessTest,
::testing::Values(ImgData{
.name = "2D array storage image",
.spirv_type = "%float 2D 0 1 0 2 Rgba32f",
.spirv_fn = "%99 = OpImageQuerySize %v3int %im",
.wgsl_type = "texture_storage_2d_array<rgba32float, read_write>",
.wgsl_fn = R"(
%4:vec2<u32> = textureDimensions %3
%5:u32 = textureNumLayers %3
%6:vec3<u32> = construct %4, %5
%7:vec3<i32> = convert %6)",
}));
// ImageQuerySize requires storage image or multisampled
// For storage image, use another instruction to indicate whether it is readonly or writeonly.
INSTANTIATE_TEST_SUITE_P(SpirvReaderTest_ImageQuerySize_NonArrayed_UnsignedResult,
NoSamplerImageAccessTest,
::testing::Values(
ImgData{
.name = "1D storage image",
.spirv_type = "%float 1D 0 0 0 2 Rgba32f",
.spirv_fn = "%99 = OpImageQuerySize %uint %im",
.wgsl_type = "texture_storage_1d<rgba32float, read_write>",
.wgsl_fn = R"(
%4:u32 = textureDimensions %3)",
},
ImgData{
.name = "2D storage image",
.spirv_type = "%float 2D 0 0 0 2 Rgba32f",
.spirv_fn = "%99 = OpImageQuerySize %v2uint %im",
.wgsl_type = "texture_storage_2d<rgba32float, read_write>",
.wgsl_fn = R"(
%4:vec2<u32> = textureDimensions %3)",
},
ImgData{
.name = "3D storage image",
.spirv_type = "%float 3D 0 0 0 2 Rgba32f",
.spirv_fn = "%99 = OpImageQuerySize %v3uint %im",
.wgsl_type = "texture_storage_3d<rgba32float, read_write>",
.wgsl_fn = R"(
%4:vec3<u32> = textureDimensions %3)",
},
ImgData{
.name = "Multisampled",
.spirv_type = "%float 2D 0 0 1 1 Unknown",
.spirv_fn = "%99 = OpImageQuerySize %v2uint %im",
.wgsl_type = "texture_multisampled_2d<f32>",
.wgsl_fn = R"(
%4:vec2<u32> = textureDimensions %3)",
}));
INSTANTIATE_TEST_SUITE_P(SpirvReaderTest_ImageQuerySize_Arrayed_UnsignedResult,
NoSamplerImageAccessTest,
::testing::Values(ImgData{
.name = "2D array storage image",
.spirv_type = "%float 2D 0 1 0 2 Rgba32f",
.spirv_fn = "%99 = OpImageQuerySize %v3uint %im",
.wgsl_type = "texture_storage_2d_array<rgba32float, read_write>",
.wgsl_fn = R"(
%4:vec2<u32> = textureDimensions %3
%5:u32 = textureNumLayers %3
%6:vec3<u32> = construct %4, %5)",
}));
INSTANTIATE_TEST_SUITE_P(SpirvReaderTest_ImageQuerySamples_SignedResult,
NoSamplerImageAccessTest,
::testing::Values(ImgData{
.name = "Multsample 2D",
.spirv_type = "%float 2D 0 0 1 1 Unknown",
.spirv_fn = "%99 = OpImageQuerySamples %int %im",
.wgsl_type = "texture_multisampled_2d<f32>",
.wgsl_fn = R"(
%4:u32 = textureNumSamples %3
%5:i32 = convert %4)",
}));
// Translation must inject a type coercion from unsigned to signed.
INSTANTIATE_TEST_SUITE_P(SpirvReaderTest_ImageQuerySamples_UnsignedResult,
NoSamplerImageAccessTest,
::testing::Values(ImgData{
.name = "Multisample 2D",
.spirv_type = "%float 2D 0 0 1 1 Unknown",
.spirv_fn = "%99 = OpImageQuerySamples %uint %im",
.wgsl_type = "texture_multisampled_2d<f32>",
.wgsl_fn = R"(
%4:u32 = textureNumSamples %3)",
}));
using SampledImageCoordsTest = SpirvReaderTestWithParam<ImgData>;
TEST_P(SampledImageCoordsTest, MakeCoordinateOperandsForImageAccess) {
auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpCapability Sampled1D
OpMemoryModel Logical Simple
OpEntryPoint Fragment %100 "main"
OpExecutionMode %100 OriginUpperLeft
OpName %20 "wg"
OpDecorate %10 DescriptorSet 0
OpDecorate %10 Binding 0
OpDecorate %20 DescriptorSet 2
OpDecorate %20 Binding 0
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%float = OpTypeFloat 32
%int = OpTypeInt 32 1
%uint = OpTypeInt 32 0
%v2int = OpTypeVector %int 2
%v3int = OpTypeVector %int 3
%v2uint = OpTypeVector %uint 2
%v3uint = OpTypeVector %uint 3
%v2float = OpTypeVector %float 2
%v3float = OpTypeVector %float 3
%v4float = OpTypeVector %float 4
%float_null = OpConstantNull %float
%int_1 = OpConstant %int 1
%int_2 = OpConstant %int 2
%int_3 = OpConstant %int 3
%uint_1 = OpConstant %uint 1
%uint_2 = OpConstant %uint 2
%uint_3 = OpConstant %uint 3
%float_0 = OpConstant %float 0
%float_1 = OpConstant %float 1
%float_2 = OpConstant %float 2
%float_3 = OpConstant %float 3
%float_4 = OpConstant %float 4
%vi12 = OpConstantComposite %v2int %int_1 %int_2
%vi123 = OpConstantComposite %v3int %int_1 %int_2 %int_3
%vu12 = OpConstantComposite %v2uint %uint_1 %uint_2
%vu123 = OpConstantComposite %v3uint %uint_1 %uint_2 %uint_3
%vf12 = OpConstantComposite %v2float %float_1 %float_2
%vf123 = OpConstantComposite %v3float %float_1 %float_2 %float_3
%vf1234 = OpConstantComposite %v4float %float_1 %float_2 %float_3 %float_4
%sampler = OpTypeSampler
%ptr_sampler = OpTypePointer UniformConstant %sampler
%im_ty = OpTypeImage )" +
params.spirv_type + R"(
%ptr_im_ty = OpTypePointer UniformConstant %im_ty
%si_ty = OpTypeSampledImage %im_ty
%10 = OpVariable %ptr_sampler UniformConstant
%20 = OpVariable %ptr_im_ty UniformConstant
%100 = OpFunction %void None %voidfn
%entry = OpLabel
%sam = OpLoad %sampler %10
%im = OpLoad %im_ty %20
%sampled_image = OpSampledImage %si_ty %im %sam
)" + params.spirv_fn +
R"(
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%1:ptr<handle, sampler, read> = var undef @binding_point(0, 0)
%wg:ptr<handle, )" +
params.wgsl_type +
R"(, read> = var undef @binding_point(2, 0)
}
%main = @fragment func():void {
$B2: {
%4:sampler = load %1
%5:)" + params.wgsl_type +
R"( = load %wg)" + params.wgsl_fn +
R"(
ret
}
}
)");
}
INSTANTIATE_TEST_SUITE_P(
SpirvReaderTest_ImageSampleImplicitLod,
SampledImageCoordsTest,
::testing::Values(
ImgData{
.name = "1D",
.spirv_type = "%float 1D 0 0 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %float_1",
.wgsl_type = "texture_1d<f32>",
.wgsl_fn = R"(
%6:vec4<f32> = textureSample %5, %4, 1.0f)",
},
ImgData{
.name = "1D one extra arg",
.spirv_type = "%float 1D 0 0 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf12",
.wgsl_type = "texture_1d<f32>",
.wgsl_fn = R"(
%6:f32 = swizzle vec2<f32>(1.0f, 2.0f), x
%7:vec4<f32> = textureSample %5, %4, %6)",
},
ImgData{
.name = "1d two extra args",
.spirv_type = "%float 1D 0 0 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf123",
.wgsl_type = "texture_1d<f32>",
.wgsl_fn = R"(
%6:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), x
%7:vec4<f32> = textureSample %5, %4, %6)",
},
ImgData{
.name = "1D three extra args",
.spirv_type = "%float 1D 0 0 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf1234",
.wgsl_type = "texture_1d<f32>",
.wgsl_fn = R"(
%6:f32 = swizzle vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f), x
%7:vec4<f32> = textureSample %5, %4, %6)",
},
ImgData{
.name = "2D",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf12",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%6:vec4<f32> = textureSample %5, %4, vec2<f32>(1.0f, 2.0f))",
},
ImgData{
.name = "2D one excess arg",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf123",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:vec4<f32> = textureSample %5, %4, %6)",
},
ImgData{
.name = "2D two excess args",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf1234",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f), xy
%7:vec4<f32> = textureSample %5, %4, %6)",
},
ImgData{
.name = "2D array",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf123",
.wgsl_type = "texture_2d_array<f32>",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), xy
%7:f32 = swizzle vec3<f32>(1.0f, 2.0f, 3.0f), z
%8:i32 = convert %7
%9:vec4<f32> = textureSample %5, %4, %6, %8)",
},
ImgData{
.name = "2D array one excess arg",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf1234",
.wgsl_type = "texture_2d_array<f32>",
.wgsl_fn = R"(
%6:vec2<f32> = swizzle vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f), xy
%7:f32 = swizzle vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f), z
%8:i32 = convert %7
%9:vec4<f32> = textureSample %5, %4, %6, %8)",
},
ImgData{
.name = "3D",
.spirv_type = "%float 3D 0 0 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf123",
.wgsl_type = "texture_3d<f32>",
.wgsl_fn = R"(
%6:vec4<f32> = textureSample %5, %4, vec3<f32>(1.0f, 2.0f, 3.0f))",
},
ImgData{
.name = "3D one excess arg",
.spirv_type = "%float 3D 0 0 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf1234",
.wgsl_type = "texture_3d<f32>",
.wgsl_fn = R"(
%6:vec3<f32> = swizzle vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f), xyz
%7:vec4<f32> = textureSample %5, %4, %6)",
},
ImgData{
.name = "Cube",
.spirv_type = "%float Cube 0 0 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf123",
.wgsl_type = "texture_cube<f32>",
.wgsl_fn = R"(
%6:vec4<f32> = textureSample %5, %4, vec3<f32>(1.0f, 2.0f, 3.0f))",
},
ImgData{
.name = "Cube one excess arg",
.spirv_type = "%float Cube 0 0 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf1234",
.wgsl_type = "texture_cube<f32>",
.wgsl_fn = R"(
%6:vec3<f32> = swizzle vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f), xyz
%7:vec4<f32> = textureSample %5, %4, %6)",
},
ImgData{
.name = "Cube array",
.spirv_type = "%float Cube 0 1 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf1234",
.wgsl_type = "texture_cube_array<f32>",
.wgsl_fn = R"(
%6:vec3<f32> = swizzle vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f), xyz
%7:f32 = swizzle vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f), w
%8:i32 = convert %7
%9:vec4<f32> = textureSample %5, %4, %6, %8)",
},
ImgData{
.name = "2d no conversion float sampled lod",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageSampleImplicitLod %v4float %sampled_image %vf12",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%6:vec4<f32> = textureSample %5, %4, vec2<f32>(1.0f, 2.0f))",
}));
// In SPIR-V, sampling and dref sampling operations use floating point coordinates. Prove that we
// preserve floating point-ness. Test across all such instructions.
INSTANTIATE_TEST_SUITE_P(
DISABLED_SpirvReaderTest_PreserveFloatCoords_NonArrayed,
SampledImageCoordsTest,
::testing::Values(
// Scalar cases
ImgData{
.name = "1D",
.spirv_type = "%float 1D 0 0 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %float_1",
.wgsl_type = "",
.wgsl_fn = "f1",
},
ImgData{
.name = "1D lod",
.spirv_type = "%float 1D 0 0 0 1 Unknown",
.spirv_fn =
"%result = OpImageSampleExplicitLod %v4float %sampled_image %float_1 Lod %float_1",
.wgsl_type = "",
.wgsl_fn = "f1",
},
ImgData{
.name = "2D vec2",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf12",
.wgsl_type = "",
.wgsl_fn = "vf12",
},
ImgData{
.name = "2D vec2 lod",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn =
"%result = OpImageSampleExplicitLod %v4float %sampled_image %vf12 Lod %float_1",
.wgsl_type = "",
.wgsl_fn = "vf12",
},
ImgData{
.name = "2D vec2 depth",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn =
"%result = OpImageSampleDrefImplicitLod %float %sampled_image %vf12 %float_1",
.wgsl_type = "",
.wgsl_fn = "vf12",
},
ImgData{
.name = "2d vec2 depth lod",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn =
"%result = OpImageSampleDrefExplicitLod %float %sampled_image %vf12 %float_1 "
"Lod %float_0",
.wgsl_type = "",
.wgsl_fn = "vf12",
}));
// In SPIR-V, sampling and dref sampling operations use floating point coordinates. Prove that we
// preserve floating point-ness of the coordinate part, but convert the array index to signed
// integer.Test across all such instructions.
INSTANTIATE_TEST_SUITE_P(
DISABLED_SpirvReaderTest_PreserveFloatCoords_Arrayed,
SampledImageCoordsTest,
::testing::Values(
ImgData{
.name = "2D vec3",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf123",
.wgsl_type = "",
.wgsl_fn = "vf123.xy",
},
ImgData{
.name = "2D lod",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn =
"%result = OpImageSampleExplicitLod %v4float %sampled_image %vf123 Lod %float_1",
.wgsl_type = "",
.wgsl_fn = "vf123.xy",
},
ImgData{
.name = "2D Depth",
.spirv_type = "%float 2D 1 1 0 1 Unknown",
.spirv_fn =
"%result = OpImageSampleDrefImplicitLod %float %sampled_image %vf123 %float_1",
.wgsl_type = "",
.wgsl_fn = "vf123.xy",
},
ImgData{
.name = "2D depth lod",
.spirv_type = "%float 2D 1 1 0 1 Unknown",
.spirv_fn =
"%result = OpImageSampleDrefExplicitLod %float %sampled_image %vf123 %float_1 Lod "
"%float_0",
.wgsl_type = "",
.wgsl_fn = "vf123.xy",
}));
// Metal requires comparison sampling with explicit Level-of-detail to use Lod 0. The SPIR-V reader
// requires the operand to be parsed as a constant 0 value. SPIR-V validation requires the Lod
// parameter to be a floating point value for non-fetch operations. So only test float values.
INSTANTIATE_TEST_SUITE_P(
DISABLED_SpirvReaderTest_ImageSampleDrefExplicitLod_CheckForLod0,
SampledImageCoordsTest,
::testing::Values(
ImgData{
.name = "2d 0.0",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleDrefExplicitLod %float %sampled_image "
"%vf1234 %float_1 Lod %float_0",
.wgsl_type = "",
.wgsl_fn = "vf1234.xy",
},
ImgData{
.name = "2D null",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleDrefExplicitLod %float %sampled_image "
"%vf1234 %float_1 Lod %float_null",
.wgsl_type = "",
.wgsl_fn = "vf1234.xy",
}));
// This is like the previous test, but for Projection sampling.
//
// Metal requires comparison sampling with explicit Level-of-detail to use Lod 0. The SPIR-V reader
// requires the operand to be parsed as a constant 0 value. SPIR-V validation requires the Lod
// parameter to be a floating point value for non-fetch operations. So only test float values.
INSTANTIATE_TEST_SUITE_P(
DISABLED_SpirvReaderTest_ImageSampleProjDrefExplicitLod_CheckForLod0,
SampledImageCoordsTest,
::testing::Values(
ImgData{
.name = "2D 0.0",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleProjDrefExplicitLod %float %sampled_image "
"%vf1234 %float_1 Lod %float_0",
.wgsl_type = "",
.wgsl_fn = "(vf1234.xy / vf1234.z)",
},
ImgData{
.name = "2D null",
// float null works
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleProjDrefExplicitLod %float %sampled_image "
"%vf1234 %float_1 Lod %float_null",
.wgsl_type = "",
.wgsl_fn = "(vf1234.xy / vf1234.z)",
}));
using NonSampledImageCoordsTest = SpirvReaderTestWithParam<ImgData>;
TEST_P(NonSampledImageCoordsTest, MakeCoordinateOperandsForImageAccess) {
auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpCapability Image1D
OpCapability Sampled1D
OpMemoryModel Logical Simple
OpEntryPoint Fragment %100 "main"
OpExecutionMode %100 OriginUpperLeft
OpName %20 "wg"
OpDecorate %20 DescriptorSet 2
OpDecorate %20 Binding 0
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%float = OpTypeFloat 32
%int = OpTypeInt 32 1
%uint = OpTypeInt 32 0
%v2int = OpTypeVector %int 2
%v3int = OpTypeVector %int 3
%v2uint = OpTypeVector %uint 2
%v3uint = OpTypeVector %uint 3
%v4float = OpTypeVector %float 4
%int_1 = OpConstant %int 1
%int_2 = OpConstant %int 2
%int_3 = OpConstant %int 3
%uint_1 = OpConstant %uint 1
%uint_2 = OpConstant %uint 2
%uint_3 = OpConstant %uint 3
%float_1 = OpConstant %float 1
%float_2 = OpConstant %float 2
%float_3 = OpConstant %float 3
%float_4 = OpConstant %float 4
%vi12 = OpConstantComposite %v2int %int_1 %int_2
%vi123 = OpConstantComposite %v3int %int_1 %int_2 %int_3
%vu12 = OpConstantComposite %v2uint %uint_1 %uint_2
%vu123 = OpConstantComposite %v3uint %uint_1 %uint_2 %uint_3
%vf1234 = OpConstantComposite %v4float %float_1 %float_2 %float_3 %float_4
%im_ty = OpTypeImage )" +
params.spirv_type + R"(
%ptr_im_ty = OpTypePointer UniformConstant %im_ty
%20 = OpVariable %ptr_im_ty UniformConstant
%100 = OpFunction %void None %voidfn
%entry = OpLabel
%im = OpLoad %im_ty %20
)" + params.spirv_fn +
R"(
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%wg:ptr<handle, )" +
params.wgsl_type +
R"(, read> = var undef @binding_point(2, 0)
}
%main = @fragment func():void {
$B2: {
%3:)" + params.wgsl_type +
R"( = load %wg)" + params.wgsl_fn + R"(
ret
}
}
)");
}
INSTANTIATE_TEST_SUITE_P(SpirvReaderTest_ImageFetch,
NonSampledImageCoordsTest,
::testing::Values(
ImgData{
.name = "1d signed",
.spirv_type = "%float 1D 0 0 0 1 Unknown",
.spirv_fn = "%result = OpImageFetch %v4float %im %int_1",
.wgsl_type = "texture_1d<f32>",
.wgsl_fn = R"(
%4:vec4<f32> = textureLoad %3, 1i, 0i)",
},
ImgData{
.name = "2d signed",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "%result = OpImageFetch %v4float %im %vi12",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%4:vec4<f32> = textureLoad %3, vec2<i32>(1i, 2i), 0i)",
},
ImgData{
.name = "2D array signed",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn = "%result = OpImageFetch %v4float %im %vi123",
.wgsl_type = "texture_2d_array<f32>",
.wgsl_fn = R"(
%4:vec2<i32> = swizzle vec3<i32>(1i, 2i, 3i), xy
%5:i32 = swizzle vec3<i32>(1i, 2i, 3i), z
%6:vec4<f32> = textureLoad %3, %4, %5, 0i)",
},
ImgData{
.name = "1D unsigned",
.spirv_type = "%float 1D 0 0 0 1 Unknown",
.spirv_fn = "%result = OpImageFetch %v4float %im %uint_1",
.wgsl_type = "texture_1d<f32>",
.wgsl_fn = R"(
%4:vec4<f32> = textureLoad %3, 1u, 0i)",
},
ImgData{
.name = "2D unsigned",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "%result = OpImageFetch %v4float %im %vu12",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%4:vec4<f32> = textureLoad %3, vec2<u32>(1u, 2u), 0i)",
},
ImgData{
.name = "2D array unsigned",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn = "%result = OpImageFetch %v4float %im %vu123",
.wgsl_type = "texture_2d_array<f32>",
.wgsl_fn = R"(
%4:vec2<u32> = swizzle vec3<u32>(1u, 2u, 3u), xy
%5:u32 = swizzle vec3<u32>(1u, 2u, 3u), z
%6:i32 = convert %5
%7:vec4<f32> = textureLoad %3, %4, %6, 0i)",
}));
INSTANTIATE_TEST_SUITE_P(SpirvReaderTest,
NonSampledImageCoordsTest,
::testing::Values(
ImgData{
.name = "1D Read signed",
.spirv_type = "%float 1D 0 0 0 2 R32f",
.spirv_fn = "%result = OpImageRead %v4float %im %int_1",
.wgsl_type = "texture_storage_1d<r32float, read_write>",
.wgsl_fn = R"(
%4:vec4<f32> = textureLoad %3, 1i)",
},
ImgData{
.name = "1D Write signed",
.spirv_type = "%float 1D 0 0 0 2 R32f",
.spirv_fn = "OpImageWrite %im %int_1 %vf1234",
.wgsl_type = "texture_storage_1d<r32float, read_write>",
.wgsl_fn = R"(
%4:void = textureStore %3, 1i, vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f))",
},
ImgData{
.name = "2D read signed",
.spirv_type = "%float 2D 0 0 0 2 R32f",
.spirv_fn = "%result = OpImageRead %v4float %im %vi12",
.wgsl_type = "texture_storage_2d<r32float, read_write>",
.wgsl_fn = R"(
%4:vec4<f32> = textureLoad %3, vec2<i32>(1i, 2i))",
},
ImgData{
.name = "2D write signed",
.spirv_type = "%float 2D 0 0 0 2 R32f",
.spirv_fn = "OpImageWrite %im %vi12 %vf1234",
.wgsl_type = "texture_storage_2d<r32float, read_write>",
.wgsl_fn = R"(
%4:void = textureStore %3, vec2<i32>(1i, 2i), vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f))",
},
ImgData{
.name = "2D array read signed",
.spirv_type = "%float 2D 0 1 0 2 R32f",
.spirv_fn = "%result = OpImageRead %v4float %im %vi123",
.wgsl_type = "texture_storage_2d_array<r32float, read_write>",
.wgsl_fn = R"(
%4:vec2<i32> = swizzle vec3<i32>(1i, 2i, 3i), xy
%5:i32 = swizzle vec3<i32>(1i, 2i, 3i), z
%6:vec4<f32> = textureLoad %3, %4, %5)",
},
ImgData{
.name = "2D array write signed",
.spirv_type = "%float 2D 0 1 0 2 R32f",
.spirv_fn = "OpImageWrite %im %vi123 %vf1234",
.wgsl_type = "texture_storage_2d_array<r32float, read_write>",
.wgsl_fn = R"(
%4:vec2<i32> = swizzle vec3<i32>(1i, 2i, 3i), xy
%5:i32 = swizzle vec3<i32>(1i, 2i, 3i), z
%6:void = textureStore %3, %4, %5, vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f))",
},
ImgData{
.name = "1D read unsigned",
.spirv_type = "%float 1D 0 0 0 2 R32f",
.spirv_fn = "%result = OpImageRead %v4float %im %uint_1",
.wgsl_type = "texture_storage_1d<r32float, read_write>",
.wgsl_fn = R"(
%4:vec4<f32> = textureLoad %3, 1u)",
},
ImgData{
.name = "1D write unsigned",
.spirv_type = "%float 1D 0 0 0 2 R32f",
.spirv_fn = "OpImageWrite %im %uint_1 %vf1234",
.wgsl_type = "texture_storage_1d<r32float, read_write>",
.wgsl_fn = R"(
%4:void = textureStore %3, 1u, vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f))",
},
ImgData{
.name = "2D read unsigned",
.spirv_type = "%float 2D 0 0 0 2 R32f",
.spirv_fn = "%result = OpImageRead %v4float %im %vu12",
.wgsl_type = "texture_storage_2d<r32float, read_write>",
.wgsl_fn = R"(
%4:vec4<f32> = textureLoad %3, vec2<u32>(1u, 2u))",
},
ImgData{
.name = "2D write unsigned",
.spirv_type = "%float 2D 0 0 0 2 R32f",
.spirv_fn = "OpImageWrite %im %vu12 %vf1234",
.wgsl_type = "texture_storage_2d<r32float, read_write>",
.wgsl_fn = R"(
%4:void = textureStore %3, vec2<u32>(1u, 2u), vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f))",
},
ImgData{
.name = "2D array read unsigned",
.spirv_type = "%float 2D 0 1 0 2 R32f",
.spirv_fn = "%result = OpImageRead %v4float %im %vu123",
.wgsl_type = "texture_storage_2d_array<r32float, read_write>",
.wgsl_fn = R"(
%4:vec2<u32> = swizzle vec3<u32>(1u, 2u, 3u), xy
%5:u32 = swizzle vec3<u32>(1u, 2u, 3u), z
%6:i32 = convert %5
%7:vec4<f32> = textureLoad %3, %4, %6)",
},
ImgData{
.name = "2D array write unsigned",
.spirv_type = "%float 2D 0 1 0 2 R32f",
.spirv_fn = "OpImageWrite %im %vu123 %vf1234",
.wgsl_type = "texture_storage_2d_array<r32float, read_write>",
.wgsl_fn = R"(
%4:vec2<u32> = swizzle vec3<u32>(1u, 2u, 3u), xy
%5:u32 = swizzle vec3<u32>(1u, 2u, 3u), z
%6:i32 = convert %5
%7:void = textureStore %3, %4, %6, vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f))",
}));
// An ad-hoc test to prove we never had the issue feared in crbug.com/tint/265.
// Never create a const-declaration for a pointer to a texture or sampler. Code generation always
// traces back to the memory object declaration.
TEST_F(SpirvReaderTest, NeverGenerateConstDeclForHandle_UseVariableDirectly) {
EXPECT_IR(R"(
OpCapability Shader
OpCapability Sampled1D
OpCapability Image1D
OpCapability StorageImageExtendedFormats
OpCapability ImageQuery
OpMemoryModel Logical Simple
OpEntryPoint Fragment %100 "main"
OpExecutionMode %100 OriginUpperLeft
OpName %var "var"
OpDecorate %var_im DescriptorSet 0
OpDecorate %var_im Binding 0
OpDecorate %var_s DescriptorSet 0
OpDecorate %var_s Binding 1
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%v2float = OpTypeVector %float 2
%v2_0 = OpConstantNull %v2float
%im = OpTypeImage %float 2D 0 0 0 1 Unknown
%si = OpTypeSampledImage %im
%s = OpTypeSampler
%ptr_im = OpTypePointer UniformConstant %im
%ptr_s = OpTypePointer UniformConstant %s
%var_im = OpVariable %ptr_im UniformConstant
%var_s = OpVariable %ptr_s UniformConstant
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%ptr_v4 = OpTypePointer Function %v4float
%100 = OpFunction %void None %voidfn
%entry = OpLabel
%var = OpVariable %ptr_v4 Function
; Try to induce generating a const-declaration of a pointer to
; a sampler or texture.
%var_im_copy = OpCopyObject %ptr_im %var_im
%var_s_copy = OpCopyObject %ptr_s %var_s
%im0 = OpLoad %im %var_im_copy
%s0 = OpLoad %s %var_s_copy
%si0 = OpSampledImage %si %im0 %s0
%t0 = OpImageSampleImplicitLod %v4float %si0 %v2_0
%im1 = OpLoad %im %var_im_copy
%s1 = OpLoad %s %var_s_copy
%si1 = OpSampledImage %si %im1 %s1
%t1 = OpImageSampleImplicitLod %v4float %si1 %v2_0
%sum = OpFAdd %v4float %t0 %t1
OpStore %var %sum
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%1:ptr<handle, texture_2d<f32>, read> = var undef @binding_point(0, 0)
%2:ptr<handle, sampler, read> = var undef @binding_point(0, 1)
}
%main = @fragment func():void {
$B2: {
%var:ptr<function, vec4<f32>, read_write> = var undef
%5:texture_2d<f32> = load %1
%6:sampler = load %2
%7:vec4<f32> = textureSample %5, %6, vec2<f32>(0.0f)
%8:texture_2d<f32> = load %1
%9:sampler = load %2
%10:vec4<f32> = textureSample %8, %9, vec2<f32>(0.0f)
%11:vec4<f32> = add %7, %10
store %var, %11
ret
}
}
)");
}
// When a sampler is loaded in an enclosing structured construct, don't generate a variable for it.
// The ordinary tracing logic will find the originating variable anyway.
// https://crbug.com/tint/1839
TEST_F(SpirvReaderTest, SamplerLoadedInEnclosingConstruct_DontGenerateVar) {
EXPECT_IR(R"(
OpCapability Shader
OpCapability Sampled1D
OpCapability Image1D
OpCapability StorageImageExtendedFormats
OpCapability ImageQuery
OpMemoryModel Logical Simple
OpEntryPoint Fragment %100 "main"
OpExecutionMode %100 OriginUpperLeft
OpName %var_im "var_im"
OpName %var_s "var_s"
OpDecorate %var_im DescriptorSet 0
OpDecorate %var_im Binding 0
OpDecorate %var_s DescriptorSet 0
OpDecorate %var_s Binding 1
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%v2float = OpTypeVector %float 2
%v2_0 = OpConstantNull %v2float
%im = OpTypeImage %float 2D 0 0 0 1 Unknown
%si = OpTypeSampledImage %im
%s = OpTypeSampler
%ptr_im = OpTypePointer UniformConstant %im
%ptr_s = OpTypePointer UniformConstant %s
%var_im = OpVariable %ptr_im UniformConstant
%var_s = OpVariable %ptr_s UniformConstant
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%ptr_v4 = OpTypePointer Function %v4float
%bool = OpTypeBool
%true = OpConstantTrue %bool
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%100 = OpFunction %void None %voidfn
%entry = OpLabel
%1 = OpLoad %im %var_im
%2 = OpLoad %s %var_s
OpSelectionMerge %90 None
OpSwitch %int_0 %20
%20 = OpLabel
OpSelectionMerge %80 None
OpBranchConditional %true %30 %80
%30 = OpLabel
%si0 = OpSampledImage %si %1 %2
%t0 = OpImageSampleImplicitLod %v4float %si0 %v2_0
%t1 = OpImageSampleImplicitLod %v4float %si0 %v2_0
OpBranch %80
%80 = OpLabel
OpBranch %90
%90 = OpLabel
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%var_im:ptr<handle, texture_2d<f32>, read> = var undef @binding_point(0, 0)
%var_s:ptr<handle, sampler, read> = var undef @binding_point(0, 1)
}
%main = @fragment func():void {
$B2: {
%4:texture_2d<f32> = load %var_im
%5:sampler = load %var_s
switch 0i [c: (default, $B3)] { # switch_1
$B3: { # case
if true [t: $B4, f: $B5] { # if_1
$B4: { # true
%6:vec4<f32> = textureSample %4, %5, vec2<f32>(0.0f)
%7:vec4<f32> = textureSample %4, %5, vec2<f32>(0.0f)
exit_if # if_1
}
$B5: { # false
exit_if # if_1
}
}
exit_switch # switch_1
}
}
ret
}
}
)");
}
TEST_F(SpirvReaderTest, SamplerLoaded_MultiUsage) {
EXPECT_IR(R"(
OpCapability Shader
OpCapability Sampled1D
OpCapability Image1D
OpCapability StorageImageExtendedFormats
OpCapability ImageQuery
OpMemoryModel Logical Simple
OpEntryPoint Fragment %100 "main"
OpExecutionMode %100 OriginUpperLeft
OpName %var_im "var_im"
OpName %var_s "var_s"
OpDecorate %var_im DescriptorSet 0
OpDecorate %var_im Binding 0
OpDecorate %var_s DescriptorSet 0
OpDecorate %var_s Binding 1
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%v2float = OpTypeVector %float 2
%v2_0 = OpConstantNull %v2float
%im = OpTypeImage %float 2D 0 0 0 1 Unknown
%si = OpTypeSampledImage %im
%s = OpTypeSampler
%ptr_im = OpTypePointer UniformConstant %im
%ptr_s = OpTypePointer UniformConstant %s
%var_im = OpVariable %ptr_im UniformConstant
%var_s = OpVariable %ptr_s UniformConstant
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%100 = OpFunction %void None %voidfn
%entry = OpLabel
%1 = OpLoad %im %var_im
%2 = OpLoad %s %var_s
%si0 = OpSampledImage %si %1 %2
%t0 = OpImageSampleImplicitLod %v4float %si0 %v2_0
%si1 = OpSampledImage %si %1 %2
%t2 = OpImageSampleImplicitLod %v4float %si1 %v2_0
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%var_im:ptr<handle, texture_2d<f32>, read> = var undef @binding_point(0, 0)
%var_s:ptr<handle, sampler, read> = var undef @binding_point(0, 1)
}
%main = @fragment func():void {
$B2: {
%4:texture_2d<f32> = load %var_im
%5:sampler = load %var_s
%6:vec4<f32> = textureSample %4, %5, vec2<f32>(0.0f)
%7:vec4<f32> = textureSample %4, %5, vec2<f32>(0.0f)
ret
}
}
)");
}
// Demonstrates fix for crbug.com/tint/1646
// The problem is the coordinate for an image operation can be a combinatorial value that has been
// hoisted out to a 'var' declaration.
//
// In this test (and the original case form the bug), the definition for the value is in an outer
// construct, and the image operation using it is in a doubly nested construct.
//
// The coordinate handling has to unwrap the ref type it made for the 'var' declaration.
TEST_F(SpirvReaderTest, ImageCoordinateCanBeHoistedConstant) {
EXPECT_IR(R"(
OpCapability Shader
OpCapability Sampled1D
OpCapability Image1D
OpCapability StorageImageExtendedFormats
OpCapability ImageQuery
OpMemoryModel Logical Simple
OpEntryPoint Fragment %100 "main"
OpExecutionMode %100 OriginUpperLeft
OpDecorate %10 DescriptorSet 0
OpDecorate %10 Binding 0
OpDecorate %20 DescriptorSet 2
OpDecorate %20 Binding 1
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%bool = OpTypeBool
%true = OpConstantTrue %bool
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%float_null = OpConstantNull %float
%sampler = OpTypeSampler
%ptr_sampler = OpTypePointer UniformConstant %sampler
%im_ty = OpTypeImage %float 1D 0 0 0 1 Unknown
%ptr_im_ty = OpTypePointer UniformConstant %im_ty
%si_ty = OpTypeSampledImage %im_ty
%10 = OpVariable %ptr_sampler UniformConstant
%20 = OpVariable %ptr_im_ty UniformConstant
%100 = OpFunction %void None %voidfn
%entry = OpLabel
%900 = OpCopyObject %float %float_null ; definition here
OpSelectionMerge %99 None
OpBranchConditional %true %40 %99
%40 = OpLabel
OpSelectionMerge %80 None
OpBranchConditional %true %50 %80
%50 = OpLabel
%sam = OpLoad %sampler %10
%im = OpLoad %im_ty %20
%sampled_image = OpSampledImage %si_ty %im %sam
%result = OpImageSampleImplicitLod %v4float %sampled_image %900 ; usage here
OpBranch %80
%80 = OpLabel
OpBranch %99
%99 = OpLabel
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%1:ptr<handle, sampler, read> = var undef @binding_point(0, 0)
%2:ptr<handle, texture_1d<f32>, read> = var undef @binding_point(2, 1)
}
%main = @fragment func():void {
$B2: {
%4:f32 = let 0.0f
if true [t: $B3, f: $B4] { # if_1
$B3: { # true
if true [t: $B5, f: $B6] { # if_2
$B5: { # true
%5:sampler = load %1
%6:texture_1d<f32> = load %2
%7:vec4<f32> = textureSample %6, %5, %4
exit_if # if_2
}
$B6: { # false
exit_if # if_2
}
}
exit_if # if_1
}
$B4: { # false
exit_if # if_1
}
}
ret
}
}
)");
}
// Demonstrates fix for crbug.com/tint/1712
// The problem is the coordinate for an image operation can be a combinatorial value that has been
// hoisted out to a 'var' declaration.
//
// In this test (and the original case form the bug), the definition for the value is in an outer
// construct, and the image operation using it is in a doubly nested construct.
//
// The coordinate handling has to unwrap the ref type it made for the 'var' declaration.
TEST_F(SpirvReaderTest, ImageCoordinateCanBeHoistedVector) {
EXPECT_IR(R"(
OpCapability Shader
OpCapability Sampled1D
OpCapability Image1D
OpCapability StorageImageExtendedFormats
OpCapability ImageQuery
OpMemoryModel Logical Simple
OpEntryPoint Fragment %100 "main"
OpExecutionMode %100 OriginUpperLeft
OpDecorate %10 DescriptorSet 0
OpDecorate %10 Binding 0
OpDecorate %20 DescriptorSet 2
OpDecorate %20 Binding 1
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%bool = OpTypeBool
%true = OpConstantTrue %bool
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%v4float = OpTypeVector %float 4
%v2float_null = OpConstantNull %v2float
%sampler = OpTypeSampler
%ptr_sampler = OpTypePointer UniformConstant %sampler
%im_ty = OpTypeImage %float 1D 0 0 0 1 Unknown
%ptr_im_ty = OpTypePointer UniformConstant %im_ty
%si_ty = OpTypeSampledImage %im_ty
%10 = OpVariable %ptr_sampler UniformConstant
%20 = OpVariable %ptr_im_ty UniformConstant
%100 = OpFunction %void None %voidfn
%entry = OpLabel
%900 = OpCopyObject %v2float %v2float_null ; definition here
OpSelectionMerge %99 None
OpBranchConditional %true %40 %99
%40 = OpLabel
OpSelectionMerge %80 None
OpBranchConditional %true %50 %80
%50 = OpLabel
%sam = OpLoad %sampler %10
%im = OpLoad %im_ty %20
%sampled_image = OpSampledImage %si_ty %im %sam
%result = OpImageSampleImplicitLod %v4float %sampled_image %900 ; usage here
OpBranch %80
%80 = OpLabel
OpBranch %99
%99 = OpLabel
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%1:ptr<handle, sampler, read> = var undef @binding_point(0, 0)
%2:ptr<handle, texture_1d<f32>, read> = var undef @binding_point(2, 1)
}
%main = @fragment func():void {
$B2: {
%4:vec2<f32> = let vec2<f32>(0.0f)
if true [t: $B3, f: $B4] { # if_1
$B3: { # true
if true [t: $B5, f: $B6] { # if_2
$B5: { # true
%5:sampler = load %1
%6:texture_1d<f32> = load %2
%7:f32 = swizzle %4, x
%8:vec4<f32> = textureSample %6, %5, %7
exit_if # if_2
}
$B6: { # false
exit_if # if_2
}
}
exit_if # if_1
}
$B4: { # false
exit_if # if_1
}
}
ret
}
}
)");
}
// Demonstrates fix for crbug.com/tint/1642 The problem is the texel value for an image write can be
// given in 'var' declaration.
//
// The texel value handling has to unwrap the ref type first.
TEST_F(SpirvReaderTest, DISABLED_TexelTypeWhenLoop) {
EXPECT_IR(R"(
OpCapability Shader
OpCapability StorageImageExtendedFormats
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %100 "main"
OpExecutionMode %100 LocalSize 8 8 1
OpSource HLSL 600
OpName %type_2d_image "type.2d.image"
OpName %Output2Texture2D "Output2Texture2D"
OpName %100 "main"
OpDecorate %Output2Texture2D DescriptorSet 0
OpDecorate %Output2Texture2D Binding 0
%float = OpTypeFloat 32
%float_0 = OpConstant %float 0
%v2float = OpTypeVector %float 2
%7 = OpConstantComposite %v2float %float_0 %float_0
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%int_2 = OpConstant %int 2
%float_1 = OpConstant %float 1
%12 = OpConstantComposite %v2float %float_1 %float_1
%int_1 = OpConstant %int 1
%uint = OpTypeInt 32 0
%uint_1 = OpConstant %uint 1
%v2uint = OpTypeVector %uint 2
%17 = OpConstantComposite %v2uint %uint_1 %uint_1
%type_2d_image = OpTypeImage %float 2D 2 0 0 2 Rg32f
%_ptr_UniformConstant_type_2d_image = OpTypePointer UniformConstant %type_2d_image
%void = OpTypeVoid
%20 = OpTypeFunction %void
%bool = OpTypeBool
%Output2Texture2D = OpVariable %_ptr_UniformConstant_type_2d_image UniformConstant
%100 = OpFunction %void None %20
%22 = OpLabel
OpBranch %23
%23 = OpLabel
%24 = OpPhi %v2float %7 %22 %12 %25
%26 = OpPhi %int %int_0 %22 %27 %25
%28 = OpSLessThan %bool %26 %int_2
OpLoopMerge %29 %25 None
OpBranchConditional %28 %25 %29
%25 = OpLabel
%27 = OpIAdd %int %26 %int_1
OpBranch %23
%29 = OpLabel
%30 = OpLoad %type_2d_image %Output2Texture2D
OpImageWrite %30 %17 %24 None
OpReturn
OpFunctionEnd
)",
R"(
var x_24 : vec2f;
var x_26 : i32;
x_24 = vec2f();
x_26 = 0i;
loop {
var x_27 : i32;
if ((x_26 < 2i)) {
} else {
break;
}
continuing {
x_27 = (x_26 + 1i);
x_24 = vec2f(1.0f);
x_26 = x_27;
}
}
textureStore(Output2Texture2D, vec2i(vec2u(1u)), vec4f(x_24, 0.0f, 0.0f));
return;
)");
}
TEST_F(SpirvReaderTest, ReadWriteStorageTexture) {
EXPECT_IR(R"(
OpCapability Shader
OpCapability StorageImageExtendedFormats
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %100 "main"
OpExecutionMode %100 LocalSize 8 8 1
OpSource HLSL 600
OpName %type_2d_image "type.2d.image"
OpName %RWTexture2D "RWTexture2D"
OpName %100 "main"
OpDecorate %RWTexture2D DescriptorSet 0
OpDecorate %RWTexture2D Binding 0
%float = OpTypeFloat 32
%float_0 = OpConstant %float 0
%v4float = OpTypeVector %float 4
%uint = OpTypeInt 32 0
%uint_1 = OpConstant %uint 1
%v2uint = OpTypeVector %uint 2
%coord = OpConstantComposite %v2uint %uint_1 %uint_1
%type_2d_image = OpTypeImage %float 2D 2 0 0 2 Rgba32f
%_ptr_UniformConstant_type_2d_image = OpTypePointer UniformConstant %type_2d_image
%void = OpTypeVoid
%20 = OpTypeFunction %void
%RWTexture2D = OpVariable %_ptr_UniformConstant_type_2d_image UniformConstant
%100 = OpFunction %void None %20
%22 = OpLabel
%30 = OpLoad %type_2d_image %RWTexture2D
%31 = OpImageRead %v4float %30 %coord None
%32 = OpFAdd %v4float %31 %31
OpImageWrite %30 %coord %32 None
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%RWTexture2D:ptr<handle, texture_storage_2d<rgba32float, read_write>, read> = var undef @binding_point(0, 0)
}
%main = @compute @workgroup_size(8u, 8u, 1u) func():void {
$B2: {
%3:texture_storage_2d<rgba32float, read_write> = load %RWTexture2D
%4:vec4<f32> = textureLoad %3, vec2<u32>(1u)
%5:vec4<f32> = add %4, %4
%6:void = textureStore %3, vec2<u32>(1u), %5
ret
}
}
)");
}
TEST_F(SpirvReaderTest, Image_UnknownDepth_NonDepthUse) {
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %1 "main"
OpExecutionMode %1 OriginUpperLeft
OpDecorate %7 DescriptorSet 0
OpDecorate %7 Binding 0
OpDecorate %8 DescriptorSet 0
OpDecorate %8 Binding 1
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%4 = OpTypeImage %float 2D 2 0 0 1 Unknown
%5 = OpTypeSampler
%v2float = OpTypeVector %float 2
%void = OpTypeVoid
%15 = OpTypeFunction %void
%6 = OpTypeSampledImage %4
%_ptr_UniformConstant_4 = OpTypePointer UniformConstant %4
%7 = OpVariable %_ptr_UniformConstant_4 UniformConstant
%_ptr_UniformConstant_5 = OpTypePointer UniformConstant %5
%8 = OpVariable %_ptr_UniformConstant_5 UniformConstant
%const = OpConstantNull %v2float
%1 = OpFunction %void None %15
%18 = OpLabel
%20 = OpLoad %4 %7
%21 = OpLoad %5 %8
%22 = OpSampledImage %6 %20 %21
%23 = OpImageSampleImplicitLod %v4float %22 %const None
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%1:ptr<handle, texture_2d<f32>, read> = var undef @binding_point(0, 0)
%2:ptr<handle, sampler, read> = var undef @binding_point(0, 1)
}
%main = @fragment func():void {
$B2: {
%4:texture_2d<f32> = load %1
%5:sampler = load %2
%6:vec4<f32> = textureSample %4, %5, vec2<f32>(0.0f)
ret
}
}
)");
}
TEST_F(SpirvReaderTest, Image_UnknownDepth_DepthUse) {
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %1 "main"
OpExecutionMode %1 OriginUpperLeft
OpDecorate %7 DescriptorSet 0
OpDecorate %7 Binding 0
OpDecorate %8 DescriptorSet 0
OpDecorate %8 Binding 1
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%4 = OpTypeImage %float 2D 2 0 0 1 Unknown
%5 = OpTypeSampler
%v2float = OpTypeVector %float 2
%void = OpTypeVoid
%15 = OpTypeFunction %void
%6 = OpTypeSampledImage %4
%_ptr_UniformConstant_4 = OpTypePointer UniformConstant %4
%7 = OpVariable %_ptr_UniformConstant_4 UniformConstant
%_ptr_UniformConstant_5 = OpTypePointer UniformConstant %5
%8 = OpVariable %_ptr_UniformConstant_5 UniformConstant
%const = OpConstantNull %v2float
%float_0 = OpConstantNull %float
%1 = OpFunction %void None %15
%18 = OpLabel
%20 = OpLoad %4 %7
%21 = OpLoad %5 %8
%22 = OpSampledImage %6 %20 %21
%23 = OpImageSampleDrefImplicitLod %float %22 %const %float_0 None
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%1:ptr<handle, texture_depth_2d, read> = var undef @binding_point(0, 0)
%2:ptr<handle, sampler_comparison, read> = var undef @binding_point(0, 1)
}
%main = @fragment func():void {
$B2: {
%4:texture_depth_2d = load %1
%5:sampler_comparison = load %2
%6:f32 = textureSampleCompare %4, %5, vec2<f32>(0.0f), 0.0f
ret
}
}
)");
}
} // namespace
} // namespace tint::spirv::reader