Google Git Google Git
Sign in
dawn/dawn/cfef9491663b6e754c73dd9e480eadc61bdf760c/./src/tint/lang/spirv/reader/texture_test.cc
blob: bdce80cf569e90f2a0e4bb0c8c6fa6863dcd54e2 [file]
// 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_SameImageType_TwoVarsWithDifferentAccessModes) {
EXPECT_IR(R"(
OpCapability Shader
OpCapability Sampled1D
OpCapability StorageImageExtendedFormats
OpMemoryModel Logical Simple
OpEntryPoint Fragment %main "main"
OpExecutionMode %main OriginUpperLeft
OpName %write_only_image "write_only_image"
OpName %read_only_image "read_only_image"
OpDecorate %write_only_image DescriptorSet 0
OpDecorate %write_only_image Binding 0
OpDecorate %write_only_image NonReadable
OpDecorate %read_only_image DescriptorSet 0
OpDecorate %read_only_image Binding 1
OpDecorate %read_only_image NonWritable
%float = OpTypeFloat 32
%storage = OpTypeImage %float 1D 0 0 0 2 Rg32f
%ptr_storage = OpTypePointer UniformConstant %storage
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%write_only_image = OpVariable %ptr_storage UniformConstant
%read_only_image = OpVariable %ptr_storage UniformConstant
%main = OpFunction %void None %voidfn
%entry = OpLabel
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%write_only_image:ptr<handle, texture_storage_1d<rg32float, write>, read> = var undef @binding_point(0, 0)
%read_only_image:ptr<handle, texture_storage_1d<rg32float, read>, read> = var undef @binding_point(0, 1)
}
%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, Image) {
EXPECT_IR(R"(
OpCapability Shader
OpCapability ImageQuery
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
%main = OpFunction %void None %voidfn
%entry = OpLabel
%sam = OpLoad %sampler %var_sampler
%im = OpLoad %tex %10
%sampled_image = OpSampledImage %sampled_img %im %sam
%result = OpImage %tex %sampled_image
%size = OpImageQueryLevels %int %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:u32 = textureNumLevels %5
%7:i32 = convert %6
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_5 = OpConstant %int 5
%int_neg_4 = OpConstant %int -4
%offset2i = OpConstantComposite %v2int %int_5 %int_neg_4
%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>(5i, -4i)
%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_5 = OpConstant %int 5
%int_neg_4 = OpConstant %int -4
%offset2i = OpConstantComposite %v2int %int_5 %int_neg_4
%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>(5i, -4i)
%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_5 = OpConstant %int 5
%int_neg_4 = OpConstant %int -4
%offset2i = OpConstantComposite %v2int %int_5 %int_neg_4
%uint_5 = OpConstant %uint 5
%uint_6 = OpConstant %uint 6
%offset2u = OpConstantComposite %v2uint %uint_5 %uint_6
%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>(5i, -4i))",
},
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:vec4<f32> = textureGather 1i, %5, %4, vec2<f32>(1.0f, 2.0f), vec2<i32>(5i, 6i))",
},
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:vec4<f32> = textureGather 1i, %5, %4, vec2<f32>(1.0f, 2.0f), 3i)",
},
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:vec4<f32> = textureGather 1i, %5, %4, vec2<f32>(1.0f, 2.0f), 3i, vec2<i32>(5i, -4i))",
},
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:vec4<f32> = textureGather 1i, %5, %4, vec2<f32>(1.0f, 2.0f), 3i, vec2<i32>(5i, 6i))",
},
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>(5i, -4i))",
},
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:vec4<f32> = textureGather %5, %4, vec2<f32>(1.0f, 2.0f), vec2<i32>(5i, 6i))",
},
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:vec4<f32> = textureGather %5, %4, vec2<f32>(1.0f, 2.0f), 3i)",
},
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:vec4<f32> = textureGather %5, %4, vec2<f32>(1.0f, 2.0f), 3i, vec2<i32>(5i, -4i))",
},
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:vec4<f32> = textureGather %5, %4, vec2<f32>(1.0f, 2.0f), 3i, vec2<i32>(5i, 6i))",
},
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:vec4<f32> = textureGather 1i, %5, %4, vec3<f32>(1.0f, 2.0f, 3.0f), 4i)",
},
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:vec4<f32> = textureGather %5, %4, vec3<f32>(1.0f, 2.0f, 3.0f), 4i)",
}));
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>(5i, -4i))",
},
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>(5i, -4i))",
},
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:vec4<f32> = textureSampleBias %5, %4, vec2<f32>(1.0f, 2.0f), 5.0f, vec2<i32>(5i, 6i))",
},
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:vec4<f32> = textureSample %5, %4, vec2<f32>(1.0f, 2.0f), 3i)",
},
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:vec4<f32> = textureSample %5, %4, vec2<f32>(1.0f, 2.0f), 3i, vec2<i32>(5i, -4i))",
},
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:vec4<f32> = textureSampleBias %5, %4, vec2<f32>(1.0f, 2.0f), 3i, 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:vec4<f32> = textureSampleBias %5, %4, vec2<f32>(1.0f, 2.0f), 3i, 5.0f, vec2<i32>(5i, -4i))",
},
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:vec4<f32> = textureSampleBias %5, %4, vec2<f32>(1.0f, 2.0f), 3i, 5.0f, vec2<i32>(5i, -4i))",
}));
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:vec4<f32> = textureSample %5, %4, 0.5f)",
},
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:vec4<f32> = textureSample %5, %4, vec2<f32>(0.3333333432674407959f, 0.6666666865348815918f))",
},
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:vec4<f32> = textureSample %5, %4, vec3<f32>(0.25f, 0.5f, 0.75f))",
},
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:vec4<f32> = textureSample %5, %4, vec2<f32>(0.3333333432674407959f, 0.6666666865348815918f), vec2<i32>(5i, -4i))",
},
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:vec4<f32> = textureSampleBias %5, %4, vec2<f32>(0.3333333432674407959f, 0.6666666865348815918f), 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:vec4<f32> = textureSampleBias %5, %4, vec2<f32>(0.3333333432674407959f, 0.6666666865348815918f), 5.0f, vec2<i32>(5i, -4i))",
},
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:vec4<f32> = textureSampleBias %5, %4, vec2<f32>(0.3333333432674407959f, 0.6666666865348815918f), 5.0f, vec2<i32>(5i, 6i))",
},
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:f32 = textureSample %5, %4, vec2<f32>(0.3333333432674407959f, 0.6666666865348815918f)
%7:vec4<f32> = construct %6, 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:vec4<f32> = textureSampleLevel %5, %4, vec2<f32>(1.0f, 2.0f), 3i, 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>(5i, -4i))",
},
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:vec4<f32> = textureSampleLevel %5, %4, vec2<f32>(1.0f, 2.0f), 0.0f, vec2<i32>(5i, 6i))",
},
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:vec4<f32> = textureSampleLevel %5, %4, vec2<f32>(1.0f, 2.0f), 3i, 0.0f, vec2<i32>(5i, -4i))",
},
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:vec4<f32> = textureSampleGrad %5, %4, vec2<f32>(1.0f, 2.0f), 3i, 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>(5i, -4i))",
},
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: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>(5i, 6i))",
},
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:vec4<f32> = textureSampleGrad %5, %4, vec2<f32>(1.0f, 2.0f), 3i, vec2<f32>(1.0f, 2.0f), vec2<f32>(2.0f, 1.0f), vec2<i32>(5i, -4i))",
},
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:vec4<f32> = textureSampleGrad %5, %4, vec2<f32>(1.0f, 2.0f), 3i, vec2<f32>(1.0f, 2.0f), vec2<f32>(2.0f, 1.0f), vec2<i32>(5i, 6i))",
},
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:f32 = textureSampleLevel %5, %4, vec2<f32>(1.0f, 2.0f), 1i
%7:vec4<f32> = construct %6, 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:vec4<f32> = textureSampleLevel %5, %4, vec2<f32>(0.3333333432674407959f, 0.6666666865348815918f), 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:vec4<f32> = textureSampleLevel %5, %4, vec2<f32>(0.3333333432674407959f, 0.6666666865348815918f), 1.0f, vec2<i32>(5i, -4i))",
},
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:vec4<f32> = textureSampleGrad %5, %4, vec2<f32>(0.3333333432674407959f, 0.6666666865348815918f), 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:vec4<f32> = textureSampleGrad %5, %4, vec2<f32>(0.3333333432674407959f, 0.6666666865348815918f), vec2<f32>(1.0f, 2.0f), vec2<f32>(2.0f, 1.0f), vec2<i32>(5i, -4i))",
}));
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
%float_null = OpConstantNull %float
%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_5 = OpConstant %int 5
%int_neg_4 = OpConstant %int -4
%offset2i = OpConstantComposite %v2int %int_5 %int_neg_4
%uint_5 = OpConstant %uint 5
%uint_6 = OpConstant %uint 6
%offset2u = OpConstantComposite %v2uint %uint_5 %uint_6
%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_ImageDrefGather,
SamplerComparisonTest,
::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 = R"(
%6:vec4<f32> = textureGatherCompare %5, %4, 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 = R"(
%6:vec4<f32> = textureGatherCompare %5, %4, vec2<f32>(1.0f, 2.0f), 1.0f, vec2<i32>(5i, -4i))",
},
ImgData{
.name = "2d Depth ConstOffset Unsigned",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn =
"OpImageDrefGather %v4float %sampled_image %coords2 %depth ConstOffset %offset2u",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = R"(
%6:vec4<f32> = textureGatherCompare %5, %4, vec2<f32>(1.0f, 2.0f), 1.0f, vec2<i32>(5i, 6i))",
},
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 = R"(
%6:vec4<f32> = textureGatherCompare %5, %4, 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_2d_array",
.wgsl_fn = R"(
%6:vec4<f32> = textureGatherCompare %5, %4, vec2<f32>(1.0f, 2.0f), 3i, 1.0f, vec2<i32>(5i, -4i))",
},
ImgData{
.name = "2d Depth Array ConstOffset Unsigned",
.spirv_type = "%float 2D 1 1 0 1 Unknown",
.spirv_fn =
"OpImageDrefGather %v4float %sampled_image %coords3 %depth ConstOffset %offset2u",
.wgsl_type = "texture_depth_2d_array",
.wgsl_fn = R"(
%6:vec4<f32> = textureGatherCompare %5, %4, vec2<f32>(1.0f, 2.0f), 3i, 1.0f, vec2<i32>(5i, 6i))",
},
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 = R"(
%6:vec4<f32> = textureGatherCompare %5, %4, 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 = R"(
%6:vec4<f32> = textureGatherCompare %5, %4, vec3<f32>(1.0f, 2.0f, 3.0f), 4i, 1.0f)",
}));
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>(5i, -4i))",
},
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:f32 = textureSampleCompare %5, %4, vec2<f32>(1.0f, 2.0f), 3i, 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:f32 = textureSampleCompare %5, %4, vec2<f32>(1.0f, 2.0f), 3i, 1.0f, vec2<i32>(5i, -4i))",
},
ImgData{
.name = "2D vec2 depth",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "OpImageSampleDrefImplicitLod %float %sampled_image %coords2 %float_1",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = R"(
%6:f32 = textureSampleCompare %5, %4, vec2<f32>(1.0f, 2.0f), 1.0f)",
}));
INSTANTIATE_TEST_SUITE_P(
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 = R"(
%6:f32 = textureSampleCompareLevel %5, %4, 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 = R"(
%6:f32 = textureSampleCompareLevel %5, %4, 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 = R"(
%6:f32 = textureSampleCompareLevel %5, %4, vec2<f32>(1.0f, 2.0f), 1.0f, vec2<i32>(5i, -4i))",
},
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 = R"(
%6:f32 = textureSampleCompareLevel %5, %4, vec2<f32>(1.0f, 2.0f), 3i, 1.0f, vec2<i32>(5i, -4i))",
},
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 = R"(
%6:f32 = textureSampleCompareLevel %5, %4, 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 = R"(
%6:f32 = textureSampleCompareLevel %5, %4, vec3<f32>(1.0f, 2.0f, 3.0f), 4i, 1.0f)",
},
ImgData{
.name = "2d vec2 depth lod",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn =
"OpImageSampleDrefExplicitLod %float %sampled_image %coords2 %float_1 Lod %float_0",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = R"(
%6:f32 = textureSampleCompareLevel %5, %4, vec2<f32>(1.0f, 2.0f), 1.0f)",
}
));
INSTANTIATE_TEST_SUITE_P(
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 = R"(
%6:f32 = textureSampleCompare %5, %4, vec2<f32>(0.3333333432674407959f, 0.6666666865348815918f), 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 = R"(
%6:f32 = textureSampleCompare %5, %4, vec2<f32>(0.3333333432674407959f, 0.6666666865348815918f), 1.0f, vec2<i32>(5i, -4i))",
}));
INSTANTIATE_TEST_SUITE_P(
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 = R"(
%6:f32 = textureSampleCompareLevel %5, %4, vec2<f32>(0.3333333432674407959f, 0.6666666865348815918f), 1.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 = R"(
%6:f32 = textureSampleCompareLevel %5, %4, vec2<f32>(0.3333333432674407959f, 0.6666666865348815918f), 1.0f, vec2<i32>(5i, -4i))",
}));
// 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(
SpirvReaderTest_ImageSampleDrefExplicitLod_CheckForLod0,
SamplerComparisonTest,
::testing::Values(
ImgData{
.name = "2d 0.0",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn =
"OpImageSampleDrefExplicitLod %float %sampled_image %coords4 %float_1 Lod %float_0",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = R"(
%6:f32 = textureSampleCompareLevel %5, %4, vec2<f32>(1.0f, 2.0f), 1.0f)",
},
ImgData{
.name = "2D null",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "OpImageSampleDrefExplicitLod %float %sampled_image %coords4 %float_1 Lod "
"%float_null",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = R"(
%6:f32 = textureSampleCompareLevel %5, %4, vec2<f32>(1.0f, 2.0f), 1.0f)",
}));
// 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(
SpirvReaderTest_ImageSampleProjDrefExplicitLod_CheckForLod0,
SamplerComparisonTest,
::testing::Values(
ImgData{
.name = "2D 0.0",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "OpImageSampleProjDrefExplicitLod %float %sampled_image %coords4 %float_1 "
"Lod %float_0",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = R"(
%6:f32 = textureSampleCompareLevel %5, %4, vec2<f32>(0.3333333432674407959f, 0.6666666865348815918f), 1.0f)",
},
ImgData{
.name = "2D null",
// float null works
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "OpImageSampleProjDrefExplicitLod %float %sampled_image %coords4 %float_1 "
"Lod %float_null",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = R"(
%6:f32 = textureSampleCompareLevel %5, %4, vec2<f32>(0.3333333432674407959f, 0.6666666865348815918f), 1.0f)",
}));
// 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:void = textureStore %3, vec2<i32>(1i, 2i), vec4<f32>(1.0f))"},
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:void = textureStore %3, vec2<i32>(1i, 2i), vec4<f32>(1.0f, 2.0f, 0.0f, 0.0f))",
},
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:void = textureStore %3, vec2<i32>(1i, 2i), vec4<f32>(1.0f, 2.0f, 3.0f, 0.0f))",
},
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:void = textureStore %3, vec2<i32>(1i, 2i), vec4<f32>(1.0f, 2.0f, 0.0f, 0.0f))",
},
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:void = textureStore %3, vec2<i32>(1i, 2i), vec4<f32>(1.0f, 2.0f, 3.0f, 0.0f))",
},
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:void = textureStore %3, vec2<i32>(1i, 2i), vec4<u32>(1u))",
},
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:void = textureStore %3, vec2<i32>(1i, 2i), vec4<u32>(1u, 2u, 0u, 0u))",
},
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:void = textureStore %3, vec2<i32>(1i, 2i), vec4<u32>(1u, 2u, 3u, 0u))",
},
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:void = textureStore %3, vec2<i32>(1i, 2i), vec4<u32>(1u, 2u, 0u, 0u))",
},
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:void = textureStore %3, vec2<i32>(1i, 2i), vec4<u32>(1u, 2u, 3u, 0u))",
},
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:void = textureStore %3, vec2<i32>(1i, 2i), vec4<i32>(1i))",
},
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:void = textureStore %3, vec2<i32>(1i, 2i), vec4<i32>(1i, 2i, 0i, 0i))",
},
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:void = textureStore %3, vec2<i32>(1i, 2i), vec4<i32>(1i, 2i, 3i, 0i))",
},
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:void = textureStore %3, vec2<i32>(1i, 2i), vec4<i32>(1i, 2i, 0i, 0i))",
},
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:void = textureStore %3, vec2<i32>(1i, 2i), vec4<i32>(1i, 2i, 3i, 0i))",
},
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:void = textureStore %3, vec2<i32>(1i, 2i), vec4<u32>(1u, 2u, 0u, 0u))",
},
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:void = textureStore %3, vec2<i32>(1i, 2i), vec4<i32>(1i, 2i, 0i, 0i))",
}));
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:vec4<f32> = textureLoad %3, vec2<i32>(1i, 2i), 1i)",
}));
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:vec4<f32> = textureSample %5, %4, 1.0f)",
},
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:vec4<f32> = textureSample %5, %4, 1.0f)",
},
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:vec4<f32> = textureSample %5, %4, 1.0f)",
},
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:vec4<f32> = textureSample %5, %4, vec2<f32>(1.0f, 2.0f))",
},
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:vec4<f32> = textureSample %5, %4, vec2<f32>(1.0f, 2.0f))",
},
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:vec4<f32> = textureSample %5, %4, vec2<f32>(1.0f, 2.0f), 3i)",
},
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:vec4<f32> = textureSample %5, %4, vec2<f32>(1.0f, 2.0f), 3i)",
},
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:vec4<f32> = textureSample %5, %4, vec3<f32>(1.0f, 2.0f, 3.0f))",
},
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:vec4<f32> = textureSample %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 = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf1234",
.wgsl_type = "texture_cube_array<f32>",
.wgsl_fn = R"(
%6:vec4<f32> = textureSample %5, %4, vec3<f32>(1.0f, 2.0f, 3.0f), 4i)",
},
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(
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 = "texture_1d<f32>",
.wgsl_fn = R"(
%6:vec4<f32> = textureSample %5, %4, 1.0f)",
},
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 = "texture_1d<f32>",
.wgsl_fn = R"(
%6:vec4<f32> = textureSampleLevel %5, %4, 1.0f, 1.0f)",
},
ImgData{
.name = "2D vec2",
.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 vec2 lod",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn =
"%result = OpImageSampleExplicitLod %v4float %sampled_image %vf12 Lod %float_1",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = R"(
%6:vec4<f32> = textureSampleLevel %5, %4, vec2<f32>(1.0f, 2.0f), 1.0f)",
}));
// 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(
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 %vf1234",
.wgsl_type = "texture_2d_array<f32>",
.wgsl_fn = R"(
%6:vec4<f32> = textureSample %5, %4, vec2<f32>(1.0f, 2.0f), 3i)",
},
ImgData{
.name = "2D lod",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn =
"%result = OpImageSampleExplicitLod %v4float %sampled_image %vf1234 Lod %float_1",
.wgsl_type = "texture_2d_array<f32>",
.wgsl_fn = R"(
%6:vec4<f32> = textureSampleLevel %5, %4, vec2<f32>(1.0f, 2.0f), 3i, 1.0f)",
}));
// 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(
SpirvReaderTest_PreserveFloatCoords_Arrayed,
SamplerComparisonTest,
::testing::Values(
ImgData{
.name = "2D Depth",
.spirv_type = "%float 2D 1 1 0 1 Unknown",
.spirv_fn = "OpImageSampleDrefImplicitLod %float %sampled_image %coords3 %float_1",
.wgsl_type = "texture_depth_2d_array",
.wgsl_fn = R"(
%6:f32 = textureSampleCompare %5, %4, vec2<f32>(1.0f, 2.0f), 3i, 1.0f)",
},
ImgData{
.name = "2D depth lod",
.spirv_type = "%float 2D 1 1 0 1 Unknown",
.spirv_fn = "OpImageSampleDrefExplicitLod %float %sampled_image %coords3 %float_1 Lod "
"%float_0",
.wgsl_type = "texture_depth_2d_array",
.wgsl_fn = R"(
%6:f32 = textureSampleCompareLevel %5, %4, vec2<f32>(1.0f, 2.0f), 3i, 1.0f)",
}));
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:vec4<f32> = textureLoad %3, vec2<i32>(1i, 2i), 3i, 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:vec4<f32> = textureLoad %3, vec2<u32>(1u, 2u), 3i, 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:vec4<f32> = textureLoad %3, vec2<i32>(1i, 2i), 3i)",
},
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:void = textureStore %3, vec2<i32>(1i, 2i), 3i, 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:vec4<f32> = textureLoad %3, vec2<u32>(1u, 2u), 3i)",
},
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:void = textureStore %3, vec2<u32>(1u, 2u), 3i, 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, 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"(
$B1: { # root
%Output2Texture2D:ptr<handle, texture_storage_2d<rg32float, read_write>, read> = var undef @binding_point(0, 0)
}
%main = @compute @workgroup_size(8u, 8u, 1u) func():void {
$B2: {
%3:ptr<function, vec2<f32>, read_write> = var undef
loop [i: $B3, b: $B4, c: $B5] { # loop_1
$B3: { # initializer
%4:ptr<function, vec2<f32>, read_write> = var undef
store %4, vec2<f32>(0.0f)
%5:ptr<function, i32, read_write> = var undef
store %5, 0i
next_iteration # -> $B4
}
$B4: { # body
%6:i32 = load %5
%7:vec2<f32> = load %4
%8:bool = lt %6, 2i
if %8 [t: $B6, f: $B7] { # if_1
$B6: { # true
continue # -> $B5
}
$B7: { # false
store %3, %7
exit_loop # loop_1
}
}
unreachable
}
$B5: { # continuing
%9:i32 = add %6, 1i
store %4, vec2<f32>(1.0f)
store %5, %9
next_iteration # -> $B4
}
}
%10:vec2<f32> = load %3
%11:texture_storage_2d<rg32float, read_write> = load %Output2Texture2D
%12:vec4<f32> = construct %10, vec2<f32>(0.0f)
%13:void = textureStore %11, vec2<u32>(1u), %12
ret
}
)");
}
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
}
}
)");
}
TEST_F(SpirvReaderTest, Let_Sampler) {
EXPECT_IR(R"(
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
OpName %main "main"
OpName %t "t"
OpName %s "s"
OpDecorate %t Binding 0
OpDecorate %t DescriptorSet 0
OpDecorate %s Binding 1
OpDecorate %s DescriptorSet 0
%void = OpTypeVoid
%3 = OpTypeFunction %void
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%10 = OpTypeImage %float 2D 0 0 0 1 Unknown
%_ptr_UniformConstant_10 = OpTypePointer UniformConstant %10
%t = OpVariable %_ptr_UniformConstant_10 UniformConstant
%14 = OpTypeSampler
%_ptr_UniformConstant_14 = OpTypePointer UniformConstant %14
%s = OpVariable %_ptr_UniformConstant_14 UniformConstant
%18 = OpTypeSampledImage %10
%v2float = OpTypeVector %float 2
%float_0 = OpConstant %float 0
%22 = OpConstantComposite %v2float %float_0 %float_0
%main = OpFunction %void None %3
%5 = OpLabel
%13 = OpLoad %10 %t
%17 = OpLoad %14 %s
%100 = OpCopyObject %14 %17
%19 = OpSampledImage %18 %13 %100
%23 = OpImageSampleExplicitLod %v4float %19 %22 Lod %float_0
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%t:ptr<handle, texture_2d<f32>, read> = var undef @binding_point(0, 0)
%s:ptr<handle, sampler, read> = var undef @binding_point(0, 1)
}
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B2: {
%4:texture_2d<f32> = load %t
%5:sampler = load %s
%6:vec4<f32> = textureSampleLevel %4, %5, vec2<f32>(0.0f), 0.0f
ret
}
}
)");
}
TEST_F(SpirvReaderTest, Let_SampledImage) {
EXPECT_IR(R"(
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
OpName %main "main"
OpName %t "t"
OpName %s "s"
OpDecorate %t Binding 0
OpDecorate %t DescriptorSet 0
OpDecorate %s Binding 1
OpDecorate %s DescriptorSet 0
%void = OpTypeVoid
%3 = OpTypeFunction %void
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%10 = OpTypeImage %float 2D 0 0 0 1 Unknown
%_ptr_UniformConstant_10 = OpTypePointer UniformConstant %10
%t = OpVariable %_ptr_UniformConstant_10 UniformConstant
%14 = OpTypeSampler
%_ptr_UniformConstant_14 = OpTypePointer UniformConstant %14
%s = OpVariable %_ptr_UniformConstant_14 UniformConstant
%18 = OpTypeSampledImage %10
%v2float = OpTypeVector %float 2
%float_0 = OpConstant %float 0
%22 = OpConstantComposite %v2float %float_0 %float_0
%main = OpFunction %void None %3
%5 = OpLabel
%13 = OpLoad %10 %t
%17 = OpLoad %14 %s
%19 = OpSampledImage %18 %13 %17
%100 = OpCopyObject %18 %19
%23 = OpImageSampleExplicitLod %v4float %100 %22 Lod %float_0
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%t:ptr<handle, texture_2d<f32>, read> = var undef @binding_point(0, 0)
%s:ptr<handle, sampler, read> = var undef @binding_point(0, 1)
}
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B2: {
%4:texture_2d<f32> = load %t
%5:sampler = load %s
%6:vec4<f32> = textureSampleLevel %4, %5, vec2<f32>(0.0f), 0.0f
ret
}
}
)");
}
// https://crbug.com/430358345
TEST_F(SpirvReaderTest, Image_UserCall_Params) {
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %main "main" %main_position_Input
OpExecutionMode %main OriginUpperLeft
OpName %sampler0 "sampler0"
OpName %texture0 "texture0"
OpName %main_position_Input "main_position_Input"
OpName %foo "foo"
OpName %t "t"
OpName %s "s"
OpName %main "main"
OpDecorate %sampler0 DescriptorSet 0
OpDecorate %sampler0 Binding 0
OpDecorate %texture0 DescriptorSet 0
OpDecorate %texture0 Binding 1
OpDecorate %main_position_Input BuiltIn FragCoord
%3 = OpTypeSampler
%_ptr_UniformConstant_3 = OpTypePointer UniformConstant %3
%sampler0 = OpVariable %_ptr_UniformConstant_3 UniformConstant
%float = OpTypeFloat 32
%6 = OpTypeImage %float 2D 0 0 0 1 Unknown
%_ptr_UniformConstant_6 = OpTypePointer UniformConstant %6
%texture0 = OpVariable %_ptr_UniformConstant_6 UniformConstant
%v4float = OpTypeVector %float 4
%_ptr_Input_v4float = OpTypePointer Input %v4float
%main_position_Input = OpVariable %_ptr_Input_v4float Input
%17 = OpTypeFunction %v4float %6 %3
%v2float = OpTypeVector %float 2
%float_2 = OpConstant %float 2
%22 = OpConstantComposite %v2float %float_2 %float_2
%25 = OpTypeSampledImage %6
%29 = OpTypeFunction %v4float %v4float
%void = OpTypeVoid
%36 = OpTypeFunction %void
%foo = OpFunction %v4float None %17
%t = OpFunctionParameter %6
%s = OpFunctionParameter %3
%18 = OpLabel
%24 = OpSampledImage %25 %t %s
%26 = OpImageSampleImplicitLod %v4float %24 %22 None
OpReturnValue %26
OpFunctionEnd
%main = OpFunction %void None %36
%37 = OpLabel
%31 = OpLoad %6 %texture0 None
%32 = OpLoad %3 %sampler0 None
%33 = OpFunctionCall %v4float %foo %31 %32
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%sampler0:ptr<handle, sampler, read> = var undef @binding_point(0, 0)
%texture0:ptr<handle, texture_2d<f32>, read> = var undef @binding_point(0, 1)
}
%main = @fragment func(%main_position_Input:vec4<f32> [@position]):void {
$B2: {
%5:texture_2d<f32> = load %texture0
%6:sampler = load %sampler0
%7:vec4<f32> = call %foo, %5, %6
ret
}
}
%foo = func(%t:texture_2d<f32>, %s:sampler):vec4<f32> {
$B3: {
%11:vec4<f32> = textureSample %t, %s, vec2<f32>(2.0f)
ret %11
}
}
)");
}
// As above, but without a name attached to the function (crbug.com/438719644).
TEST_F(SpirvReaderTest, Image_UserCall_Params_NoName) {
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %main "main" %main_position_Input
OpExecutionMode %main OriginUpperLeft
OpDecorate %sampler0 DescriptorSet 0
OpDecorate %sampler0 Binding 0
OpDecorate %texture0 DescriptorSet 0
OpDecorate %texture0 Binding 1
OpDecorate %main_position_Input BuiltIn FragCoord
%3 = OpTypeSampler
%_ptr_UniformConstant_3 = OpTypePointer UniformConstant %3
%sampler0 = OpVariable %_ptr_UniformConstant_3 UniformConstant
%float = OpTypeFloat 32
%6 = OpTypeImage %float 2D 0 0 0 1 Unknown
%_ptr_UniformConstant_6 = OpTypePointer UniformConstant %6
%texture0 = OpVariable %_ptr_UniformConstant_6 UniformConstant
%v4float = OpTypeVector %float 4
%_ptr_Input_v4float = OpTypePointer Input %v4float
%main_position_Input = OpVariable %_ptr_Input_v4float Input
%17 = OpTypeFunction %v4float %6 %3
%v2float = OpTypeVector %float 2
%float_2 = OpConstant %float 2
%22 = OpConstantComposite %v2float %float_2 %float_2
%25 = OpTypeSampledImage %6
%29 = OpTypeFunction %v4float %v4float
%void = OpTypeVoid
%36 = OpTypeFunction %void
%foo = OpFunction %v4float None %17
%t = OpFunctionParameter %6
%s = OpFunctionParameter %3
%18 = OpLabel
%24 = OpSampledImage %25 %t %s
%26 = OpImageSampleImplicitLod %v4float %24 %22 None
OpReturnValue %26
OpFunctionEnd
%main = OpFunction %void None %36
%37 = OpLabel
%31 = OpLoad %6 %texture0 None
%32 = OpLoad %3 %sampler0 None
%33 = OpFunctionCall %v4float %foo %31 %32
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%1:ptr<handle, sampler, read> = var undef @binding_point(0, 0)
%2:ptr<handle, texture_2d<f32>, read> = var undef @binding_point(0, 1)
}
%main = @fragment func(%4:vec4<f32> [@position]):void {
$B2: {
%5:texture_2d<f32> = load %2
%6:sampler = load %1
%7:vec4<f32> = call %8, %5, %6
ret
}
}
%8 = func(%9:texture_2d<f32>, %10:sampler):vec4<f32> {
$B3: {
%11:vec4<f32> = textureSample %9, %10, vec2<f32>(2.0f)
ret %11
}
}
)");
}
// https://crbug.com/441874372
TEST_F(SpirvReaderTest, TexSamplerAsParams) {
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %main "main"
OpExecutionMode %main OriginUpperLeft
OpName %main "main"
OpName %helper "helper"
OpName %t "t"
OpName %s "s"
OpName %tex "tex"
OpName %samp "samp"
OpDecorate %tex DescriptorSet 1
OpDecorate %tex Binding 1
OpDecorate %samp DescriptorSet 2
OpDecorate %samp Binding 1
%void = OpTypeVoid
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%v4float = OpTypeVector %float 4
%texture = OpTypeImage %float 2D 0 0 0 1 Unknown
%ptr_tex = OpTypePointer UniformConstant %texture
%sampler = OpTypeSampler
%ptr_sampler = OpTypePointer UniformConstant %sampler
%26 = OpTypeSampledImage %texture
%float_1 = OpConstant %float 1
%23 = OpConstantComposite %v2float %float_1 %float_1
%3 = OpTypeFunction %void
%12 = OpTypeFunction %void %ptr_tex %ptr_sampler
%tex = OpVariable %ptr_tex UniformConstant
%samp = OpVariable %ptr_sampler UniformConstant
%main = OpFunction %void None %3
%5 = OpLabel
%39 = OpFunctionCall %void %helper %tex %samp
OpReturn
OpFunctionEnd
%helper = OpFunction %void None %12
%t = OpFunctionParameter %ptr_tex
%s = OpFunctionParameter %ptr_sampler
%16 = OpLabel
%24 = OpLoad %texture %t
%25 = OpLoad %sampler %s
%27 = OpSampledImage %26 %24 %25
%31 = OpImageSampleExplicitLod %v4float %27 %23 Lod %float_1
%50 = OpCopyObject %ptr_tex %t
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%tex:ptr<handle, texture_2d<f32>, read> = var undef @binding_point(1, 1)
%samp:ptr<handle, sampler, read> = var undef @binding_point(2, 1)
}
%main = @fragment func():void {
$B2: {
%4:texture_2d<f32> = load %tex
%5:sampler = load %samp
%6:void = call %helper, %4, %5
ret
}
}
%helper = func(%t:texture_2d<f32>, %s:sampler):void {
$B3: {
%10:vec4<f32> = textureSampleLevel %t, %s, vec2<f32>(1.0f), 1.0f
ret
}
}
)");
}
TEST_F(SpirvReaderTest, TexSamplerAsParams_Dref) {
EXPECT_IR(R"(
OpCapability Shader
OpCapability ImageQuery
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %main "main"
OpExecutionMode %main OriginUpperLeft
OpName %main "main"
OpDecorate %var_img Binding 12
OpDecorate %var_img DescriptorSet 0
OpDecorate %var_samp Binding 13
OpDecorate %var_samp DescriptorSet 0
%void = OpTypeVoid
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%img = OpTypeImage %float 2D 1 1 0 1 Unknown
%ptr_img = OpTypePointer UniformConstant %img
%samp = OpTypeSampler
%ptr_samp = OpTypePointer UniformConstant %samp
%sampled_img = OpTypeSampledImage %img
%1282 = OpTypeFunction %void
%2564 = OpTypeFunction %float %ptr_img %ptr_samp
%ptr_v4float = OpTypePointer Function %v4float
%ptr_float = OpTypePointer Function %float
%float_0 = OpConstant %float 0
%v4float_0 = OpConstantComposite %v4float %float_0 %float_0 %float_0 %float_0
%var_samp = OpVariable %ptr_samp UniformConstant
%var_img = OpVariable %ptr_img UniformConstant
%main = OpFunction %void None %1282
%12243 = OpLabel
%param = OpVariable %ptr_v4float Function
%param_0 = OpVariable %ptr_float Function
%22971 = OpFunctionCall %float %bounce %var_img %var_samp
OpReturn
OpFunctionEnd
%bounce = OpFunction %float None %2564
%p1 = OpFunctionParameter %ptr_img
%p2 = OpFunctionParameter %ptr_samp
%90 = OpLabel
%80 = OpFunctionCall %float %shadow_fn %p1 %p2
OpReturnValue %80
OpFunctionEnd
%shadow_fn = OpFunction %float None %2564
%param_img = OpFunctionParameter %ptr_img
%param_samp = OpFunctionParameter %ptr_samp
%5000 = OpLabel
%6208 = OpLoad %img %param_img
%25104 = OpLoad %samp %param_samp
%8422 = OpSampledImage %sampled_img %6208 %25104
%7977 = OpImageSampleDrefImplicitLod %float %8422 %v4float_0 %float_0
OpReturnValue %float_0
OpFunctionEnd
)",
R"(
$B1: { # root
%1:ptr<handle, sampler_comparison, read> = var undef @binding_point(0, 13)
%2:ptr<handle, texture_depth_2d_array, read> = var undef @binding_point(0, 12)
}
%main = @fragment func():void {
$B2: {
%4:ptr<function, vec4<f32>, read_write> = var undef
%5:ptr<function, f32, read_write> = var undef
%6:texture_depth_2d_array = load %2
%7:sampler_comparison = load %1
%8:f32 = call %9, %6, %7
ret
}
}
%9 = func(%10:texture_depth_2d_array, %11:sampler_comparison):f32 {
$B3: {
%12:f32 = call %13, %10, %11
ret %12
}
}
%13 = func(%14:texture_depth_2d_array, %15:sampler_comparison):f32 {
$B4: {
%16:f32 = textureSampleCompare %14, %15, vec2<f32>(0.0f), 0i, 0.0f
ret 0.0f
}
}
)");
}
} // namespace
} // namespace tint::spirv::reader