Google Git
Sign in
dawn/dawn/76d2ddd42e9db254dcd263a397790db9c19e8127/./src/tint/lang/spirv/reader/texture_test.cc
blob: 3f91ca144ecab66ac22db2b02756a9e4c568ac79 [file] [log] [blame]
// Copyright 2025 The Dawn & Tint Authors
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are met:
//
// 1. Redistributions of source code must retain the above copyright notice, this
// list of conditions and the following disclaimer.
//
// 2. Redistributions in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// 3. Neither the name of the copyright holder nor the names of its
// contributors may be used to endorse or promote products derived from
// this software without specific prior written permission.
//
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
#include "src/tint/lang/spirv/reader/helper_test.h"
namespace tint::spirv::reader {
namespace {
TEST_F(SpirvReaderTest, Handle_Sampler) {
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical Simple
OpEntryPoint Fragment %main "main"
OpExecutionMode %main OriginUpperLeft
OpName %10 "wg"
OpDecorate %10 DescriptorSet 0
OpDecorate %10 Binding 0
%sampler = OpTypeSampler
%ptr = OpTypePointer UniformConstant %sampler
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%10 = OpVariable %ptr UniformConstant
%main = OpFunction %void None %voidfn
%entry = OpLabel
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%wg:ptr<handle, sampler, read> = var undef @binding_point(0, 0)
}
%main = @fragment func():void {
$B2: {
ret
}
}
)");
}
TEST_F(SpirvReaderTest, DISABLED_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, DISABLED_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_1d<f32>, read> = var undef @binding_point(0, 0)
}
%main = @fragment func():void {
$B2: {
ret
}
}
)");
}
TEST_F(SpirvReaderTest, DISABLED_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>, write> = var undef @binding_point(0, 0)
}
%main = @fragment func():void {
$B2: {
ret
}
}
)");
}
TEST_F(SpirvReaderTest, DISABLED_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
%im = OpLoad %tex %10
%result = OpImageQuerySamples %uint %im
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
}
}
)");
}
struct ImgData {
std::string name;
std::string spirv_type;
std::string spirv_fn;
std::string wgsl_type;
std::string wgsl_fn;
};
[[maybe_unused]] std::ostream& operator<<(std::ostream& out, const ImgData& d) {
out << d.name;
return out;
}
using SamplerTest = SpirvReaderTestWithParam<ImgData>;
TEST_P(SamplerTest, Handle) {
auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpCapability Sampled1D
OpMemoryModel Logical Simple
OpEntryPoint Fragment %main "main"
OpExecutionMode %main OriginUpperLeft
OpName %10 "wg"
OpDecorate %var_sampler DescriptorSet 0
OpDecorate %var_sampler Binding 0
OpDecorate %10 DescriptorSet 2
OpDecorate %10 Binding 0
%int = OpTypeInt 32 1
%uint = OpTypeInt 32 0
%float = OpTypeFloat 32
%v2int = OpTypeVector %int 2
%v2uint = OpTypeVector %uint 2
%v2float = OpTypeVector %float 2
%v3float = OpTypeVector %float 3
%v4float = OpTypeVector %float 4
%tex = OpTypeImage )" +
params.spirv_type + R"(
%ptr_tex = OpTypePointer UniformConstant %tex
%sampled_img = OpTypeSampledImage %tex
%sampler = OpTypeSampler
%ptr_sampler = OpTypePointer UniformConstant %sampler
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%int_1 = OpConstant %int 1
%float_null = OpConstantNull %float
%float_1 = OpConstant %float 1
%float_2 = OpConstant %float 2
%float_3 = OpConstant %float 3
%float_4 = OpConstant %float 4
%float_5 = OpConstant %float 5
%coords2 = OpConstantComposite %v2float %float_1 %float_2
%coords3 = OpConstantComposite %v3float %float_1 %float_2 %float_3
%coords4 = OpConstantComposite %v4float %float_1 %float_2 %float_3 %float_4
%vf12 = OpConstantComposite %v2float %float_1 %float_2
%vf21 = OpConstantComposite %v2float %float_2 %float_1
%int_10 = OpConstant %int 10
%int_11 = OpConstant %int 11
%offset2i = OpConstantComposite %v2int %int_10 %int_11
%uint_20 = OpConstant %uint 20
%uint_21 = OpConstant %uint 21
%offset2u = OpConstantComposite %v2uint %uint_20 %uint_21
%depth = OpConstant %float 1
%var_sampler = OpVariable %ptr_sampler UniformConstant
%10 = OpVariable %ptr_tex UniformConstant
%main = OpFunction %void None %voidfn
%entry = OpLabel
%sam = OpLoad %sampler %var_sampler
%im = OpLoad %tex %10
%sampled_image = OpSampledImage %sampled_img %im %sam
%result = )" + params.spirv_fn +
R"(
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%1:ptr<handle, sampler, read> = var undef @binding_point(0, 0)
%wg:ptr<handle, )" +
params.wgsl_type + R"(<f32>, read> = var undef @binding_point(2, 0)
}
%main = @fragment func():void {
$B2: {
%4:)" + params.wgsl_type +
R"(<f32> = load %wg
%5:sampler = load %1
%6:vec4<f32> = )" +
params.wgsl_fn + R"(
ret
}
}
)");
}
INSTANTIATE_TEST_SUITE_P(
DISABLED_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 = "textureGather 1i, %4, %5, 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 = "textureGather 1i, %4, %5, vec2<f32>(1.0f, 2.0f), vec2<i32>(10i, 11i)",
},
ImgData{
.name = "2D ConstOffset Unsigned",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn =
"OpImageGather %v4float %sampled_image %coords2 %int_1 ConstOffset %offset2u",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = "textureGather 1i, %4, %5, vec2<f32>(1.0f, 2.0f), vec2<u32>(20u, 21u)",
},
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 = "textureGather 1i, %4, %5, 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 = "textureGather 1i, %4, %5, vec2<f32>(1.0f, 2.0f), 3i, vec2<i32>(10i, 11i)",
},
ImgData{
.name = "2D Array ConstOffset Unsigned",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn =
"OpImageGather %v4float %sampled_image %coords3 %int_1 ConstOffset %offset2u",
.wgsl_type = "texture_2d_array<f32>",
.wgsl_fn = "textureGather 1i, %4, %5, vec2<f32>(1.0f, 2.0f), 3i, vec2<u32>(20u, 21u)",
},
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 = "textureGather 1i, %4, %5, 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 = "textureGather 1i, %4, %5, vec3<f32>(1.0f, 2.0f, 3.0f), 4i",
},
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 = "textureGather 1i, %4, %5, 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 = "textureGather 1i, %4, %5, vec2<f32>(1.0f, 2.0f), vec2<i32>(10i, 11i)",
},
ImgData{
.name = "2D Depth ConstOffset Unigned",
.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 = "textureGather 1i, %4, %5, vec2<f32>(1.0f, 2.0f), vec2<u32>(20u, 21u)",
},
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 = "textureGather 1i, %4, %5, 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 = "textureGather 1i, %4, %5, vec2<f32>(1.0f, 2.0f), 3i, vec2<i32>(10i, 11i)",
},
ImgData{
.name = "2D Depth Array ConstOffset Unigned",
.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 = "textureGather 1i, %4, %5, vec2<f32>(1.0f, 2.0f), 3i, vec2<u32>(20u, 21u)",
},
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 = "textureGather 1i, %4, %5, vec3<f32>(1.0f, 2.0f, 3.0f)",
},
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_2d_array",
.wgsl_fn = "textureGather 1i, %4, %5, vec3<f32>(1.0f, 2.0f, 3.0f), 4i",
}));
INSTANTIATE_TEST_SUITE_P(
DISABLED_SpirvReaderTest_ImageDrefGather,
SamplerTest,
::testing::Values(
ImgData{
.name = "2d Depth",
.spirv_type = "%float 2d 1 0 0 1 Unknown",
.spirv_fn = "OpImageDrefGather %v4float %sampled_image %coords2 %depth",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = "textureGatherCompare %4, %5, vec2<f32>(1.0f, 2.0f), 1.0f",
},
ImgData{
.name = "2d Depth ConstOffset Signed",
.spirv_type = "%float 2d 1 0 0 1 Unknown",
.spirv_fn =
"OpImageDrefGather %v4float %sampled_image %coords2 %depth ConstOffset %offset2i",
.wgsl_type = "texture_depth_2d",
.wgsl_fn =
"textureGatherCompare %4, %5, vec2<f32>(1.0f, 2.0f), 1.0f, vec2<i32>(10i, 11i)",
},
ImgData{
.name = "2d Depth ConstOffset Unsigned",
.spirv_type = "%float 2d 1 0 0 1 Unknown",
.spirv_fn =
"OpImageDrefGather %v4float %sampled_image %coords2 %depth ConstOfset %offset2u",
.wgsl_type = "texture_depth_2d",
.wgsl_fn =
"textureGatherCompare %4, %5, vec2<f32>(1.0f, 2.0f), 1.0f, vec2<u32>(20u, 21u)",
},
ImgData{
.name = "2d Depth Array",
.spirv_type = "%float 2d 1 1 0 1 Unknown",
.spirv_fn = "OpImageDrefGather %v4float %sampled_image %coords3 %depth",
.wgsl_type = "texture_depth_2d_array",
.wgsl_fn = "textureGatherCompare %4, %5, vec2<f32>(1.0f, 2.0f), 3i, 1.0f",
},
ImgData{
.name = "2d Depth Array ConstOffset Signed",
.spirv_type = "%float 2d 1 1 0 1 Unknown",
.spirv_fn =
"OpImageDrefGather %v4float %sampled_image %coords3 %depth ConstOffset %offset2i",
.wgsl_type = "texture_depth_2_arrayd",
.wgsl_fn =
"textureGatherCompare %4, %5, vec2<f32>(1.0f, 2.0f), 3i, 1.0f, vec2<i32>(10i, 11i)",
},
ImgData{
.name = "2d Depth Array ConstOffset Unsigned",
.spirv_type = "%float 2d 1 1 0 1 Unknown",
.spirv_fn =
"OpImageDrefGather %v4float %sampled_image %coords3 %depth ConstOfset %offset2u",
.wgsl_type = "texture_depth_2d_array",
.wgsl_fn =
"textureGatherCompare %4, %5, vec2<f32>(1.0f, 2.0f), 3i, 1.0f, vec2<u32>(20u, 21u)",
},
ImgData{
.name = "Cube Depth",
.spirv_type = "%float Cube 1 0 0 1 Unknown",
.spirv_fn = "OpImageDrefGather %v4float %sampled_image %coords3 %depth",
.wgsl_type = "texture_depth_cube",
.wgsl_fn = "textureGatherCompare %4, %5, vec3<f32>(1.0f, 2.0f, 3.0f), 1.0f",
},
ImgData{
.name = "Cube Depth Array",
.spirv_type = "%float Cube 1 1 0 1 Unknown",
.spirv_fn = "OpImageDrefGather %v4float %sampled_image %coords4 %depth",
.wgsl_type = "texture_depth_cube_array",
.wgsl_fn = "textureGatherCompare %4, %5, vec3<f32>(1.0f, 2.0f, 3.0f), 4i, 1.0f",
}));
INSTANTIATE_TEST_SUITE_P(
DISABLED_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 = "textureSample %4, %5, 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 = "textureSample %4, %5, vec2<f32>(1.0f, 2.0f), vec2<i32>(10i, 11i)",
},
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 = "textureSample %4, %5, 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 = "textureSample %4, %5, vec2<f32>(1.0f, 2.0f), 3i, vec2<i32>(10i, 11i)",
},
ImgData{
.name = "2D Bias",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "OpImageSampleImplicitLod %v4float %sampled_image %coords2 Bias %float_5",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = "textureSampleBias %4, %5, vec2<f32>(1.0f, 2.0f), 5.0f",
},
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 = "textureSampleBias %4, %5, vec2<f32>(1.0f, 2.0f), 3i, 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 = "textureSampleBias %4, %5, vec2<f32>(1.0f, 2.0f), 5.0f, vec2<i32>(10i, 11i)",
},
ImgData{
.name = "2D Bias ConstOffset Unigned",
.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 = "textureSampleBias %4, %5, vec2<f32>(1.0f, 2.0f), 5.0f, vec2<u32>(20i, 21i)",
},
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 =
"textureSampleBias %4, %5, vec2<f32>(1.0f, 2.0f), 3i, 5.0f, vec2<i32>(10i, 11i)",
},
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 = "textureSample %4, %5, vec2<f32>(1.0f, 2.0f)",
},
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 = "textureSample %4, %5, vec2<f32>(1.0f, 2.0f, 3.0f)",
},
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 = "textureSample %4, %5, vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f)",
},
ImgData{
.name = "2D ConstOffset",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "OpImageSampleProjImplicitLod %v4float %sampled_image %coords3 ConstOffset "
"%offsets2d",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = "textureSample %4, %5, vec3<f32>(1.0f, 2.0f, 3.0f), vec2<i32>(10i, 11i)",
},
ImgData{
.name = "2D Bias",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn =
"OpImageSampleProjImplicitLod %v4float %sampled_image %coords3 Bias %float_5",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = "textureSampleBias %4, %5, vec3<f32>(1.0, 2.0, 3.0), 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 =
"textureSampleBias %4, %5, vec3<f32>(1.0f, 2.0f, 3.0f), 5.0f, vec2<i32>(10i, 11i)",
},
ImgData{
.name = "2D Bias ConstOffset Unsigned",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "OpImageSampleProjImplicitLod %v4float %sampled_image %coords3 "
"Bias|ConstOffset %float_5 %offset2u",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn =
"textureSampleBias %4, %5, vec3<f32>(1.0f, 2.0f, 3.0f), 5.0f, vec2<u32>(20u, 21u)",
},
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 = "textureSample %4, %5, vec3<f32>(1.0f, 2.0f, 3.0f)",
}
));
INSTANTIATE_TEST_SUITE_P(
DISABLED_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 = "textureSampleLevel %4, %5, 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 = "textureSampleLevel %4, %5, 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 =
"textureSampleLevel %4, %5, vec2<f32>(1.0f, 2.0f), 0.0f, vec2<i32>(10i, 11i)",
},
ImgData{
.name = "2D lod ConstOffset unsigned",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "OpImageSampleExplicitLod %v4float %sampled_image %coords2 Lod|ConstOffset "
"%float_null %offset2u",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn =
"textureSampleLevel %4, %5, vec2<f32>(1.0f, 2.0f), 0.0f, vec2<u32>(20u, 21u)",
},
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 =
"textureSampleLevel %4, %5, vec2<f32>(1.0f, 2.0f), 3i, 0.0f, vec2<i32>(10i, 11i)",
},
ImgData{
.name = "2D Grad",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn =
"OpImageSampleExplicitLod %v4float %sampled_image %coords12 Grad %vf12 %vf21",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = "textureSampleGrad %4, %5, vec2<f32>(1.0f, 2.0f), vec2<f32>(1.0f, 2.0f), "
"vec2<f32>(2.0f, 1.0)",
},
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 = "textureSampleGrad %4, %5, 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 = "textureSampleGrad %4, %5, vec2<f32>(1.0f, 2.0f), vec2<f32>(1.0f, 2.0f), "
"vec2<f32>(2.0f, 1.0), vec2<i32>(10i, 11i)",
},
ImgData{
.name = "2D Grad ConstOffset Unsigned",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "OpImageSampleExplicitLod %v4float %sampled_image %coords2 "
"Grad|ConstOffset %vf12 %vf21 %offset2u",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = "textureSampleGrad %4, %5, vec2<f32>(1.0f, 2.0f), vec2<f32>(1.0f, 2.0f), "
"vec2<f32>(2.0f, 1.0f), vec2<u32>(20u, 21u)",
},
ImgData{
.name = "2D Arrayed 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 = "textureSampleGrad %4, %5, vec2<f32>(1.0f, 2.0f), 3i, vec2<f32>(1.0f, "
"2.0f), vec2<f32>(2.0f, 1.0f, vec2<i32>(10i, 11i)",
},
ImgData{
.name = "2D arrayed 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 = "textureSampleGrad %4, %5, vec2<f32>(1.0f, 2.0f), 3i, vec2<f32>(1.0f, 2.0), "
"vec2<f32>(2.0f, 1.0f), vec2<u32>(20u, 21u)",
},
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 = "textureSampleLevel %4, %5, vec2<f23>(1.0f, 2.0f), 1i",
}));
INSTANTIATE_TEST_SUITE_P(
DISABLED_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 = "textureSampleLevel %4, %5, vec3<f32>(1.0f, 2.0f, 3.0f), 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 =
"textureSampleLevel %4, %5, vec3<f32>(1.0f, 2.0f, 3.0f), 1.0f, vec2<i32>(10i, 11i)",
},
ImgData{
.name = "2D Grad",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn =
"OpImageSampleProjExplicitLod %v4float %sampled_image %coords3 Grad %vf12 %vf21",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = "textureSampleGrad %4, %5, vec3<f32>(1.0f, 2.0f, 3.0f), 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 = "textureSampleGrad %4, %5, vec3<f32>(1.0f, 2.0f, 3.0f), vec2<f32>(1.0f, "
"2.0f), vec2<f32>(2.0f, 1.0f), vec2<i32>(10i, 20i)",
}
));
using SamplerComparisonTest = SpirvReaderTestWithParam<ImgData>;
TEST_P(SamplerComparisonTest, Handle) {
auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpCapability Sampled1D
OpMemoryModel Logical Simple
OpEntryPoint Fragment %main "main"
OpExecutionMode %main OriginUpperLeft
OpName %10 "wg"
OpDecorate %var_sampler DescriptorSet 0
OpDecorate %var_sampler Binding 0
OpDecorate %10 DescriptorSet 2
OpDecorate %10 Binding 0
%int = OpTypeInt 32 1
%uint = OpTypeInt 32 0
%float = OpTypeFloat 32
%v2int = OpTypeVector %int 2
%v2uint = OpTypeVector %uint 2
%v2float = OpTypeVector %float 2
%v3float = OpTypeVector %float 3
%v4float = OpTypeVector %float 4
%tex = OpTypeImage )" +
params.spirv_type +
R"(
%ptr_tex = OpTypePointer UniformConstant %tex
%sampled_img = OpTypeSampledImage %tex
%sampler = OpTypeSampler
%ptr_sampler = OpTypePointer UniformConstant %sampler
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%int_1 = OpConstant %int 1
%float_0 = OpConstant %float 0
%float_1 = OpConstant %float 1
%float_2 = OpConstant %float 2
%float_3 = OpConstant %float 3
%float_4 = OpConstant %float 4
%float_5 = OpConstant %float 5
%coords2 = OpConstantComposite %v2float %float_1 %float_2
%coords3 = OpConstantComposite %v3float %float_1 %float_2 %float_3
%coords4 = OpConstantComposite %v4float %float_1 %float_2 %float_3 %float_4
%int_10 = OpConstant %int 10
%int_11 = OpConstant %int 11
%offset2i = OpConstantComposite %v2int %int_10 %int_11
%uint_20 = OpConstant %uint 20
%uint_21 = OpConstant %uint 21
%offset2u = OpConstantComposite %v2uint %uint_20 %uint_21
%depth = OpConstant %float 1
%var_sampler = OpVariable %ptr_sampler UniformConstant
%10 = OpVariable %ptr_tex UniformConstant
%main = OpFunction %void None %voidfn
%entry = OpLabel
%sam = OpLoad %sampler %var_sampler
%im = OpLoad %tex %10
%sampled_image = OpSampledImage %sampled_img %im %sam
%result = )" + params.spirv_fn +
R"(
OpReturn
OpFunctionEnd
)",
R"(
$B1: { # root
%1:ptr<handle, sampler_comparison, read> = var undef @binding_point(0, 0)
%wg:ptr<handle, )" +
params.wgsl_type +
R"(<f32>, read> = var undef @binding_point(2, 0)
}
%main = @fragment func():void {
$B2: {
%4:)" + params.wgsl_type +
R"(<f32> = load %wg
%5:sampler_comparison = load %1
%6:vec4<f32> = )" +
params.wgsl_fn +
R"(
ret
}
}
)");
}
INSTANTIATE_TEST_SUITE_P(
DISABLED_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 = "textureSampleCompare %4, %5, 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 = "textureSampleCompare %4, %5, vec2<f32>(1.0f, 2.0f), 1.0f, "
"vec2<i32>(10i, 11i)",
},
ImgData{
.name = "2D Array",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn = "OpImageSampleDrefImplicitLod %float %sampled_image %coords3 %depth",
.wgsl_type = "texture_depth_2d_array",
.wgsl_fn = "textureSampleCompare %4, %5, 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 = "textureSampleCompare %4, %5, vec2<f32>(1.0f, 2.0f), 3i, 1.0f, "
"vec2<i32>(10i, 11i)",
}));
INSTANTIATE_TEST_SUITE_P(
DISABLED_SpirvReaderTest_ImageSampleDrefExplicitLod,
SamplerComparisonTest,
::testing::Values(
ImgData{
.name = "2D",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "OpImageSampleDrefExplicitLod %float %sampled_image %coords2 "
"%depth Lod %float_0",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = "textureSampleCompareLevel %4, %5, vec2<f32>(1.0f, 2.0f), 1.0f",
},
ImgData{
.name = "2D array",
.spirv_type = "%float 2D 1 1 0 1 Unknown",
.spirv_fn = "OpImageSampleDrefExplicitLod %float %sampled_image %coords3 "
"%depth Lod %float_0",
.wgsl_type = "texture_depth_2d_array",
.wgsl_fn = "textureSampleCompareLevel %4, %5, vec2<f32>(1.0f, 2.0f), 3i, 1.0f",
},
ImgData{
.name = "2D ConstOffset",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "OpImageSampleDrefExplicitLod %float %sampled_image %coords2 %depth "
"Lod|ConstOffset %float_0 %offset2i",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = "textureSampleCompareLevel %4, %5, vec2<f32>(1.0f, 2.0f), 1.0f, "
"vec2<i32>(10i, 11i)",
},
ImgData{
.name = "2D array ConstOffset",
.spirv_type = "%float 2D 1 1 0 1 Unknown",
.spirv_fn = "OpImageSampleDrefExplicitLod %float %sampled_image %coords3 %depth "
"Lod|ConstOffset %float_0 %offset2i",
.wgsl_type = "texture_depth_2d_array",
.wgsl_fn = "textureSampleCompareLevel %4, %5, vec2<f32>(1.0f, 2.0f), 3i, 1.0f, "
"vec2<i32>(10i, 11i)",
},
ImgData{
.name = "Cube",
.spirv_type = "%float Cube 1 0 0 1 Unknown",
.spirv_fn = "OpImageSampleDrefExplicitLod %float %sampled_image %coords3 "
"%depth Lod %float_0",
.wgsl_type = "texture_depth_cube",
.wgsl_fn = "textureSampleCompareLevel %4, %5, vec3<f32>(1.0f, 2.0f, 3.0f), 1.0f",
},
ImgData{
.name = "Cube array",
.spirv_type = "%float Cube 1 1 0 1 Unknown",
.spirv_fn = "OpImageSampleDrefExplicitLod %float %sampled_image %coords4 "
"%depth Lod %float_0",
.wgsl_type = "texture_depth_cube_array",
.wgsl_fn = "textureSampleCompareLevel %4, %5, vec3<f32>(1.0f, 2.0f, 3.0f), "
"4i, 1.0f",
}));
INSTANTIATE_TEST_SUITE_P(
DISABLED_SpirvReaderTest_ImageSampleProjDrefImplicitLod,
SamplerComparisonTest,
::testing::Values(
ImgData{
.name = "2D Depth",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "OpImageSampleProjDrefImplicitLod %float %sampled_image %coords3 %float_1",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = "textureSampleCompare %4, %5, vec3<f32>(1.0f, 2.0f, 3.0f), 1.0f",
},
ImgData{
.name = "2D Depth ConstOffset",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "OpImageSampleProjDrefImplicitLod %float %sampled_image %coords3 %float_1 "
"ConstOffset %offset2i",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = "textureSampleCompare %4, %5, vec3<f32>(1.0f, 2.0f, 3.0f), 1.0f, "
"vec2<i32>(10i, 11i)",
}));
INSTANTIATE_TEST_SUITE_P(
DISABLED_SpirvReaderTest_ImageSampleProjDrefExplicitLod,
SamplerComparisonTest,
::testing::Values(
ImgData{
.name = "2D Depth Lod",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "OpImageSampleProjDrefExplicitLod %float %sampled_image %coords3 %float_1 "
"Lod %float_0",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = "textureSampleCompare %4, %5, vec3<f32>(1.0f, 2.0f, 3.0f), 1.0f, 0.0f",
},
ImgData{
.name = "2D Depth Lod ConstOffset",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "OpImageSampleProjDrefExplicitLod %float %sampled_image %coords3 %float_1 "
"Lod|ConstOffset %float_0 %offset2i",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = "textureSampleCompareLevel %4, %5, vec3<f32>(1.0f, 2.0f, 3.0f), 1.0f, 0.0f, "
"vec2<i32>(10i, 11i)",
}));
// This test shows the use of a sampled image used with both regular
// sampling and depth-reference sampling. The texture is a depth-texture,
// and we use builtins textureSample and textureSampleCompare
TEST_F(SpirvReaderTest, DISABLED_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: {
%4:sampler = load %1
%5:sampler_comparison = load %2
%6:texture_depth_2d = load %wg
%7:vec4<f32> = textureSample %4, %6, vec2<f32>(1.0f, 2.0f)
%8:vec2<f32> = textureSampleCompare %6, %5, 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
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
%v2uint = OpTypeVector %uint 2
%v3uint = OpTypeVector %uint 3
%v4uint = OpTypeVector %uint 4
%v2float = OpTypeVector %float 2
%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
%vf12 = OpConstantComposite %v2float %float_1 %float_2
%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
%1:ptr<handle, sampler_comparison, read> = var undef @binding_point(0, 0)
%wg:ptr<handle, )" +
params.wgsl_type +
R"(<f32>, read> = var undef @binding_point(2, 0)
}
%main = @fragment func():void {
$B2: {
%4:)" + params.wgsl_type +
R"(<f32> = load %wg
%5:sampler_comparison = load %1
%6:vec4<f32> = )" +
params.wgsl_fn +
R"(
ret
}
}
)");
}
INSTANTIATE_TEST_SUITE_P(DISABLED_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, write>;",
.wgsl_fn = "textureStore(x_20, vi12, vf1234);",
}));
// 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(
DISABLED_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, write>",
.wgsl_fn = "textureStore(x_20, vi12, vec4f(f1, 0.0f, 0.0f, 0.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, write>",
.wgsl_fn = "textureStore(x_20, vi12, vec4f(vf12, 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, write>;)",
.wgsl_fn = "textureStore(x_20, vi12, vec4f(vf123, 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, write>",
.wgsl_fn = "textureStore(x_20, vi12, vf1234);",
},
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, write>",
.wgsl_fn = "textureStore(x_20, vi12, vec4f(vf12, 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, write>",
.wgsl_fn = "textureStore(x_20, vi12, vec4f(vf123, 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, write>",
.wgsl_fn = "textureStore(x_20, vi12, vf1234);",
},
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, write>",
.wgsl_fn = "textureStore(x_20, vi12, vf1234);",
}));
// As above, but unsigned integer.
INSTANTIATE_TEST_SUITE_P(DISABLED_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, write>",
.wgsl_fn = "textureStore(x_20, vi12, vec4u(u1, 0u, 0u, 0u));",
},
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, write>",
.wgsl_fn = "textureStore(x_20, vi12, vec4u(vu12, 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, write>",
.wgsl_fn = "textureStore(x_20, vi12, vec4u(vu123, 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, write>",
.wgsl_fn = "textureStore(x_20, vi12, vu1234);",
},
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, write>",
.wgsl_fn = "textureStore(x_20, vi12, vec4u(vu12, 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, write>",
.wgsl_fn = "textureStore(x_20, vi12, vec4u(vu123, 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, write>",
.wgsl_fn = "textureStore(x_20, vi12, vu1234);",
},
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, write>",
.wgsl_fn = "textureStore(x_20, vi12, vu1234);",
}));
// As above, but signed integer.
INSTANTIATE_TEST_SUITE_P(DISABLED_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, write>",
.wgsl_fn = "textureStore(x_20, vi12, vec4i(i1, 0i, 0i, 0i));",
},
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, write>",
.wgsl_fn = "textureStore(x_20, vi12, vec4i(vi12, 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, write>",
.wgsl_fn = "textureStore(x_20, vi12, vec4i(vi123, 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, write>",
.wgsl_fn = "textureStore(x_20, vi12, vi1234);",
},
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, write>",
.wgsl_fn = "textureStore(x_20, vi12, vec4i(vi12, 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, write>",
.wgsl_fn = "textureStore(x_20, vi12, vec4i(vi123, 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, write>",
.wgsl_fn = "textureStore(x_20, vi12, vi1234);",
},
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, write>",
.wgsl_fn = "textureStore(x_20, vi12, vi1234);",
}));
// 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(DISABLED_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, write>",
.wgsl_fn = "textureStore(x_20, vi12, vu1234)",
},
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, write>",
.wgsl_fn = "textureStore(x_20, vi12, vi1234)",
}));
// Show that zeros of the correct integer signedness are
// created when expanding an integer vector.
INSTANTIATE_TEST_SUITE_P(
DISABLED_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, write>",
.wgsl_fn = "textureStore(x_20, vi12, vec4u(vu12, 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, write>",
.wgsl_fn = "textureStore(x_20, vi12, vec4i(vi12, 0i, 0i))",
}));
INSTANTIATE_TEST_SUITE_P(
DISABLED_SpirvReaderTest_ImageFetch_OptionalParams,
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 = "textureLoad(x_20, vi12, 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 = "let x_99 = textureLoad(x_20, vi12, 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 = "let x_99 = vec4f(textureLoad(x_20, vi12, 0i), 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 = "let x_99 = vec4f(textureLoad(x_20, vi12, 3i), 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
INSTANTIATE_TEST_SUITE_P(
DISABLED_SpirvReaderTest_ImageFetch_Depth,
ImageAccessTest,
::testing::Values(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 = "let x_99 = vec4f(textureLoad(x_20, vi12, 0i), 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
INSTANTIATE_TEST_SUITE_P(
DISABLED_SpirvReaderTest_ImageFetch_DepthMultisampled,
ImageAccessTest,
::testing::Values(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 = "let x_99 = vec4f(textureLoad(x_20, vi12, i1), 0.0f, 0.0f, 0.0f)",
}));
// SPIR-V requires a Sample image operand when operating on a multisampled image.
INSTANTIATE_TEST_SUITE_P(DISABLED_SpirvReaderTest_ImageFetch_Multisampled,
ImageAccessTest,
::testing::Values(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 = "let x_99 = textureLoad(x_20, vi12, i1)",
}));
INSTANTIATE_TEST_SUITE_P(DISABLED_SpirvReaderTest_ImageFetch_Multisampled_ConvertSampleOperand,
ImageAccessTest,
::testing::Values(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 = "let x_99 = textureLoad(x_20, vi12, i32(u1))",
}));
using SampledImageAccessTest = SpirvReaderTestWithParam<ImgData>;
TEST_P(SampledImageAccessTest, Variable) {
EXPECT_IR(R"(
OpCapability Shader
OpCapability ImageQuery
OpCapability Sampled1D
OpMemoryModel Logical Simple
OpEntryPoint Fragment %main "main"
OpExecutionMode %main OriginUpperLeft
OpDecorate %10 DescriptorSet 0
OpDecorate %10 Binding 0
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
%sampler = OpTypeSampler
%ptr_sampler = OpTypePointer UniformConstant %sampler
%im_ty = OpTypeImage )" +
GetParam().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
%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
%sam = OpLoad %sampler %10
%im = OpLoad %im_ty %20
%sampled_image = OpSampledImage %si_ty %im %sam
)" + GetParam().spirv_fn +
R"(
OpReturn
OpFunctionEnd
)",
R"( ... )");
}
INSTANTIATE_TEST_SUITE_P(
DISABLED_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 = "let x_99 = textureLoad(x_20, vi12, 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 = "let x_99 = textureLoad(x_20, vi12, 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 = "let x_99 = textureLoad(x_20, vi12, 0i)",
},
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 = "let x_99 = textureSample(x_20, x_10, vf12)",
}));
// From VUID-StandaloneSpirv-OpImageQuerySizeLod-04659:
// ImageQuerySizeLod requires Sampled=1
INSTANTIATE_TEST_SUITE_P(
DISABLED_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 = "let x_99 = i32(textureDimensions(x_20, i1)))",
},
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 = "let x_99 = vec2i(textureDimensions(x_20, i1))",
},
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 = "let x_99 = vec3i(textureDimensions(x_20, i1))",
},
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 = "let x_99 = vec2i(textureDimensions(x_20, i1).xy)",
},
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 = "let x_99 = vec2i(textureDimensions(x_20, i1))",
},
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 = "let x_99 = vec2i(textureDimensions(x_20, i1).xy)",
}));
// ImageQuerySize requires storage image or multisampled
// For storage image, use another instruction to indicate whether it
// is readonly or writeonly.
INSTANTIATE_TEST_SUITE_P(
DISABLED_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 =
"let x_99 = vec3i(vec3u(textureDimensions(x_20, i1), textureNumLayers(x_20)))",
},
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 =
"let x_99 = vec3i(vec3u(textureDimensions(x_20, i1).xy, textureNumLayers(x_20)))",
},
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 =
"let x_99 = vec3i(vec3u(textureDimensions(x_20, i1), textureNumLayers(x_20)))",
},
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 =
"let x_99 = vec3i(vec3u(textureDimensions(x_20, i1).xy, textureNumLayers(x_20)))",
}));
// textureDimensions accepts both signed and unsigned the level-of-detail values.
INSTANTIATE_TEST_SUITE_P(
DISABLED_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 = "let x_99 = i32(textureDimensions(x_20, u1))",
}));
// When SPIR-V wants the result type to be unsigned, we have to insert a value constructor or
// bitcast for WGSL to do the type coercion. But the algorithm already does that as a matter of
// course.
INSTANTIATE_TEST_SUITE_P(
DISABLED_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\n",
.wgsl_type = "texture_1d<f32>",
.wgsl_fn = "let x_99 = textureDimensions(x_20, i1)",
}));
INSTANTIATE_TEST_SUITE_P(DISABLED_SpirvReaderTest_ImageQueryLevels_SignedResult,
SampledImageAccessTest,
::testing::Values(
ImgData{
.name = "2D",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageQueryLevels %int %im\n",
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = "let x_99 = i32(textureNumLevels(x_20))",
},
ImgData{
.name = "2D array",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn = "%99 = OpImageQueryLevels %int %im\n",
.wgsl_type = "texture_2d_array<f32>",
.wgsl_fn = "let x_99 = i32(textureNumLevels(x_20))",
},
ImgData{
.name = "3D",
.spirv_type = "%float 3D 0 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageQueryLevels %int %im\n",
.wgsl_type = "texture_3d<f32>",
.wgsl_fn = "let x_99 = i32(textureNumLevels(x_20))",
},
ImgData{
.name = "Cube",
.spirv_type = "%float Cube 0 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageQueryLevels %int %im\n",
.wgsl_type = "texture_cube<f32>",
.wgsl_fn = "let x_99 = i32(textureNumLevels(x_20))",
},
ImgData{
.name = "Cube array",
.spirv_type = "%float Cube 0 1 0 1 Unknown",
.spirv_fn = "%99 = OpImageQueryLevels %int %im\n",
.wgsl_type = "texture_cube_array<f32>",
.wgsl_fn = "let x_99 = i32(textureNumLevels(x_20))",
},
ImgData{
.name = "depth 2d",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageQueryLevels %int %im\n",
.wgsl_type = "texture_depth_2d",
.wgsl_fn = "let x_99 = i32(textureNumLevels(x_20))",
},
ImgData{
.name = "depth 2d array",
.spirv_type = "%float 2D 1 1 0 1 Unknown",
.spirv_fn = "%99 = OpImageQueryLevels %int %im\n",
.wgsl_type = "texture_depth_2d_array",
.wgsl_fn = "let x_99 = i32(textureNumLevels(x_20))",
},
ImgData{
.name = "depth cube",
.spirv_type = "%float Cube 1 0 0 1 Unknown",
.spirv_fn = "%99 = OpImageQueryLevels %int %im\n",
.wgsl_type = "texture_depth_cube",
.wgsl_fn = "let x_99 = i32(textureNumLevels(x_20))",
},
ImgData{
.name = "depth cube array",
.spirv_type = "%float Cube 1 1 0 1 Unknown",
.spirv_fn = "%99 = OpImageQueryLevels %int %im\n",
.wgsl_type = "texture_depth_cube_array",
.wgsl_fn = "let x_99 = i32(textureNumLevels(x_20))",
}));
// Spot check that a value conversion is inserted when SPIR-V asks for an unsigned int result.
INSTANTIATE_TEST_SUITE_P(DISABLED_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 = "let x_99 = textureNumLevels(x_20)",
}));
using MultiSampledImageAccessTest = SpirvReaderTestWithParam<ImgData>;
TEST_P(MultiSampledImageAccessTest, 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
OpDecorate %10 DescriptorSet 0
OpDecorate %10 Binding 0
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
%v2uint = OpTypeVector %uint 2
%v3uint = OpTypeVector %uint 3
%v4uint = OpTypeVector %uint 4
%v4float = OpTypeVector %float 4
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
%sampler = OpTypeSampler
%ptr_sampler = OpTypePointer UniformConstant %sampler
%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
%10 = OpVariable %ptr_sampler UniformConstant
%20 = OpVariable %ptr_im_ty UniformConstant
%main = OpFunction %void None %voidfn
%entry = OpLabel
%sam = OpLoad %sampler %10
%im = OpLoad %im_ty %20
)" + params.spirv_fn +
R"(
OpReturn
OpFunctionEnd
)",
R"(
...
)");
}
INSTANTIATE_TEST_SUITE_P(DISABLED_SpirvReaderTest_ConvertResultSignedness,
MultiSampledImageAccessTest,
::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_2d<f32>",
.wgsl_fn = "let x_99 = textureLoad(x_20, vi12, 0i)",
},
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_2d<u32>",
.wgsl_fn = "let x_99 = textureLoad(x_20, vi12, 0i)",
},
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_2d<i32>",
.wgsl_fn = "let x_99 = textureLoad(x_20, vi12, 0i)",
}));
// ImageQuerySize requires storage image or multisampled
// For storage image, use another instruction to indicate whether it is readonly or writeonly.
INSTANTIATE_TEST_SUITE_P(
DISABLED_SpirvReaderTest_ImageQuerySize_NonArrayed_SignedResult,
MultiSampledImageAccessTest,
::testing::Values(
ImgData{
.name = "1D storage image",
.spirv_type = "%float 1D 0 0 0 2 Rgba32f",
.spirv_fn = "%99 = OpImageQuerySize %int %im\n"
"%98 = OpImageRead %v4float %im %int_1", // Implicitly mark as
.wgsl_type = "texture_1d<f32>",
.wgsl_fn = "let x_99 = i32(textureDimensions(x_20))",
},
ImgData{
.name = "2D storage image",
.spirv_type = "%float 2D 0 0 0 2 Rgba32f",
.spirv_fn = "%99 = OpImageQuerySize %v2int %im\n"
"%98 = OpImageRead %v4float %im %vi12", // Implicitly mark as
// NonWritable
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = "let x_99 = vec2i(textureDimensions(x_20))",
},
ImgData{
.name = "3D storage image",
.spirv_type = "%float 3D 0 0 0 2 Rgba32f",
.spirv_fn = "%99 = OpImageQuerySize %v3int %im\n"
"%98 = OpImageRead %v4float %im %vi123", // Implicitly mark as
// NonWritable
.wgsl_type = "texture_3d<f32>",
.wgsl_fn = "let x_99 = vec3i(textureDimensions(x_20))",
},
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 = "let x_99 = vec2i(textureDimensions(x_20))",
}));
// ImageQuerySize requires storage image or multisampled
// For storage image, use another instruction to indicate whether it is readonly or writeonly.
INSTANTIATE_TEST_SUITE_P(
DISABLED_SpirvReaderTest_ImageQuerySize_Arrayed_SignedResult,
MultiSampledImageAccessTest,
::testing::Values(ImgData{
.name = "2D array storage image",
.spirv_type = "%float 2D 0 1 0 2 Rgba32f",
.spirv_fn = "%99 = OpImageQuerySize %v3int %im\n"
"%98 = OpImageRead %v4float %im %vi123",
.wgsl_type = "texture_2d_array<f32>",
.wgsl_fn = "let x_99 = vec3i(vec3u(textureDimensions(x_20), textureNumLayers(x_20)))",
}));
// ImageQuerySize requires storage image or multisampled
// For storage image, use another instruction to indicate whether it is readonly or writeonly.
INSTANTIATE_TEST_SUITE_P(
DISABLED_SpirvReaderTest_ImageQuerySize_NonArrayed_UnsignedResult,
MultiSampledImageAccessTest,
::testing::Values(
ImgData{
.name = "1D storage image",
.spirv_type = "%float 1D 0 0 0 2 Rgba32f",
.spirv_fn = "%99 = OpImageQuerySize %uint %im\n"
"%98 = OpImageRead %v4float %im %int_1", // Implicitly mark as
// NonWritable
.wgsl_type = "texture_1d<f32>",
.wgsl_fn = "let x_99 = textureDimensions(x_20)",
},
ImgData{
.name = "2D storage image",
.spirv_type = "%float 2D 0 0 0 2 Rgba32f",
.spirv_fn = "%99 = OpImageQuerySize %v2uint %im\n"
"%98 = OpImageRead %v4float %im %vi12", // Implicitly mark as
// NonWritable
.wgsl_type = "texture_2d<f32>",
.wgsl_fn = "let x_99 = textureDimensions(x_20)",
},
ImgData{
.name = "3D storage image",
.spirv_type = "%float 3D 0 0 0 2 Rgba32f",
.spirv_fn = "%99 = OpImageQuerySize %v3uint %im\n"
"%98 = OpImageRead %v4float %im %vi123", // Implicitly mark as
// NonWritable
.wgsl_type = "texture_3d<f32>",
.wgsl_fn = "let x_99 = textureDimensions(x_20)",
},
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 = "let x_99 = textureDimensions(x_20)",
}));
INSTANTIATE_TEST_SUITE_P(
DISABLED_SpirvReaderTest_ImageQuerySize_Arrayed_UnsignedResult,
MultiSampledImageAccessTest,
::testing::Values(ImgData{
.name = "2D array storage image",
.spirv_type = "%float 2D 0 1 0 2 Rgba32f",
.spirv_fn = "%99 = OpImageQuerySize %v3uint %im\n"
"%98 = OpImageRead %v4float %im %vi123",
.wgsl_type = "texture_2d_array<f32>",
.wgsl_fn = "let x_99 = vec3u(textureDimensions(x_20), textureNumLayers(x_20))",
}));
INSTANTIATE_TEST_SUITE_P(DISABLED_SpirvReaderTest_ImageQuerySamples_SignedResult,
MultiSampledImageAccessTest,
::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 = "let x_99 = i32(textureNumSamples(x_20))",
}));
// Translation must inject a type coercion from unsigned to signed.
INSTANTIATE_TEST_SUITE_P(DISABLED_SpirvReaderTest_ImageQuerySamples_UnsignedResult,
MultiSampledImageAccessTest,
::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 = "let x_99 = textureNumSamples(x_20)",
}));
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
OpDecorate %10 DescriptorSet 0
OpDecorate %10 Binding 0
OpDecorate %20 DescriptorSet 2
OpDecorate %20 Binding 1
%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"( ... )");
}
INSTANTIATE_TEST_SUITE_P(
DISABLED_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 = "",
.wgsl_fn = "f1",
},
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 = "",
.wgsl_fn = "vf12.x",
},
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 = "",
.wgsl_fn = "vf123.x",
},
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 = "",
.wgsl_fn = "vf1234.x",
},
ImgData{
.name = "1D array",
.spirv_type = "%float 1D 0 1 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf12",
.wgsl_type = "",
.wgsl_fn = "vf12.x",
},
ImgData{
.name = "1D array one excess arg",
.spirv_type = "%float 1D 0 1 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf123",
.wgsl_type = "",
.wgsl_fn = "vf123.x",
},
ImgData{
.name = "1D array two excess args",
.spirv_type = "%float 1D 0 1 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf1234",
.wgsl_type = "",
.wgsl_fn = "vf1234.x",
},
ImgData{
.name = "2D",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf12",
.wgsl_type = "",
.wgsl_fn = {"vf12"},
},
ImgData{
.name = "2D one excess arg",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf123",
.wgsl_type = "",
.wgsl_fn = "vf123.xy",
},
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 = "",
.wgsl_fn = "vf1234.xy",
},
ImgData{
.name = "2D array",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf123",
.wgsl_type = "",
.wgsl_fn = "vf123.xy",
},
ImgData{
.name = "2D array one excess arg",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf1234",
.wgsl_type = "",
.wgsl_fn = "vf1234.xy",
},
ImgData{
.name = "3D",
.spirv_type = "%float 3D 0 0 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf123",
.wgsl_type = "",
.wgsl_fn = "vf123",
},
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 = "",
.wgsl_fn = "vf1234.xyz",
},
ImgData{
.name = "Cube",
.spirv_type = "%float Cube 0 0 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf123",
.wgsl_type = "",
.wgsl_fn = "vf123",
},
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 = "",
.wgsl_fn = "vf1234.xyz",
},
ImgData{
.name = "Cube array",
.spirv_type = "%float Cube 0 1 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf1234",
.wgsl_type = "",
.wgsl_fn = "vf1234.xyz",
}));
// In SPIR-V, sampling and dref sampling operations use floating point coordinates. Prove that we
// preserve floating point-ness. Test across all such instructions.
INSTANTIATE_TEST_SUITE_P(
DISABLED_SpirvReaderTest_PreserveFloatCoords_NonArrayed,
SampledImageCoordsTest,
::testing::Values(
// Scalar cases
ImgData{
.name = "1D",
.spirv_type = "%float 1D 0 0 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %float_1",
.wgsl_type = "",
.wgsl_fn = "f1",
},
ImgData{
.name = "1D lod",
.spirv_type = "%float 1D 0 0 0 1 Unknown",
.spirv_fn =
"%result = OpImageSampleExplicitLod %v4float %sampled_image %float_1 Lod %float_1",
.wgsl_type = "",
.wgsl_fn = "f1",
},
ImgData{
.name = "2D vec2",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf12",
.wgsl_type = "",
.wgsl_fn = "vf12",
},
ImgData{
.name = "2D vec2 lod",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn =
"%result = OpImageSampleExplicitLod %v4float %sampled_image %vf12 Lod %float_1",
.wgsl_type = "",
.wgsl_fn = "vf12",
},
ImgData{
.name = "2D vec2 depth",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn =
"%result = OpImageSampleDrefImplicitLod %float %sampled_image %vf12 %float_1",
.wgsl_type = "",
.wgsl_fn = "vf12",
},
ImgData{
.name = "2d vec2 depth lod",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn =
"%result = OpImageSampleDrefExplicitLod %float %sampled_image %vf12 %float_1 "
"Lod %float_0",
.wgsl_type = "",
.wgsl_fn = "vf12",
}));
// In SPIR-V, sampling and dref sampling operations use floating point coordinates. Prove that we
// preserve floating point-ness of the coordinate part, but convert the array index to signed
// integer.Test across all such instructions.
INSTANTIATE_TEST_SUITE_P(
DISABLED_SpirvReaderTest_PreserveFloatCoords_Arrayed,
SampledImageCoordsTest,
::testing::Values(
ImgData{
.name = "2D vec3",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleImplicitLod %v4float %sampled_image %vf123",
.wgsl_type = "",
.wgsl_fn = "vf123.xy",
},
ImgData{
.name = "2D lod",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn =
"%result = OpImageSampleExplicitLod %v4float %sampled_image %vf123 Lod %float_1",
.wgsl_type = "",
.wgsl_fn = "vf123.xy",
},
ImgData{
.name = "2D Depth",
.spirv_type = "%float 2D 1 1 0 1 Unknown",
.spirv_fn =
"%result = OpImageSampleDrefImplicitLod %float %sampled_image %vf123 %float_1",
.wgsl_type = "",
.wgsl_fn = "vf123.xy",
},
ImgData{
.name = "2D depth lod",
.spirv_type = "%float 2D 1 1 0 1 Unknown",
.spirv_fn =
"%result = OpImageSampleDrefExplicitLod %float %sampled_image %vf123 %float_1 Lod "
"%float_0",
.wgsl_type = "",
.wgsl_fn = "vf123.xy",
}));
INSTANTIATE_TEST_SUITE_P(DISABLED_SpirvReaderTest_ImageFetch,
SampledImageCoordsTest,
::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 = "",
.wgsl_fn = "i1",
},
ImgData{
.name = "2d signed",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "%result = OpImageFetch %v4float %im %vi12",
.wgsl_type = "",
.wgsl_fn = "vi12",
},
ImgData{
.name = "2D array signed",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn = "%result = OpImageFetch %v4float %im %vi123",
.wgsl_type = "",
.wgsl_fn = "vi123.xy",
},
ImgData{
.name = "1D unsigned",
.spirv_type = "%float 1D 0 0 0 1 Unknown",
.spirv_fn = "%result = OpImageFetch %v4float %im %uint_1",
.wgsl_type = "",
.wgsl_fn = "i32(u1)",
},
ImgData{
.name = "2D unsigned",
.spirv_type = "%float 2D 0 0 0 1 Unknown",
.spirv_fn = "%result = OpImageFetch %v4float %im %vu12",
.wgsl_type = "",
.wgsl_fn = "vec2i(vu12)",
},
ImgData{
.name = "2D array unsigned",
.spirv_type = "%float 2D 0 1 0 1 Unknown",
.spirv_fn = "%result = OpImageFetch %v4float %im %vu123",
.wgsl_type = "",
.wgsl_fn = "vec2i(vu123.xy)",
}));
// Metal requires comparison sampling with explicit Level-of-detail to use Lod 0. The SPIR-V reader
// requires the operand to be parsed as a constant 0 value. SPIR-V validation requires the Lod
// parameter to be a floating point value for non-fetch operations. So only test float values.
INSTANTIATE_TEST_SUITE_P(
DISABLED_SpirvReaderTest_ImageSampleDrefExplicitLod_CheckForLod0,
SampledImageCoordsTest,
::testing::Values(
ImgData{
.name = "2d 0.0",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleDrefExplicitLod %float %sampled_image "
"%vf1234 %float_1 Lod %float_0",
.wgsl_type = "",
.wgsl_fn = "vf1234.xy",
},
ImgData{
.name = "2D null",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleDrefExplicitLod %float %sampled_image "
"%vf1234 %float_1 Lod %float_null",
.wgsl_type = "",
.wgsl_fn = "vf1234.xy",
}));
// This is like the previous test, but for Projection sampling.
//
// Metal requires comparison sampling with explicit Level-of-detail to use Lod 0. The SPIR-V reader
// requires the operand to be parsed as a constant 0 value. SPIR-V validation requires the Lod
// parameter to be a floating point value for non-fetch operations. So only test float values.
INSTANTIATE_TEST_SUITE_P(
DISABLED_SpirvReaderTest_ImageSampleProjDrefExplicitLod_CheckForLod0,
SampledImageCoordsTest,
::testing::Values(
ImgData{
.name = "2D 0.0",
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleProjDrefExplicitLod %float %sampled_image "
"%vf1234 %float_1 Lod %float_0",
.wgsl_type = "",
.wgsl_fn = "(vf1234.xy / vf1234.z)",
},
ImgData{
.name = "2D null",
// float null works
.spirv_type = "%float 2D 1 0 0 1 Unknown",
.spirv_fn = "%result = OpImageSampleProjDrefExplicitLod %float %sampled_image "
"%vf1234 %float_1 Lod %float_null",
.wgsl_type = "",
.wgsl_fn = "(vf1234.xy / vf1234.z)",
}));
using NonSampledImageCoordsTest = SpirvReaderTestWithParam<ImgData>;
TEST_P(NonSampledImageCoordsTest, MakeCoordinateOperandsForImageAccess) {
auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
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
%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
%sampler = OpTypeSampler
%ptr_sampler = OpTypePointer UniformConstant %sampler
%im_ty = OpTypeImage )" +
params.spirv_type + R"(
%ptr_im_ty = OpTypePointer UniformConstant %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
)" + params.spirv_fn +
R"(
OpReturn
OpFunctionEnd
)",
R"( ... )");
}
// In SPIR-V, image read, fetch, and write use integer coordinates. Prove that we preserve signed
// integer coordinates.
INSTANTIATE_TEST_SUITE_P(DISABLED_SpirvReaderTest_PreserveIntCoords_NonArrayed,
NonSampledImageCoordsTest,
::testing::Values(
ImgData{
.name = "1D Read",
.spirv_type = "%float 1D 0 0 0 2 R32f",
.spirv_fn = "%result = OpImageRead %v4float %im %int_1",
.wgsl_type = "",
.wgsl_fn = "i1",
},
ImgData{
.name = "1D Write",
.spirv_type = "%float 1D 0 0 0 2 R32f",
.spirv_fn = "OpImageWrite %im %int_1 %vf1234",
.wgsl_type = "",
.wgsl_fn = "i1",
},
ImgData{
.name = "2D read",
.spirv_type = "%float 2D 0 0 0 2 R32f",
.spirv_fn = "%result = OpImageRead %v4float %im %vi12",
.wgsl_type = "",
.wgsl_fn = "vi12",
},
ImgData{
.name = "2D write",
.spirv_type = "%float 2D 0 0 0 2 R32f",
.spirv_fn = "OpImageWrite %im %vi12 %vf1234",
.wgsl_type = "",
.wgsl_fn = "vi12",
}));
// In SPIR-V, image read, fetch, and write use integer coordinates. Prove that we preserve signed
// integer coordinates.
INSTANTIATE_TEST_SUITE_P(DISABLED_SpirvReaderTest_PreserveIntCoords_Arrayed,
NonSampledImageCoordsTest,
::testing::Values(
ImgData{
.name = "2D read",
.spirv_type = "%float 2D 0 1 0 2 R32f",
.spirv_fn = "%result = OpImageRead %v4float %im %vi123",
.wgsl_type = "",
.wgsl_fn = "vi123.xy",
},
ImgData{
.name = "2D write",
.spirv_type = "%float 2D 0 1 0 2 R32f",
.spirv_fn = "OpImageWrite %im %vi123 %vf1234",
.wgsl_type = "",
.wgsl_fn = "vi123.xy",
}));
// In SPIR-V, image read, fetch, and write use integer coordinates. Prove that we convert unsigned
// integer coordinates to signed.
INSTANTIATE_TEST_SUITE_P(DISABLED_SpirvReaderTest_ConvertUintCoords_NonArrayed,
NonSampledImageCoordsTest,
::testing::Values(
ImgData{
.name = "1D read",
.spirv_type = "%float 1D 0 0 0 2 R32f",
.spirv_fn = "%result = OpImageRead %v4float %im %uint_1",
.wgsl_type = "",
.wgsl_fn = "i32(u1)",
},
ImgData{
.name = "1D write",
.spirv_type = "%float 1D 0 0 0 2 R32f",
.spirv_fn = "OpImageWrite %im %uint_1 %vf1234",
.wgsl_type = "",
.wgsl_fn = "i32(u1)",
},
ImgData{
.name = "2D read",
.spirv_type = "%float 2D 0 0 0 2 R32f",
.spirv_fn = "%result = OpImageRead %v4float %im %vu12",
.wgsl_type = "",
.wgsl_fn = "vec2i(vu12)",
},
ImgData{
.name = "2D write",
.spirv_type = "%float 2D 0 0 0 2 R32f",
.spirv_fn = "OpImageWrite %im %vu12 %vf1234",
.wgsl_type = "",
.wgsl_fn = "vec2i(vu12)",
}));
// In SPIR-V, image read, fetch, and write use integer coordinates. Prove that we convert unsigned
// integer coordinates to signed.
INSTANTIATE_TEST_SUITE_P(DISABLED_SpirvReaderTest_ConvertUintCoords_Arrayed,
NonSampledImageCoordsTest,
::testing::Values(
ImgData{
.name = "2D read",
.spirv_type = "%float 2D 0 1 0 2 R32f",
.spirv_fn = "%result = OpImageRead %v4float %im %vu123",
.wgsl_type = "",
.wgsl_fn = "vec2i(vu123.xy)",
},
ImgData{
.name = "2D write",
.spirv_type = "%float 2D 0 1 0 2 R32f",
.spirv_fn = "OpImageWrite %im %vu123 %vf1234",
.wgsl_type = "",
.wgsl_fn = "vec2i(vu123.xy)",
}));
// 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, DISABLED_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"(var var_1 : vec4f;
let x_22 = textureSample(x_2, x_3, vec2f());
let x_26 = textureSample(x_2, x_3, vec2f());
var_1 = (x_22 + x_26);
return;
)");
}
// 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, DISABLED_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
OpBranch %80
%80 = OpLabel
OpBranch %90
%90 = OpLabel
OpReturn
OpFunctionEnd
)",
R"(
switch(0i) {
default: {
if (true) {
let x_24 = textureSample(var_im, var_s, vec2f());
}
}
}
return;
)");
}
// 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, DISABLED_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"(
var x_900 : f32;
x_900 = 0.0f;
if (true) {
if (true) {
let x_18 = textureSample(x_20, x_10, x_900);
}
}
return;
)");
}
// 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, DISABLED_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"(
var x_900 : vec2f;
x_900 = vec2f();
if (true) {
if (true) {
let x_19 = textureSample(x_20, x_10, x_900.x);
}
}
return;
)");
}
// Demonstrates fix for crbug.com/tint/1642 The problem is the texel value for an image write can be
// given in 'var' declaration.
//
// The texel value handling has to unwrap the ref type first.
TEST_F(SpirvReaderTest, DISABLED_TexelTypeWhenLoop) {
EXPECT_IR(R"(
OpCapability Shader
OpCapability StorageImageExtendedFormats
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %100 "main"
OpExecutionMode %100 LocalSize 8 8 1
OpSource HLSL 600
OpName %type_2d_image "type.2d.image"
OpName %Output2Texture2D "Output2Texture2D"
OpName %100 "main"
OpDecorate %Output2Texture2D DescriptorSet 0
OpDecorate %Output2Texture2D Binding 0
%float = OpTypeFloat 32
%float_0 = OpConstant %float 0
%v2float = OpTypeVector %float 2
%7 = OpConstantComposite %v2float %float_0 %float_0
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%int_2 = OpConstant %int 2
%float_1 = OpConstant %float 1
%12 = OpConstantComposite %v2float %float_1 %float_1
%int_1 = OpConstant %int 1
%uint = OpTypeInt 32 0
%uint_1 = OpConstant %uint 1
%v2uint = OpTypeVector %uint 2
%17 = OpConstantComposite %v2uint %uint_1 %uint_1
%type_2d_image = OpTypeImage %float 2D 2 0 0 2 Rg32f
%_ptr_UniformConstant_type_2d_image = OpTypePointer UniformConstant %type_2d_image
%void = OpTypeVoid
%20 = OpTypeFunction %void
%bool = OpTypeBool
%Output2Texture2D = OpVariable %_ptr_UniformConstant_type_2d_image UniformConstant
%100 = OpFunction %void None %20
%22 = OpLabel
OpBranch %23
%23 = OpLabel
%24 = OpPhi %v2float %7 %22 %12 %25
%26 = OpPhi %int %int_0 %22 %27 %25
%28 = OpSLessThan %bool %26 %int_2
OpLoopMerge %29 %25 None
OpBranchConditional %28 %25 %29
%25 = OpLabel
%27 = OpIAdd %int %26 %int_1
OpBranch %23
%29 = OpLabel
%30 = OpLoad %type_2d_image %Output2Texture2D
OpImageWrite %30 %17 %24 None
OpReturn
OpFunctionEnd
)",
R"(
var x_24 : vec2f;
var x_26 : i32;
x_24 = vec2f();
x_26 = 0i;
loop {
var x_27 : i32;
if ((x_26 < 2i)) {
} else {
break;
}
continuing {
x_27 = (x_26 + 1i);
x_24 = vec2f(1.0f);
x_26 = x_27;
}
}
textureStore(Output2Texture2D, vec2i(vec2u(1u)), vec4f(x_24, 0.0f, 0.0f));
return;
)");
}
TEST_F(SpirvReaderTest, DISABLED_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"(
@group(0) @binding(0) var RWTexture2D : texture_storage_2d<rgba32float, read_write>;
const x_9 = vec2u(1u);
fn main_1() {
let x_31 = textureLoad(RWTexture2D, vec2i(x_9));
textureStore(RWTexture2D, vec2i(x_9), (x_31 + x_31));
return;
}
@compute @workgroup_size(8i, 8i, 1i)
fn main() {
main_1();
}
)");
}
// Demonstrates fix for crbug.com/tint/1642
// The problem is an operand to a simple select can be a value that is hoisted into a 'var'
// declaration.
//
// The selection-generation logic has to UnwrapRef if needed.
TEST_F(SpirvReaderTest, DISABLED_SimpleSelectCanSelectFromHoistedConstant) {
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %100 "main" %gl_Position
OpSource HLSL 600
OpName %100 "main"
OpDecorate %gl_Position BuiltIn Position
%float = OpTypeFloat 32
%float_0 = OpConstant %float 0
%float_1 = OpConstant %float 1
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
%void = OpTypeVoid
%9 = OpTypeFunction %void
%bool = OpTypeBool
%gl_Position = OpVariable %_ptr_Output_v4float Output
%11 = OpUndef %float
%100 = OpFunction %void None %9
%12 = OpLabel
OpBranch %13
%13 = OpLabel
%14 = OpPhi %float %11 %12 %15 %16
%15 = OpPhi %float %float_0 %12 %17 %16
%18 = OpFOrdLessThan %bool %15 %float_1
OpLoopMerge %19 %16 None
OpBranchConditional %18 %16 %19
%16 = OpLabel
%17 = OpFAdd %float %15 %float_1
OpBranch %13
%19 = OpLabel
%20 = OpFOrdGreaterThan %bool %14 %float_1
%21 = OpSelect %float %20 %14 %float_0
%22 = OpCompositeConstruct %v4float %21 %21 %21 %21
OpStore %gl_Position %22
OpReturn
OpFunctionEnd
)",
R"(
var x_14 : f32;
var x_15 : f32;
x_14 = 0.0f;
x_15 = 0.0f;
loop {
var x_17 : f32;
if ((x_15 < 1.0f)) {
} else {
break;
}
continuing {
x_17 = (x_15 + 1.0f);
let x_15_c16_1 = x_15;
x_14 = x_15_c16_1;
x_15 = x_17;
}
}
let x_21 = select(0.0f, x_14, (x_14 > 1.0f));
x_1 = vec4f(x_21);
return;
)");
}
TEST_F(SpirvReaderTest, DISABLED_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"(
@group(0) @binding(0) var x_7 : texture_2d<f32>;
@group(0) @binding(1) var x_8 : sampler;
fn main_1() {
let x_23 = textureSample(x_7, x_8, vec2f());
return;
}
@fragment
fn main() {
main_1();
}
)");
}
TEST_F(SpirvReaderTest, DISABLED_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"(
@group(0) @binding(0) var x_7 : texture_depth_2d;
@group(0) @binding(1) var x_8 : sampler_comparison;
fn main_1() {
let x_23 = textureSampleCompare(x_7, x_8, vec2f(), 0.0f);
return;
}
@fragment
fn main() {
main_1();
}
)");
}
} // namespace
} // namespace tint::spirv::reader