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

tint: Validate @must_use on functions and builtins 

Fixed: tint:1844
Change-Id: I812dd338925c1dccaa3029a46e10a0ea34a8ed54
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/120960
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
Commit-Queue: Ben Clayton <bclayton@chromium.org>
This commit is contained in:
Ben Clayton 2023-02-22 16:18:22 +00:00 committed by Dawn LUCI CQ
parent d84903201d
commit b549b3051e
59 changed files with 916 additions and 734 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
docs/tint/origin-trial-changes.md
View File

@ -4,6 +4,7 @@
### Breaking changes ### Breaking changes
* Most builtin functions that return a value can no longer be used as a call statement. [tint:1844](crbug.com/tint/1844)
* The `sig` member of the return type of `frexp()` has been renamed to `fract`. [tint:1766](crbug.com/tint/1766) * The `sig` member of the return type of `frexp()` has been renamed to `fract`. [tint:1766](crbug.com/tint/1766)
* Calling a function with multiple pointer arguments that alias each other is now a error. [tint:1675](crbug.com/tint/1675) * Calling a function with multiple pointer arguments that alias each other is now a error. [tint:1675](crbug.com/tint/1675)
* `type` deprecation has been removed. `alias` must be used now. [tint:1812](crbug.com/tint/1812) * `type` deprecation has been removed. `alias` must be used now. [tint:1812](crbug.com/tint/1812)

22
src/dawn/tests/unittests/validation/BindGroupValidationTests.cpp
View File

@ -2472,12 +2472,12 @@ TEST_F(BindGroupLayoutCompatibilityTest, TextureViewDimension) {
constexpr char kTexture2DShaderFS[] = R"( constexpr char kTexture2DShaderFS[] = R"(
@group(0) @binding(0) var myTexture : texture_2d<f32>; @group(0) @binding(0) var myTexture : texture_2d<f32>;
@fragment fn main() { @fragment fn main() {
textureDimensions(myTexture); _ = textureDimensions(myTexture);
})"; })";
constexpr char kTexture2DShaderCS[] = R"( constexpr char kTexture2DShaderCS[] = R"(
@group(0) @binding(0) var myTexture : texture_2d<f32>; @group(0) @binding(0) var myTexture : texture_2d<f32>;
@compute @workgroup_size(1) fn main() { @compute @workgroup_size(1) fn main() {
textureDimensions(myTexture); _ = textureDimensions(myTexture);
})"; })";
// Render: Test that 2D texture with 2D view dimension works // Render: Test that 2D texture with 2D view dimension works
@ -2511,12 +2511,12 @@ TEST_F(BindGroupLayoutCompatibilityTest, TextureViewDimension) {
constexpr char kTexture2DArrayShaderFS[] = R"( constexpr char kTexture2DArrayShaderFS[] = R"(
@group(0) @binding(0) var myTexture : texture_2d_array<f32>; @group(0) @binding(0) var myTexture : texture_2d_array<f32>;
@fragment fn main() { @fragment fn main() {
textureDimensions(myTexture); _ = textureDimensions(myTexture);
})"; })";
constexpr char kTexture2DArrayShaderCS[] = R"( constexpr char kTexture2DArrayShaderCS[] = R"(
@group(0) @binding(0) var myTexture : texture_2d_array<f32>; @group(0) @binding(0) var myTexture : texture_2d_array<f32>;
@compute @workgroup_size(1) fn main() { @compute @workgroup_size(1) fn main() {
textureDimensions(myTexture); _ = textureDimensions(myTexture);
})"; })";
// Render: Test that 2D texture array with 2D array view dimension works // Render: Test that 2D texture array with 2D array view dimension works
@ -2873,7 +2873,7 @@ TEST_F(SamplerTypeBindingTest, ShaderAndBGLMatches) {
@group(0) @binding(0) var mySampler: sampler; @group(0) @binding(0) var mySampler: sampler;
@group(0) @binding(1) var myTexture: texture_2d<f32>; @group(0) @binding(1) var myTexture: texture_2d<f32>;
@fragment fn main() { @fragment fn main() {
textureSample(myTexture, mySampler, vec2f(0.0, 0.0)); _ = textureSample(myTexture, mySampler, vec2f(0.0, 0.0));
})"); })");
} }
@ -2887,7 +2887,7 @@ TEST_F(SamplerTypeBindingTest, ShaderAndBGLMatches) {
@group(0) @binding(0) var mySampler: sampler; @group(0) @binding(0) var mySampler: sampler;
@group(0) @binding(1) var myTexture: texture_2d<f32>; @group(0) @binding(1) var myTexture: texture_2d<f32>;
@fragment fn main() { @fragment fn main() {
textureSample(myTexture, mySampler, vec2f(0.0, 0.0)); _ = textureSample(myTexture, mySampler, vec2f(0.0, 0.0));
})"); })");
} }
@ -2901,7 +2901,7 @@ TEST_F(SamplerTypeBindingTest, ShaderAndBGLMatches) {
@group(0) @binding(0) var mySampler: sampler; @group(0) @binding(0) var mySampler: sampler;
@group(0) @binding(1) var myTexture: texture_depth_2d; @group(0) @binding(1) var myTexture: texture_depth_2d;
@fragment fn main() { @fragment fn main() {
textureSample(myTexture, mySampler, vec2f(0.0, 0.0)); _ = textureSample(myTexture, mySampler, vec2f(0.0, 0.0));
})"); })");
} }
@ -2915,7 +2915,7 @@ TEST_F(SamplerTypeBindingTest, ShaderAndBGLMatches) {
@group(0) @binding(0) var mySampler: sampler; @group(0) @binding(0) var mySampler: sampler;
@group(0) @binding(1) var myTexture: texture_depth_2d; @group(0) @binding(1) var myTexture: texture_depth_2d;
@fragment fn main() { @fragment fn main() {
textureSample(myTexture, mySampler, vec2f(0.0, 0.0)); _ = textureSample(myTexture, mySampler, vec2f(0.0, 0.0));
})"); })");
} }
@ -2929,7 +2929,7 @@ TEST_F(SamplerTypeBindingTest, ShaderAndBGLMatches) {
@group(0) @binding(0) var mySampler: sampler_comparison; @group(0) @binding(0) var mySampler: sampler_comparison;
@group(0) @binding(1) var myTexture: texture_depth_2d; @group(0) @binding(1) var myTexture: texture_depth_2d;
@fragment fn main() { @fragment fn main() {
textureSampleCompare(myTexture, mySampler, vec2f(0.0, 0.0), 0.0); _ = textureSampleCompare(myTexture, mySampler, vec2f(0.0, 0.0), 0.0);
})"); })");
} }
@ -2943,7 +2943,7 @@ TEST_F(SamplerTypeBindingTest, ShaderAndBGLMatches) {
@group(0) @binding(0) var mySampler: sampler; @group(0) @binding(0) var mySampler: sampler;
@group(0) @binding(1) var myTexture: texture_2d<f32>; @group(0) @binding(1) var myTexture: texture_2d<f32>;
@fragment fn main() { @fragment fn main() {
textureSample(myTexture, mySampler, vec2f(0.0, 0.0)); _ = textureSample(myTexture, mySampler, vec2f(0.0, 0.0));
})")); })"));
} }
@ -2957,7 +2957,7 @@ TEST_F(SamplerTypeBindingTest, ShaderAndBGLMatches) {
@group(0) @binding(0) var mySampler: sampler; @group(0) @binding(0) var mySampler: sampler;
@group(0) @binding(1) var myTexture: texture_2d<f32>; @group(0) @binding(1) var myTexture: texture_2d<f32>;
@fragment fn main() { @fragment fn main() {
textureSample(myTexture, mySampler, vec2f(0.0, 0.0)); _ = textureSample(myTexture, mySampler, vec2f(0.0, 0.0));
})"); })");
} }
} }

42
src/dawn/tests/unittests/validation/GetBindGroupLayoutValidationTests.cpp
View File

@ -188,7 +188,7 @@ TEST_F(GetBindGroupLayoutTests, DefaultTextureSampleType) {
@group(0) @binding(0) var myTexture : texture_2d<f32>; @group(0) @binding(0) var myTexture : texture_2d<f32>;
@group(0) @binding(1) var mySampler : sampler; @group(0) @binding(1) var mySampler : sampler;
@vertex fn main() -> @builtin(position) vec4f { @vertex fn main() -> @builtin(position) vec4f {
textureLoad(myTexture, vec2i(), 0); _ = textureLoad(myTexture, vec2i(), 0);
_ = mySampler; _ = mySampler;
return vec4f(); return vec4f();
})"); })");
@ -197,7 +197,7 @@ TEST_F(GetBindGroupLayoutTests, DefaultTextureSampleType) {
@group(0) @binding(0) var myTexture : texture_2d<f32>; @group(0) @binding(0) var myTexture : texture_2d<f32>;
@group(0) @binding(1) var mySampler : sampler; @group(0) @binding(1) var mySampler : sampler;
@vertex fn main() -> @builtin(position) vec4f { @vertex fn main() -> @builtin(position) vec4f {
textureSampleLevel(myTexture, mySampler, vec2f(), 0.0); _ = textureSampleLevel(myTexture, mySampler, vec2f(), 0.0);
return vec4f(); return vec4f();
})"); })");
@ -213,7 +213,7 @@ TEST_F(GetBindGroupLayoutTests, DefaultTextureSampleType) {
@group(0) @binding(0) var myTexture : texture_2d<f32>; @group(0) @binding(0) var myTexture : texture_2d<f32>;
@group(0) @binding(1) var mySampler : sampler; @group(0) @binding(1) var mySampler : sampler;
@fragment fn main() { @fragment fn main() {
textureLoad(myTexture, vec2i(), 0); _ = textureLoad(myTexture, vec2i(), 0);
_ = mySampler; _ = mySampler;
})"); })");
@ -221,7 +221,7 @@ TEST_F(GetBindGroupLayoutTests, DefaultTextureSampleType) {
@group(0) @binding(0) var myTexture : texture_2d<f32>; @group(0) @binding(0) var myTexture : texture_2d<f32>;
@group(0) @binding(1) var mySampler : sampler; @group(0) @binding(1) var mySampler : sampler;
@fragment fn main() { @fragment fn main() {
textureSample(myTexture, mySampler, vec2f()); _ = textureSample(myTexture, mySampler, vec2f());
})"); })");
auto BGLFromModules = [this](wgpu::ShaderModule vertexModule, auto BGLFromModules = [this](wgpu::ShaderModule vertexModule,
@ -392,7 +392,7 @@ TEST_F(GetBindGroupLayoutTests, BindingType) {
@group(0) @binding(0) var myTexture : texture_2d<f32>; @group(0) @binding(0) var myTexture : texture_2d<f32>;
@fragment fn main() { @fragment fn main() {
textureDimensions(myTexture); _ = textureDimensions(myTexture);
})"); })");
EXPECT_TRUE(dawn::native::BindGroupLayoutBindingsEqualForTesting( EXPECT_TRUE(dawn::native::BindGroupLayoutBindingsEqualForTesting(
device.CreateBindGroupLayout(&desc).Get(), pipeline.GetBindGroupLayout(0).Get())); device.CreateBindGroupLayout(&desc).Get(), pipeline.GetBindGroupLayout(0).Get()));
@ -404,7 +404,7 @@ TEST_F(GetBindGroupLayoutTests, BindingType) {
@group(0) @binding(0) var myTexture : texture_multisampled_2d<f32>; @group(0) @binding(0) var myTexture : texture_multisampled_2d<f32>;
@fragment fn main() { @fragment fn main() {
textureDimensions(myTexture); _ = textureDimensions(myTexture);
})"); })");
EXPECT_TRUE(dawn::native::BindGroupLayoutBindingsEqualForTesting( EXPECT_TRUE(dawn::native::BindGroupLayoutBindingsEqualForTesting(
device.CreateBindGroupLayout(&desc).Get(), pipeline.GetBindGroupLayout(0).Get())); device.CreateBindGroupLayout(&desc).Get(), pipeline.GetBindGroupLayout(0).Get()));
@ -473,7 +473,7 @@ TEST_F(GetBindGroupLayoutTests, ViewDimension) {
@group(0) @binding(0) var myTexture : texture_1d<f32>; @group(0) @binding(0) var myTexture : texture_1d<f32>;
@fragment fn main() { @fragment fn main() {
textureDimensions(myTexture); _ = textureDimensions(myTexture);
})"); })");
EXPECT_TRUE(dawn::native::BindGroupLayoutBindingsEqualForTesting( EXPECT_TRUE(dawn::native::BindGroupLayoutBindingsEqualForTesting(
device.CreateBindGroupLayout(&desc).Get(), pipeline.GetBindGroupLayout(0).Get())); device.CreateBindGroupLayout(&desc).Get(), pipeline.GetBindGroupLayout(0).Get()));
@ -485,7 +485,7 @@ TEST_F(GetBindGroupLayoutTests, ViewDimension) {
@group(0) @binding(0) var myTexture : texture_2d<f32>; @group(0) @binding(0) var myTexture : texture_2d<f32>;
@fragment fn main() { @fragment fn main() {
textureDimensions(myTexture); _ = textureDimensions(myTexture);
})"); })");
EXPECT_TRUE(dawn::native::BindGroupLayoutBindingsEqualForTesting( EXPECT_TRUE(dawn::native::BindGroupLayoutBindingsEqualForTesting(
device.CreateBindGroupLayout(&desc).Get(), pipeline.GetBindGroupLayout(0).Get())); device.CreateBindGroupLayout(&desc).Get(), pipeline.GetBindGroupLayout(0).Get()));
@ -497,7 +497,7 @@ TEST_F(GetBindGroupLayoutTests, ViewDimension) {
@group(0) @binding(0) var myTexture : texture_2d_array<f32>; @group(0) @binding(0) var myTexture : texture_2d_array<f32>;
@fragment fn main() { @fragment fn main() {
textureDimensions(myTexture); _ = textureDimensions(myTexture);
})"); })");
EXPECT_TRUE(dawn::native::BindGroupLayoutBindingsEqualForTesting( EXPECT_TRUE(dawn::native::BindGroupLayoutBindingsEqualForTesting(
device.CreateBindGroupLayout(&desc).Get(), pipeline.GetBindGroupLayout(0).Get())); device.CreateBindGroupLayout(&desc).Get(), pipeline.GetBindGroupLayout(0).Get()));
@ -509,7 +509,7 @@ TEST_F(GetBindGroupLayoutTests, ViewDimension) {
@group(0) @binding(0) var myTexture : texture_3d<f32>; @group(0) @binding(0) var myTexture : texture_3d<f32>;
@fragment fn main() { @fragment fn main() {
textureDimensions(myTexture); _ = textureDimensions(myTexture);
})"); })");
EXPECT_TRUE(dawn::native::BindGroupLayoutBindingsEqualForTesting( EXPECT_TRUE(dawn::native::BindGroupLayoutBindingsEqualForTesting(
device.CreateBindGroupLayout(&desc).Get(), pipeline.GetBindGroupLayout(0).Get())); device.CreateBindGroupLayout(&desc).Get(), pipeline.GetBindGroupLayout(0).Get()));
@ -521,7 +521,7 @@ TEST_F(GetBindGroupLayoutTests, ViewDimension) {
@group(0) @binding(0) var myTexture : texture_cube<f32>; @group(0) @binding(0) var myTexture : texture_cube<f32>;
@fragment fn main() { @fragment fn main() {
textureDimensions(myTexture); _ = textureDimensions(myTexture);
})"); })");
EXPECT_TRUE(dawn::native::BindGroupLayoutBindingsEqualForTesting( EXPECT_TRUE(dawn::native::BindGroupLayoutBindingsEqualForTesting(
device.CreateBindGroupLayout(&desc).Get(), pipeline.GetBindGroupLayout(0).Get())); device.CreateBindGroupLayout(&desc).Get(), pipeline.GetBindGroupLayout(0).Get()));
@ -533,7 +533,7 @@ TEST_F(GetBindGroupLayoutTests, ViewDimension) {
@group(0) @binding(0) var myTexture : texture_cube_array<f32>; @group(0) @binding(0) var myTexture : texture_cube_array<f32>;
@fragment fn main() { @fragment fn main() {
textureDimensions(myTexture); _ = textureDimensions(myTexture);
})"); })");
EXPECT_TRUE(dawn::native::BindGroupLayoutBindingsEqualForTesting( EXPECT_TRUE(dawn::native::BindGroupLayoutBindingsEqualForTesting(
device.CreateBindGroupLayout(&desc).Get(), pipeline.GetBindGroupLayout(0).Get())); device.CreateBindGroupLayout(&desc).Get(), pipeline.GetBindGroupLayout(0).Get()));
@ -561,7 +561,7 @@ TEST_F(GetBindGroupLayoutTests, TextureComponentType) {
@group(0) @binding(0) var myTexture : texture_2d<f32>; @group(0) @binding(0) var myTexture : texture_2d<f32>;
@fragment fn main() { @fragment fn main() {
textureDimensions(myTexture); _ = textureDimensions(myTexture);
})"); })");
EXPECT_TRUE(dawn::native::BindGroupLayoutBindingsEqualForTesting( EXPECT_TRUE(dawn::native::BindGroupLayoutBindingsEqualForTesting(
device.CreateBindGroupLayout(&desc).Get(), pipeline.GetBindGroupLayout(0).Get())); device.CreateBindGroupLayout(&desc).Get(), pipeline.GetBindGroupLayout(0).Get()));
@ -573,7 +573,7 @@ TEST_F(GetBindGroupLayoutTests, TextureComponentType) {
@group(0) @binding(0) var myTexture : texture_2d<i32>; @group(0) @binding(0) var myTexture : texture_2d<i32>;
@fragment fn main() { @fragment fn main() {
textureDimensions(myTexture); _ = textureDimensions(myTexture);
})"); })");
EXPECT_TRUE(dawn::native::BindGroupLayoutBindingsEqualForTesting( EXPECT_TRUE(dawn::native::BindGroupLayoutBindingsEqualForTesting(
device.CreateBindGroupLayout(&desc).Get(), pipeline.GetBindGroupLayout(0).Get())); device.CreateBindGroupLayout(&desc).Get(), pipeline.GetBindGroupLayout(0).Get()));
@ -585,7 +585,7 @@ TEST_F(GetBindGroupLayoutTests, TextureComponentType) {
@group(0) @binding(0) var myTexture : texture_2d<u32>; @group(0) @binding(0) var myTexture : texture_2d<u32>;
@fragment fn main() { @fragment fn main() {
textureDimensions(myTexture); _ = textureDimensions(myTexture);
})"); })");
EXPECT_TRUE(dawn::native::BindGroupLayoutBindingsEqualForTesting( EXPECT_TRUE(dawn::native::BindGroupLayoutBindingsEqualForTesting(
device.CreateBindGroupLayout(&desc).Get(), pipeline.GetBindGroupLayout(0).Get())); device.CreateBindGroupLayout(&desc).Get(), pipeline.GetBindGroupLayout(0).Get()));
@ -898,7 +898,7 @@ TEST_F(GetBindGroupLayoutTests, ConflictingBindingTextureMultisampling) {
@group(0) @binding(0) var myTexture : texture_2d<f32>; @group(0) @binding(0) var myTexture : texture_2d<f32>;
@vertex fn main() -> @builtin(position) vec4f { @vertex fn main() -> @builtin(position) vec4f {
textureDimensions(myTexture); _ = textureDimensions(myTexture);
return vec4f(); return vec4f();
})"); })");
@ -906,7 +906,7 @@ TEST_F(GetBindGroupLayoutTests, ConflictingBindingTextureMultisampling) {
@group(0) @binding(0) var myTexture : texture_multisampled_2d<f32>; @group(0) @binding(0) var myTexture : texture_multisampled_2d<f32>;
@fragment fn main() { @fragment fn main() {
textureDimensions(myTexture); _ = textureDimensions(myTexture);
})"); })");
utils::ComboRenderPipelineDescriptor descriptor; utils::ComboRenderPipelineDescriptor descriptor;
@ -923,7 +923,7 @@ TEST_F(GetBindGroupLayoutTests, ConflictingBindingViewDimension) {
@group(0) @binding(0) var myTexture : texture_2d<f32>; @group(0) @binding(0) var myTexture : texture_2d<f32>;
@vertex fn main() -> @builtin(position) vec4f { @vertex fn main() -> @builtin(position) vec4f {
textureDimensions(myTexture); _ = textureDimensions(myTexture);
return vec4f(); return vec4f();
})"); })");
@ -931,7 +931,7 @@ TEST_F(GetBindGroupLayoutTests, ConflictingBindingViewDimension) {
@group(0) @binding(0) var myTexture : texture_3d<f32>; @group(0) @binding(0) var myTexture : texture_3d<f32>;
@fragment fn main() { @fragment fn main() {
textureDimensions(myTexture); _ = textureDimensions(myTexture);
})"); })");
utils::ComboRenderPipelineDescriptor descriptor; utils::ComboRenderPipelineDescriptor descriptor;
@ -948,7 +948,7 @@ TEST_F(GetBindGroupLayoutTests, ConflictingBindingTextureComponentType) {
@group(0) @binding(0) var myTexture : texture_2d<f32>; @group(0) @binding(0) var myTexture : texture_2d<f32>;
@vertex fn main() -> @builtin(position) vec4f { @vertex fn main() -> @builtin(position) vec4f {
textureDimensions(myTexture); _ = textureDimensions(myTexture);
return vec4f(); return vec4f();
})"); })");
@ -956,7 +956,7 @@ TEST_F(GetBindGroupLayoutTests, ConflictingBindingTextureComponentType) {
@group(0) @binding(0) var myTexture : texture_2d<i32>; @group(0) @binding(0) var myTexture : texture_2d<i32>;
@fragment fn main() { @fragment fn main() {
textureDimensions(myTexture); _ = textureDimensions(myTexture);
})"); })");
utils::ComboRenderPipelineDescriptor descriptor; utils::ComboRenderPipelineDescriptor descriptor;

4
src/dawn/tests/unittests/validation/RenderPipelineValidationTests.cpp
View File

@ -929,7 +929,7 @@ TEST_F(RenderPipelineValidationTest, TextureComponentTypeCompatibility) {
<< kScalarTypes[i] << R"(>; << kScalarTypes[i] << R"(>;
@fragment fn main() { @fragment fn main() {
textureDimensions(myTexture); _ = textureDimensions(myTexture);
})"; })";
descriptor.cFragment.module = utils::CreateShaderModule(device, stream.str().c_str()); descriptor.cFragment.module = utils::CreateShaderModule(device, stream.str().c_str());
descriptor.cTargets[0].writeMask = wgpu::ColorWriteMask::None; descriptor.cTargets[0].writeMask = wgpu::ColorWriteMask::None;
@ -978,7 +978,7 @@ TEST_F(RenderPipelineValidationTest, TextureViewDimensionCompatibility) {
@group(0) @binding(0) var myTexture : )" @group(0) @binding(0) var myTexture : )"
<< kTextureKeywords[i] << R"(<f32>; << kTextureKeywords[i] << R"(<f32>;
@fragment fn main() { @fragment fn main() {
textureDimensions(myTexture); _ = textureDimensions(myTexture);
})"; })";
descriptor.cFragment.module = utils::CreateShaderModule(device, stream.str().c_str()); descriptor.cFragment.module = utils::CreateShaderModule(device, stream.str().c_str());
descriptor.cTargets[0].writeMask = wgpu::ColorWriteMask::None; descriptor.cTargets[0].writeMask = wgpu::ColorWriteMask::None;

8
src/dawn/tests/unittests/validation/StorageTextureValidationTests.cpp
View File

@ -85,7 +85,7 @@ class StorageTextureValidationTests : public ValidationTest {
<< imageFormatQualifier << ", " << access << imageFormatQualifier << ", " << access
<< ">;\n" << ">;\n"
"@compute @workgroup_size(1) fn main() {\n" "@compute @workgroup_size(1) fn main() {\n"
" textureDimensions(image0);\n" " _ = textureDimensions(image0);\n"
"}\n"; "}\n";
return ostream.str(); return ostream.str();
@ -178,7 +178,7 @@ TEST_F(StorageTextureValidationTests, ReadWriteStorageTexture) {
ASSERT_DEVICE_ERROR(utils::CreateShaderModule(device, R"( ASSERT_DEVICE_ERROR(utils::CreateShaderModule(device, R"(
@group(0) @binding(0) var image0 : texture_storage_2d<rgba8unorm, read_write>; @group(0) @binding(0) var image0 : texture_storage_2d<rgba8unorm, read_write>;
@vertex fn main() { @vertex fn main() {
textureDimensions(image0); _ = textureDimensions(image0);
})")); })"));
} }
@ -187,7 +187,7 @@ TEST_F(StorageTextureValidationTests, ReadWriteStorageTexture) {
ASSERT_DEVICE_ERROR(utils::CreateShaderModule(device, R"( ASSERT_DEVICE_ERROR(utils::CreateShaderModule(device, R"(
@group(0) @binding(0) var image0 : texture_storage_2d<rgba8unorm, read_write>; @group(0) @binding(0) var image0 : texture_storage_2d<rgba8unorm, read_write>;
@fragment fn main() { @fragment fn main() {
textureDimensions(image0); _ = textureDimensions(image0);
})")); })"));
} }
@ -196,7 +196,7 @@ TEST_F(StorageTextureValidationTests, ReadWriteStorageTexture) {
ASSERT_DEVICE_ERROR(utils::CreateShaderModule(device, R"( ASSERT_DEVICE_ERROR(utils::CreateShaderModule(device, R"(
@group(0) @binding(0) var image0 : texture_storage_2d<rgba8unorm, read_write>; @group(0) @binding(0) var image0 : texture_storage_2d<rgba8unorm, read_write>;
@compute @workgroup_size(1) fn main() { @compute @workgroup_size(1) fn main() {
textureDimensions(image0); _ = textureDimensions(image0);
})")); })"));
} }
} }

156
src/tint/ast/builtin_texture_helper_test.cc
View File

File diff suppressed because it is too large Load Diff

11
src/tint/ast/builtin_texture_helper_test.h
View File

@ -196,7 +196,8 @@ struct TextureOverloadCase {
type::TextureDimension, type::TextureDimension,
TextureDataType, TextureDataType,
const char*, const char*,
std::function<Args(ProgramBuilder*)>); std::function<Args(ProgramBuilder*)>,
bool /* returns_value */);
/// Constructor for textureLoad() functions with non-storage textures /// Constructor for textureLoad() functions with non-storage textures
TextureOverloadCase(ValidTextureOverload, TextureOverloadCase(ValidTextureOverload,
const char*, const char*,
@ -204,7 +205,8 @@ struct TextureOverloadCase {
type::TextureDimension, type::TextureDimension,
TextureDataType, TextureDataType,
const char*, const char*,
std::function<Args(ProgramBuilder*)>); std::function<Args(ProgramBuilder*)>,
bool /* returns_value */);
/// Constructor for textureLoad() with storage textures /// Constructor for textureLoad() with storage textures
TextureOverloadCase(ValidTextureOverload, TextureOverloadCase(ValidTextureOverload,
const char*, const char*,
@ -213,7 +215,8 @@ struct TextureOverloadCase {
type::TextureDimension, type::TextureDimension,
TextureDataType, TextureDataType,
const char*, const char*,
std::function<Args(ProgramBuilder*)>); std::function<Args(ProgramBuilder*)>,
bool /* returns_value */);
/// Copy constructor /// Copy constructor
TextureOverloadCase(const TextureOverloadCase&); TextureOverloadCase(const TextureOverloadCase&);
/// Destructor /// Destructor
@ -258,6 +261,8 @@ struct TextureOverloadCase {
const char* const function; const char* const function;
/// A function that builds the AST arguments for the overload /// A function that builds the AST arguments for the overload
std::function<Args(ProgramBuilder*)> const args; std::function<Args(ProgramBuilder*)> const args;
/// True if the function returns a value
const bool returns_value;
}; };
std::ostream& operator<<(std::ostream& out, const TextureOverloadCase& data); std::ostream& operator<<(std::ostream& out, const TextureOverloadCase& data);

2
src/tint/fuzzers/tint_ast_fuzzer/mutations/delete_statement_test.cc
View File

@ -214,7 +214,7 @@ fn main() {
TEST(DeleteStatementTest, DeleteCall) { TEST(DeleteStatementTest, DeleteCall) {
auto original = R"( auto original = R"(
fn main() { fn main() {
sin(1.0); workgroupBarrier();
})"; })";
auto expected = R"( auto expected = R"(
fn main() { fn main() {

8
src/tint/inspector/inspector_test.cc
View File

@ -2899,7 +2899,7 @@ TEST_P(InspectorGetMultisampledTextureResourceBindingsTestWithParam, textureLoad
Func("ep", utils::Empty, ty.void_(), Func("ep", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(Call("textureLoad", "foo_texture", "foo_coords", "foo_sample_index")), Assign(Phony(), Call("textureLoad", "foo_texture", "foo_coords", "foo_sample_index")),
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
@ -3083,7 +3083,7 @@ TEST_P(InspectorGetDepthTextureResourceBindingsTestWithParam, textureDimensions)
Func("ep", utils::Empty, ty.void_(), Func("ep", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(Call("textureDimensions", "dt")), Assign(Phony(), Call("textureDimensions", "dt")),
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
@ -3120,7 +3120,7 @@ TEST_F(InspectorGetDepthMultisampledTextureResourceBindingsTest, textureDimensio
Func("ep", utils::Empty, ty.void_(), Func("ep", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(Call("textureDimensions", "tex")), Assign(Phony(), Call("textureDimensions", "tex")),
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
@ -3144,7 +3144,7 @@ TEST_F(InspectorGetExternalTextureResourceBindingsTest, Simple) {
Func("ep", utils::Empty, ty.void_(), Func("ep", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(Call("textureDimensions", "et")), Assign(Phony(), Call("textureDimensions", "et")),
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),

2
src/tint/resolver/alias_analysis_test.cc
View File

@ -558,7 +558,7 @@ TEST_P(Use, Read_UnaryMinus) {
TEST_P(Use, Read_FunctionCallArg) { TEST_P(Use, Read_FunctionCallArg) {
// abs(*p2); // abs(*p2);
Run(CallStmt(Call("abs", Deref("p2"))), R"(56:78 error: invalid aliased pointer argument Run(Assign(Phony(), Call("abs", Deref("p2"))), R"(56:78 error: invalid aliased pointer argument
12:34 note: aliases with another argument passed here)"); 12:34 note: aliases with another argument passed here)");
} }

24
src/tint/resolver/attribute_validation_test.cc
View File

@ -1482,9 +1482,7 @@ TEST_F(WorkgroupAttribute, NotAnEntryPoint) {
}); });
EXPECT_FALSE(r()->Resolve()); EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(), EXPECT_EQ(r()->error(), "12:34 error: @workgroup_size is only valid for compute stages");
"12:34 error: the workgroup_size attribute is only valid for "
"compute stages");
} }
TEST_F(WorkgroupAttribute, NotAComputeShader) { TEST_F(WorkgroupAttribute, NotAComputeShader) {
@ -1495,9 +1493,7 @@ TEST_F(WorkgroupAttribute, NotAComputeShader) {
}); });
EXPECT_FALSE(r()->Resolve()); EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(), EXPECT_EQ(r()->error(), "12:34 error: @workgroup_size is only valid for compute stages");
"12:34 error: the workgroup_size attribute is only valid for "
"compute stages");
} }
TEST_F(WorkgroupAttribute, DuplicateAttribute) { TEST_F(WorkgroupAttribute, DuplicateAttribute) {
@ -1964,4 +1960,20 @@ INSTANTIATE_TEST_SUITE_P(LocationTest,
} // namespace } // namespace
} // namespace InterpolateTests } // namespace InterpolateTests
namespace MustUseTests {
namespace {
using MustUseAttributeTest = ResolverTest;
TEST_F(MustUseAttributeTest, UsedOnFnWithNoReturnValue) {
Func("fn_must_use", utils::Empty, ty.void_(), utils::Empty,
utils::Vector{MustUse(Source{{12, 34}})});
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
R"(12:34 error: @must_use can only be applied to functions that return a value)");
}
} // namespace
} // namespace MustUseTests
} // namespace tint::resolver } // namespace tint::resolver

3
src/tint/resolver/builtin_test.cc
View File

@ -2438,7 +2438,8 @@ TEST_P(ResolverBuiltinTest_Texture, Call) {
param.BuildSamplerVariable(this); param.BuildSamplerVariable(this);
auto* call = Call(param.function, param.args(this)); auto* call = Call(param.function, param.args(this));
auto* stmt = CallStmt(call); auto* stmt = param.returns_value ? static_cast<const ast::Statement*>(Assign(Phony(), call))
: static_cast<const ast::Statement*>(CallStmt(call));
Func("func", utils::Empty, ty.void_(), utils::Vector{stmt}, Func("func", utils::Empty, ty.void_(), utils::Vector{stmt},
utils::Vector{Stage(ast::PipelineStage::kFragment)}); utils::Vector{Stage(ast::PipelineStage::kFragment)});

22
src/tint/resolver/builtin_validation_test.cc
View File

@ -42,7 +42,7 @@ TEST_F(ResolverBuiltinValidationTest, InvalidPipelineStageDirect) {
auto* dpdx = Call(Source{{3, 4}}, "dpdx", 1_f); auto* dpdx = Call(Source{{3, 4}}, "dpdx", 1_f);
Func(Source{{1, 2}}, "func", utils::Empty, ty.void_(), Func(Source{{1, 2}}, "func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(dpdx), Assign(Phony(), dpdx),
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kCompute), Stage(ast::PipelineStage::kCompute),
@ -62,7 +62,7 @@ TEST_F(ResolverBuiltinValidationTest, InvalidPipelineStageIndirect) {
auto* dpdx = Call(Source{{3, 4}}, "dpdx", 1_f); auto* dpdx = Call(Source{{3, 4}}, "dpdx", 1_f);
Func(Source{{1, 2}}, "f0", utils::Empty, ty.void_(), Func(Source{{1, 2}}, "f0", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(dpdx), Assign(Phony(), dpdx),
}); });
Func(Source{{3, 4}}, "f1", utils::Empty, ty.void_(), Func(Source{{3, 4}}, "f1", utils::Empty, ty.void_(),
@ -318,10 +318,14 @@ TEST_P(BuiltinTextureConstExprArgValidationTest, Immediate) {
arg_to_replace = expr(Source{{12, 34}}, *this); arg_to_replace = expr(Source{{12, 34}}, *this);
auto* call = Call(overload.function, args);
auto* stmt = overload.returns_value ? static_cast<const ast::Statement*>(Assign(Phony(), call))
: static_cast<const ast::Statement*>(CallStmt(call));
// Call the builtin with the constexpr argument replaced // Call the builtin with the constexpr argument replaced
Func("func", utils::Empty, ty.void_(), Func("func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(Call(overload.function, args)), stmt,
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
@ -371,10 +375,14 @@ TEST_P(BuiltinTextureConstExprArgValidationTest, GlobalConst) {
arg_to_replace = Expr(Source{{12, 34}}, "G"); arg_to_replace = Expr(Source{{12, 34}}, "G");
auto* call = Call(overload.function, args);
auto* stmt = overload.returns_value ? static_cast<const ast::Statement*>(Assign(Phony(), call))
: static_cast<const ast::Statement*>(CallStmt(call));
// Call the builtin with the constant-expression argument replaced // Call the builtin with the constant-expression argument replaced
Func("func", utils::Empty, ty.void_(), Func("func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(Call(overload.function, args)), stmt,
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
@ -420,10 +428,14 @@ TEST_P(BuiltinTextureConstExprArgValidationTest, GlobalVar) {
arg_to_replace = Expr(Source{{12, 34}}, "G"); arg_to_replace = Expr(Source{{12, 34}}, "G");
auto* call = Call(overload.function, args);
auto* stmt = overload.returns_value ? static_cast<const ast::Statement*>(Assign(Phony(), call))
: static_cast<const ast::Statement*>(CallStmt(call));
// Call the builtin with the constant-expression argument replaced // Call the builtin with the constant-expression argument replaced
Func("func", utils::Empty, ty.void_(), Func("func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(Call(overload.function, args)), stmt,
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),

8
src/tint/resolver/builtins_validation_test.cc
View File

@ -1161,7 +1161,7 @@ TEST_P(FloatAllMatching, Scalar) {
auto* builtin = Call(name, params); auto* builtin = Call(name, params);
Func("func", utils::Empty, ty.void_(), Func("func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(builtin), Assign(Phony(), builtin),
}, },
utils::Vector{ utils::Vector{
create<ast::StageAttribute>(ast::PipelineStage::kFragment), create<ast::StageAttribute>(ast::PipelineStage::kFragment),
@ -1182,7 +1182,7 @@ TEST_P(FloatAllMatching, Vec2) {
auto* builtin = Call(name, params); auto* builtin = Call(name, params);
Func("func", utils::Empty, ty.void_(), Func("func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(builtin), Assign(Phony(), builtin),
}, },
utils::Vector{ utils::Vector{
create<ast::StageAttribute>(ast::PipelineStage::kFragment), create<ast::StageAttribute>(ast::PipelineStage::kFragment),
@ -1203,7 +1203,7 @@ TEST_P(FloatAllMatching, Vec3) {
auto* builtin = Call(name, params); auto* builtin = Call(name, params);
Func("func", utils::Empty, ty.void_(), Func("func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(builtin), Assign(Phony(), builtin),
}, },
utils::Vector{ utils::Vector{
create<ast::StageAttribute>(ast::PipelineStage::kFragment), create<ast::StageAttribute>(ast::PipelineStage::kFragment),
@ -1224,7 +1224,7 @@ TEST_P(FloatAllMatching, Vec4) {
auto* builtin = Call(name, params); auto* builtin = Call(name, params);
Func("func", utils::Empty, ty.void_(), Func("func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(builtin), Assign(Phony(), builtin),
}, },
utils::Vector{ utils::Vector{
create<ast::StageAttribute>(ast::PipelineStage::kFragment), create<ast::StageAttribute>(ast::PipelineStage::kFragment),

21
src/tint/resolver/call_validation_test.cc
View File

@ -457,5 +457,26 @@ TEST_F(ResolverCallValidationTest, ComplexPointerChain_NotWholeVar_WithFullPtrPa
EXPECT_TRUE(r()->Resolve()); EXPECT_TRUE(r()->Resolve());
} }
TEST_F(ResolverCallValidationTest, MustUseFunction) {
Func(Source{{56, 78}}, "fn_must_use", utils::Empty, ty.i32(), utils::Vector{Return(1_i)},
utils::Vector{MustUse()});
Func("f", utils::Empty, ty.void_(),
utils::Vector{CallStmt(Call(Source{{12, 34}}, "fn_must_use"))});
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(
r()->error(),
R"(12:34 error: ignoring return value of function 'fn_must_use' annotated with @must_use
56:78 note: function 'fn_must_use' declared here)");
}
TEST_F(ResolverCallValidationTest, MustUseBuiltin) {
Func("f", utils::Empty, ty.void_(),
utils::Vector{CallStmt(Call(Source{{12, 34}}, "max", 1_a, 2_a))});
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(), "12:34 error: ignoring return value of builtin 'max'");
}
} // namespace } // namespace
} // namespace tint::resolver } // namespace tint::resolver

16
src/tint/resolver/intrinsic_table.cc
View File

@ -1280,20 +1280,22 @@ Impl::Builtin Impl::Lookup(sem::BuiltinType builtin_type,
builtin::AddressSpace::kUndefined, builtin::Access::kUndefined, p.usage)); builtin::AddressSpace::kUndefined, builtin::Access::kUndefined, p.usage));
} }
sem::PipelineStageSet supported_stages; sem::PipelineStageSet supported_stages;
if (match.overload->flags.Contains(OverloadFlag::kSupportsVertexPipeline)) { auto& overload = *match.overload;
if (overload.flags.Contains(OverloadFlag::kSupportsVertexPipeline)) {
supported_stages.Add(ast::PipelineStage::kVertex); supported_stages.Add(ast::PipelineStage::kVertex);
} }
if (match.overload->flags.Contains(OverloadFlag::kSupportsFragmentPipeline)) { if (overload.flags.Contains(OverloadFlag::kSupportsFragmentPipeline)) {
supported_stages.Add(ast::PipelineStage::kFragment); supported_stages.Add(ast::PipelineStage::kFragment);
} }
if (match.overload->flags.Contains(OverloadFlag::kSupportsComputePipeline)) { if (overload.flags.Contains(OverloadFlag::kSupportsComputePipeline)) {
supported_stages.Add(ast::PipelineStage::kCompute); supported_stages.Add(ast::PipelineStage::kCompute);
} }
auto eval_stage = match.overload->const_eval_fn ? sem::EvaluationStage::kConstant auto eval_stage = overload.const_eval_fn ? sem::EvaluationStage::kConstant
: sem::EvaluationStage::kRuntime; : sem::EvaluationStage::kRuntime;
return builder.create<sem::Builtin>( return builder.create<sem::Builtin>(builtin_type, match.return_type, std::move(params),
builtin_type, match.return_type, std::move(params), eval_stage, supported_stages, eval_stage, supported_stages,
match.overload->flags.Contains(OverloadFlag::kIsDeprecated)); overload.flags.Contains(OverloadFlag::kIsDeprecated),
overload.flags.Contains(OverloadFlag::kMustUse));
}); });
return Builtin{sem, match.overload->const_eval_fn}; return Builtin{sem, match.overload->const_eval_fn};
} }

2
src/tint/resolver/materialize_test.cc
View File

@ -340,7 +340,7 @@ TEST_P(MaterializeAbstractNumericToConcreteType, Test) {
WrapInFunction(CallStmt(Call("F", abstract_expr))); WrapInFunction(CallStmt(Call("F", abstract_expr)));
break; break;
case Method::kBuiltinArg: case Method::kBuiltinArg:
WrapInFunction(CallStmt(Call("min", target_expr(), abstract_expr))); WrapInFunction(Assign(Phony(), Call("min", target_expr(), abstract_expr)));
break; break;
case Method::kReturn: case Method::kReturn:
Func("F", utils::Empty, target_ty(), utils::Vector{Return(abstract_expr)}); Func("F", utils::Empty, target_ty(), utils::Vector{Return(abstract_expr)});

15
src/tint/resolver/resolver_test.cc
View File

@ -2128,8 +2128,9 @@ TEST_F(ResolverTest, TextureSampler_TextureSample) {
Binding(1_a)); Binding(1_a));
GlobalVar("s", ty.sampler(type::SamplerKind::kSampler), Group(1_a), Binding(2_a)); GlobalVar("s", ty.sampler(type::SamplerKind::kSampler), Group(1_a), Binding(2_a));
auto* call = CallStmt(Call("textureSample", "t", "s", vec2<f32>(1_f, 2_f))); auto* call = Call("textureSample", "t", "s", vec2<f32>(1_f, 2_f));
const ast::Function* f = Func("test_function", utils::Empty, ty.void_(), utils::Vector{call}, const ast::Function* f =
Func("test_function", utils::Empty, ty.void_(), utils::Vector{Assign(Phony(), call)},
utils::Vector{Stage(ast::PipelineStage::kFragment)}); utils::Vector{Stage(ast::PipelineStage::kFragment)});
EXPECT_TRUE(r()->Resolve()) << r()->error(); EXPECT_TRUE(r()->Resolve()) << r()->error();
@ -2146,7 +2147,7 @@ TEST_F(ResolverTest, TextureSampler_TextureSampleInFunction) {
Binding(1_a)); Binding(1_a));
GlobalVar("s", ty.sampler(type::SamplerKind::kSampler), Group(1_a), Binding(2_a)); GlobalVar("s", ty.sampler(type::SamplerKind::kSampler), Group(1_a), Binding(2_a));
auto* inner_call = CallStmt(Call("textureSample", "t", "s", vec2<f32>(1_f, 2_f))); auto* inner_call = Assign(Phony(), Call("textureSample", "t", "s", vec2<f32>(1_f, 2_f)));
const ast::Function* inner_func = const ast::Function* inner_func =
Func("inner_func", utils::Empty, ty.void_(), utils::Vector{inner_call}); Func("inner_func", utils::Empty, ty.void_(), utils::Vector{inner_call});
auto* outer_call = CallStmt(Call("inner_func")); auto* outer_call = CallStmt(Call("inner_func"));
@ -2172,10 +2173,10 @@ TEST_F(ResolverTest, TextureSampler_TextureSampleFunctionDiamondSameVariables) {
Binding(1_a)); Binding(1_a));
GlobalVar("s", ty.sampler(type::SamplerKind::kSampler), Group(1_a), Binding(2_a)); GlobalVar("s", ty.sampler(type::SamplerKind::kSampler), Group(1_a), Binding(2_a));
auto* inner_call_1 = CallStmt(Call("textureSample", "t", "s", vec2<f32>(1_f, 2_f))); auto* inner_call_1 = Assign(Phony(), Call("textureSample", "t", "s", vec2<f32>(1_f, 2_f)));
const ast::Function* inner_func_1 = const ast::Function* inner_func_1 =
Func("inner_func_1", utils::Empty, ty.void_(), utils::Vector{inner_call_1}); Func("inner_func_1", utils::Empty, ty.void_(), utils::Vector{inner_call_1});
auto* inner_call_2 = CallStmt(Call("textureSample", "t", "s", vec2<f32>(3_f, 4_f))); auto* inner_call_2 = Assign(Phony(), Call("textureSample", "t", "s", vec2<f32>(3_f, 4_f)));
const ast::Function* inner_func_2 = const ast::Function* inner_func_2 =
Func("inner_func_2", utils::Empty, ty.void_(), utils::Vector{inner_call_2}); Func("inner_func_2", utils::Empty, ty.void_(), utils::Vector{inner_call_2});
auto* outer_call_1 = CallStmt(Call("inner_func_1")); auto* outer_call_1 = CallStmt(Call("inner_func_1"));
@ -2209,10 +2210,10 @@ TEST_F(ResolverTest, TextureSampler_TextureSampleFunctionDiamondDifferentVariabl
Binding(2_a)); Binding(2_a));
GlobalVar("s", ty.sampler(type::SamplerKind::kSampler), Group(1_a), Binding(3_a)); GlobalVar("s", ty.sampler(type::SamplerKind::kSampler), Group(1_a), Binding(3_a));
auto* inner_call_1 = CallStmt(Call("textureSample", "t1", "s", vec2<f32>(1_f, 2_f))); auto* inner_call_1 = Assign(Phony(), Call("textureSample", "t1", "s", vec2<f32>(1_f, 2_f)));
const ast::Function* inner_func_1 = const ast::Function* inner_func_1 =
Func("inner_func_1", utils::Empty, ty.void_(), utils::Vector{inner_call_1}); Func("inner_func_1", utils::Empty, ty.void_(), utils::Vector{inner_call_1});
auto* inner_call_2 = CallStmt(Call("textureSample", "t2", "s", vec2<f32>(3_f, 4_f))); auto* inner_call_2 = Assign(Phony(), Call("textureSample", "t2", "s", vec2<f32>(3_f, 4_f)));
const ast::Function* inner_func_2 = const ast::Function* inner_func_2 =
Func("inner_func_2", utils::Empty, ty.void_(), utils::Vector{inner_call_2}); Func("inner_func_2", utils::Empty, ty.void_(), utils::Vector{inner_call_2});
auto* outer_call_1 = CallStmt(Call("inner_func_1")); auto* outer_call_1 = CallStmt(Call("inner_func_1"));

232
src/tint/resolver/side_effects_test.cc
View File

@ -134,16 +134,19 @@ struct Case {
const char* name; const char* name;
utils::Vector<const char*, 3> args; utils::Vector<const char*, 3> args;
bool has_side_effects; bool has_side_effects;
bool returns_value;
ast::PipelineStage pipeline_stage; ast::PipelineStage pipeline_stage;
}; };
static Case C(const char* name, static Case C(const char* name,
utils::VectorRef<const char*> args, utils::VectorRef<const char*> args,
bool has_side_effects, bool has_side_effects,
bool returns_value,
ast::PipelineStage stage = ast::PipelineStage::kFragment) { ast::PipelineStage stage = ast::PipelineStage::kFragment) {
Case c; Case c;
c.name = name; c.name = name;
c.args = std::move(args); c.args = std::move(args);
c.has_side_effects = has_side_effects; c.has_side_effects = has_side_effects;
c.returns_value = returns_value;
c.pipeline_stage = stage; c.pipeline_stage = stage;
return c; return c;
} }
@ -155,8 +158,7 @@ static std::ostream& operator<<(std::ostream& o, const Case& c) {
o << ", "; o << ", ";
} }
} }
o << "), "; o << ")";
o << "has_side_effects = " << c.has_side_effects;
return o; return o;
} }
@ -222,7 +224,11 @@ TEST_P(SideEffectsBuiltinTest, Test) {
attrs.Push(WorkgroupSize(Expr(1_u))); attrs.Push(WorkgroupSize(Expr(1_u)));
} }
if (c.returns_value) {
stmts.Push(Assign(Phony(), expr));
} else {
stmts.Push(CallStmt(expr)); stmts.Push(CallStmt(expr));
}
Func("func", utils::Empty, ty.void_(), stmts, attrs); Func("func", utils::Empty, ty.void_(), stmts, attrs);
@ -237,120 +243,124 @@ INSTANTIATE_TEST_SUITE_P(
SideEffectsBuiltinTest, SideEffectsBuiltinTest,
testing::ValuesIn(std::vector<Case>{ testing::ValuesIn(std::vector<Case>{
// No side-effect builts // No side-effect builts
C("abs", utils::Vector{"f"}, false), // C("abs", utils::Vector{"f"}, false, true), //
C("acos", utils::Vector{"f"}, false), // C("acos", utils::Vector{"f"}, false, true), //
C("acosh", utils::Vector{"f"}, false), // C("acosh", utils::Vector{"f"}, false, true), //
C("all", utils::Vector{"vb"}, false), // C("all", utils::Vector{"vb"}, false, true), //
C("any", utils::Vector{"vb"}, false), // C("any", utils::Vector{"vb"}, false, true), //
C("arrayLength", utils::Vector{"pstorage_arr"}, false), // C("arrayLength", utils::Vector{"pstorage_arr"}, false, true), //
C("asin", utils::Vector{"f"}, false), // C("asin", utils::Vector{"f"}, false, true), //
C("asinh", utils::Vector{"f"}, false), // C("asinh", utils::Vector{"f"}, false, true), //
C("atan", utils::Vector{"f"}, false), // C("atan", utils::Vector{"f"}, false, true), //
C("atan2", utils::Vector{"f", "f"}, false), // C("atan2", utils::Vector{"f", "f"}, false, true), //
C("atanh", utils::Vector{"f"}, false), // C("atanh", utils::Vector{"f"}, false, true), //
C("atomicLoad", utils::Vector{"pa"}, false), // C("atomicLoad", utils::Vector{"pa"}, false, true), //
C("ceil", utils::Vector{"f"}, false), // C("ceil", utils::Vector{"f"}, false, true), //
C("clamp", utils::Vector{"f", "f", "f"}, false), // C("clamp", utils::Vector{"f", "f", "f"}, false, true), //
C("cos", utils::Vector{"f"}, false), // C("cos", utils::Vector{"f"}, false, true), //
C("cosh", utils::Vector{"f"}, false), // C("cosh", utils::Vector{"f"}, false, true), //
C("countLeadingZeros", utils::Vector{"i"}, false), // C("countLeadingZeros", utils::Vector{"i"}, false, true), //
C("countOneBits", utils::Vector{"i"}, false), // C("countOneBits", utils::Vector{"i"}, false, true), //
C("countTrailingZeros", utils::Vector{"i"}, false), // C("countTrailingZeros", utils::Vector{"i"}, false, true), //
C("cross", utils::Vector{"vf", "vf"}, false), // C("cross", utils::Vector{"vf", "vf"}, false, true), //
C("degrees", utils::Vector{"f"}, false), // C("degrees", utils::Vector{"f"}, false, true), //
C("determinant", utils::Vector{"m"}, false), // C("determinant", utils::Vector{"m"}, false, true), //
C("distance", utils::Vector{"f", "f"}, false), // C("distance", utils::Vector{"f", "f"}, false, true), //
C("dot", utils::Vector{"vf", "vf"}, false), // C("dot", utils::Vector{"vf", "vf"}, false, true), //
C("dot4I8Packed", utils::Vector{"u", "u"}, false), // C("dot4I8Packed", utils::Vector{"u", "u"}, false, true), //
C("dot4U8Packed", utils::Vector{"u", "u"}, false), // C("dot4U8Packed", utils::Vector{"u", "u"}, false, true), //
C("exp", utils::Vector{"f"}, false), // C("exp", utils::Vector{"f"}, false, true), //
C("exp2", utils::Vector{"f"}, false), // C("exp2", utils::Vector{"f"}, false, true), //
C("extractBits", utils::Vector{"i", "u", "u"}, false), // C("extractBits", utils::Vector{"i", "u", "u"}, false, true), //
C("faceForward", utils::Vector{"vf", "vf", "vf"}, false), // C("faceForward", utils::Vector{"vf", "vf", "vf"}, false, true), //
C("firstLeadingBit", utils::Vector{"u"}, false), // C("firstLeadingBit", utils::Vector{"u"}, false, true), //
C("firstTrailingBit", utils::Vector{"u"}, false), // C("firstTrailingBit", utils::Vector{"u"}, false, true), //
C("floor", utils::Vector{"f"}, false), // C("floor", utils::Vector{"f"}, false, true), //
C("fma", utils::Vector{"f", "f", "f"}, false), // C("fma", utils::Vector{"f", "f", "f"}, false, true), //
C("fract", utils::Vector{"vf"}, false), // C("fract", utils::Vector{"vf"}, false, true), //
C("frexp", utils::Vector{"f"}, false), // C("frexp", utils::Vector{"f"}, false, true), //
C("insertBits", utils::Vector{"i", "i", "u", "u"}, false), // C("insertBits", utils::Vector{"i", "i", "u", "u"}, false, true), //
C("inverseSqrt", utils::Vector{"f"}, false), // C("inverseSqrt", utils::Vector{"f"}, false, true), //
C("ldexp", utils::Vector{"f", "i"}, false), // C("ldexp", utils::Vector{"f", "i"}, false, true), //
C("length", utils::Vector{"vf"}, false), // C("length", utils::Vector{"vf"}, false, true), //
C("log", utils::Vector{"f"}, false), // C("log", utils::Vector{"f"}, false, true), //
C("log2", utils::Vector{"f"}, false), // C("log2", utils::Vector{"f"}, false, true), //
C("max", utils::Vector{"f", "f"}, false), // C("max", utils::Vector{"f", "f"}, false, true), //
C("min", utils::Vector{"f", "f"}, false), // C("min", utils::Vector{"f", "f"}, false, true), //
C("mix", utils::Vector{"f", "f", "f"}, false), // C("mix", utils::Vector{"f", "f", "f"}, false, true), //
C("modf", utils::Vector{"f"}, false), // C("modf", utils::Vector{"f"}, false, true), //
C("normalize", utils::Vector{"vf"}, false), // C("normalize", utils::Vector{"vf"}, false, true), //
C("pack2x16float", utils::Vector{"vf2"}, false), // C("pack2x16float", utils::Vector{"vf2"}, false, true), //
C("pack2x16snorm", utils::Vector{"vf2"}, false), // C("pack2x16snorm", utils::Vector{"vf2"}, false, true), //
C("pack2x16unorm", utils::Vector{"vf2"}, false), // C("pack2x16unorm", utils::Vector{"vf2"}, false, true), //
C("pack4x8snorm", utils::Vector{"vf4"}, false), // C("pack4x8snorm", utils::Vector{"vf4"}, false, true), //
C("pack4x8unorm", utils::Vector{"vf4"}, false), // C("pack4x8unorm", utils::Vector{"vf4"}, false, true), //
C("pow", utils::Vector{"f", "f"}, false), // C("pow", utils::Vector{"f", "f"}, false, true), //
C("radians", utils::Vector{"f"}, false), // C("radians", utils::Vector{"f"}, false, true), //
C("reflect", utils::Vector{"vf", "vf"}, false), // C("reflect", utils::Vector{"vf", "vf"}, false, true), //
C("refract", utils::Vector{"vf", "vf", "f"}, false), // C("refract", utils::Vector{"vf", "vf", "f"}, false, true), //
C("reverseBits", utils::Vector{"u"}, false), // C("reverseBits", utils::Vector{"u"}, false, true), //
C("round", utils::Vector{"f"}, false), // C("round", utils::Vector{"f"}, false, true), //
C("select", utils::Vector{"f", "f", "b"}, false), // C("select", utils::Vector{"f", "f", "b"}, false, true), //
C("sign", utils::Vector{"f"}, false), // C("sign", utils::Vector{"f"}, false, true), //
C("sin", utils::Vector{"f"}, false), // C("sin", utils::Vector{"f"}, false, true), //
C("sinh", utils::Vector{"f"}, false), // C("sinh", utils::Vector{"f"}, false, true), //
C("smoothstep", utils::Vector{"f", "f", "f"}, false), // C("smoothstep", utils::Vector{"f", "f", "f"}, false, true), //
C("sqrt", utils::Vector{"f"}, false), // C("sqrt", utils::Vector{"f"}, false, true), //
C("step", utils::Vector{"f", "f"}, false), // C("step", utils::Vector{"f", "f"}, false, true), //
C("tan", utils::Vector{"f"}, false), // C("tan", utils::Vector{"f"}, false, true), //
C("tanh", utils::Vector{"f"}, false), // C("tanh", utils::Vector{"f"}, false, true), //
C("textureDimensions", utils::Vector{"t2d"}, false), // C("textureDimensions", utils::Vector{"t2d"}, false, true), //
C("textureGather", utils::Vector{"tdepth2d", "s2d", "vf2"}, false), // C("textureGather", utils::Vector{"tdepth2d", "s2d", "vf2"}, false, true), //
C("textureGatherCompare", utils::Vector{"tdepth2d", "scomp", "vf2", "f"}, false), // C("textureGatherCompare", utils::Vector{"tdepth2d", "scomp", "vf2", "f"}, false, true), //
C("textureLoad", utils::Vector{"t2d", "vi2", "i"}, false), // C("textureLoad", utils::Vector{"t2d", "vi2", "i"}, false, true), //
C("textureNumLayers", utils::Vector{"t2d_arr"}, false), // C("textureNumLayers", utils::Vector{"t2d_arr"}, false, true), //
C("textureNumLevels", utils::Vector{"t2d"}, false), // C("textureNumLevels", utils::Vector{"t2d"}, false, true), //
C("textureNumSamples", utils::Vector{"t2d_multi"}, false), // C("textureNumSamples", utils::Vector{"t2d_multi"}, false, true), //
C("textureSampleCompareLevel", utils::Vector{"tdepth2d", "scomp", "vf2", "f"}, false), // C("textureSampleCompareLevel",
C("textureSampleGrad", utils::Vector{"t2d", "s2d", "vf2", "vf2", "vf2"}, false), // utils::Vector{"tdepth2d", "scomp", "vf2", "f"},
C("textureSampleLevel", utils::Vector{"t2d", "s2d", "vf2", "f"}, false), // false,
C("transpose", utils::Vector{"m"}, false), // true), //
C("trunc", utils::Vector{"f"}, false), // C("textureSampleGrad", utils::Vector{"t2d", "s2d", "vf2", "vf2", "vf2"}, false, true), //
C("unpack2x16float", utils::Vector{"u"}, false), // C("textureSampleLevel", utils::Vector{"t2d", "s2d", "vf2", "f"}, false, true), //
C("unpack2x16snorm", utils::Vector{"u"}, false), // C("transpose", utils::Vector{"m"}, false, true), //
C("unpack2x16unorm", utils::Vector{"u"}, false), // C("trunc", utils::Vector{"f"}, false, true), //
C("unpack4x8snorm", utils::Vector{"u"}, false), // C("unpack2x16float", utils::Vector{"u"}, false, true), //
C("unpack4x8unorm", utils::Vector{"u"}, false), // C("unpack2x16snorm", utils::Vector{"u"}, false, true), //
C("storageBarrier", utils::Empty, false, ast::PipelineStage::kCompute), // C("unpack2x16unorm", utils::Vector{"u"}, false, true), //
C("workgroupBarrier", utils::Empty, false, ast::PipelineStage::kCompute), // C("unpack4x8snorm", utils::Vector{"u"}, false, true), //
C("textureSample", utils::Vector{"t2d", "s2d", "vf2"}, false), // C("unpack4x8unorm", utils::Vector{"u"}, false, true), //
C("textureSampleBias", utils::Vector{"t2d", "s2d", "vf2", "f"}, false), // C("storageBarrier", utils::Empty, false, false, ast::PipelineStage::kCompute), //
C("textureSampleCompare", utils::Vector{"tdepth2d", "scomp", "vf2", "f"}, false), // C("workgroupBarrier", utils::Empty, false, false, ast::PipelineStage::kCompute), //
C("dpdx", utils::Vector{"f"}, false), // C("textureSample", utils::Vector{"t2d", "s2d", "vf2"}, false, true), //
C("dpdxCoarse", utils::Vector{"f"}, false), // C("textureSampleBias", utils::Vector{"t2d", "s2d", "vf2", "f"}, false, true), //
C("dpdxFine", utils::Vector{"f"}, false), // C("textureSampleCompare", utils::Vector{"tdepth2d", "scomp", "vf2", "f"}, false, true), //
C("dpdy", utils::Vector{"f"}, false), // C("dpdx", utils::Vector{"f"}, false, true), //
C("dpdyCoarse", utils::Vector{"f"}, false), // C("dpdxCoarse", utils::Vector{"f"}, false, true), //
C("dpdyFine", utils::Vector{"f"}, false), // C("dpdxFine", utils::Vector{"f"}, false, true), //
C("fwidth", utils::Vector{"f"}, false), // C("dpdy", utils::Vector{"f"}, false, true), //
C("fwidthCoarse", utils::Vector{"f"}, false), // C("dpdyCoarse", utils::Vector{"f"}, false, true), //
C("fwidthFine", utils::Vector{"f"}, false), // C("dpdyFine", utils::Vector{"f"}, false, true), //
C("fwidth", utils::Vector{"f"}, false, true), //
C("fwidthCoarse", utils::Vector{"f"}, false, true), //
C("fwidthFine", utils::Vector{"f"}, false, true), //
// Side-effect builtins // Side-effect builtins
C("atomicAdd", utils::Vector{"pa", "i"}, true), // C("atomicAdd", utils::Vector{"pa", "i"}, true, true), //
C("atomicAnd", utils::Vector{"pa", "i"}, true), // C("atomicAnd", utils::Vector{"pa", "i"}, true, true), //
C("atomicCompareExchangeWeak", utils::Vector{"pa", "i", "i"}, true), // C("atomicCompareExchangeWeak", utils::Vector{"pa", "i", "i"}, true, true), //
C("atomicExchange", utils::Vector{"pa", "i"}, true), // C("atomicExchange", utils::Vector{"pa", "i"}, true, true), //
C("atomicMax", utils::Vector{"pa", "i"}, true), // C("atomicMax", utils::Vector{"pa", "i"}, true, true), //
C("atomicMin", utils::Vector{"pa", "i"}, true), // C("atomicMin", utils::Vector{"pa", "i"}, true, true), //
C("atomicOr", utils::Vector{"pa", "i"}, true), // C("atomicOr", utils::Vector{"pa", "i"}, true, true), //
C("atomicStore", utils::Vector{"pa", "i"}, true), // C("atomicStore", utils::Vector{"pa", "i"}, true, false), //
C("atomicSub", utils::Vector{"pa", "i"}, true), // C("atomicSub", utils::Vector{"pa", "i"}, true, true), //
C("atomicXor", utils::Vector{"pa", "i"}, true), // C("atomicXor", utils::Vector{"pa", "i"}, true, true), //
C("textureStore", utils::Vector{"tstorage2d", "vi2", "vf4"}, true), // C("textureStore", utils::Vector{"tstorage2d", "vi2", "vf4"}, true, false), //
C("workgroupUniformLoad", C("workgroupUniformLoad",
utils::Vector{"pworkgroup_arr"}, utils::Vector{"pworkgroup_arr"},
true, true,
true,
ast::PipelineStage::kCompute), // ast::PipelineStage::kCompute), //
// Unimplemented builtins // Unimplemented builtins

108
src/tint/resolver/uniformity_test.cc
View File

@ -168,14 +168,14 @@ class BasicTest : public UniformityAnalysisTestBase,
} }
/// Convert a function call to its string representation. /// Convert a function call to its string representation.
static std::string FunctionToStr(Function f) { static std::string FunctionCallToStr(Function f) {
switch (f) { switch (f) {
case kUserNoRestriction: case kUserNoRestriction:
return "user_no_restriction()"; return "user_no_restriction()";
case kMin: case kMin:
return "min(1, 1)"; return "_ = min(1, 1)";
case kTextureSampleLevel: case kTextureSampleLevel:
return "textureSampleLevel(t, s, vec2(0.5, 0.5), 0.0)"; return "_ = textureSampleLevel(t, s, vec2(0.5, 0.5), 0.0)";
case kUserRequiredToBeUniform: case kUserRequiredToBeUniform:
return "user_required_to_be_uniform()"; return "user_required_to_be_uniform()";
case kWorkgroupBarrier: case kWorkgroupBarrier:
@ -183,31 +183,31 @@ class BasicTest : public UniformityAnalysisTestBase,
case kStorageBarrier: case kStorageBarrier:
return "storageBarrier()"; return "storageBarrier()";
case kWorkgroupUniformLoad: case kWorkgroupUniformLoad:
return "workgroupUniformLoad(&w)"; return "_ = workgroupUniformLoad(&w)";
case kTextureSample: case kTextureSample:
return "textureSample(t, s, vec2(0.5, 0.5))"; return "_ = textureSample(t, s, vec2(0.5, 0.5))";
case kTextureSampleBias: case kTextureSampleBias:
return "textureSampleBias(t, s, vec2(0.5, 0.5), 2.0)"; return "_ = textureSampleBias(t, s, vec2(0.5, 0.5), 2.0)";
case kTextureSampleCompare: case kTextureSampleCompare:
return "textureSampleCompare(td, sc, vec2(0.5, 0.5), 0.5)"; return "_ = textureSampleCompare(td, sc, vec2(0.5, 0.5), 0.5)";
case kDpdx: case kDpdx:
return "dpdx(1.0)"; return "_ = dpdx(1.0)";
case kDpdxCoarse: case kDpdxCoarse:
return "dpdxCoarse(1.0)"; return "_ = dpdxCoarse(1.0)";
case kDpdxFine: case kDpdxFine:
return "dpdxFine(1.0)"; return "_ = dpdxFine(1.0)";
case kDpdy: case kDpdy:
return "dpdy(1.0)"; return "_ = dpdy(1.0)";
case kDpdyCoarse: case kDpdyCoarse:
return "dpdyCoarse(1.0)"; return "_ = dpdyCoarse(1.0)";
case kDpdyFine: case kDpdyFine:
return "dpdyFine(1.0)"; return "_ = dpdyFine(1.0)";
case kFwidth: case kFwidth:
return "fwidth(1.0)"; return "_ = fwidth(1.0)";
case kFwidthCoarse: case kFwidthCoarse:
return "fwidthCoarse(1.0)"; return "_ = fwidthCoarse(1.0)";
case kFwidthFine: case kFwidthFine:
return "fwidthFine(1.0)"; return "_ = fwidthFine(1.0)";
case kEndOfFunctionRange: case kEndOfFunctionRange:
return "<invalid>"; return "<invalid>";
} }
@ -314,7 +314,7 @@ fn foo() {
if ()" + ConditionToStr(condition) + if ()" + ConditionToStr(condition) +
R"() { R"() {
)" + FunctionToStr(function) + )" + FunctionCallToStr(function) +
R"(; R"(;
} }
} }
@ -610,7 +610,7 @@ TEST_P(FragmentBuiltin, AsParam) {
fn main(@builtin()" + GetParam().name + fn main(@builtin()" + GetParam().name +
R"() b : )" + GetParam().type + R"() { R"() b : )" + GetParam().type + R"() {
if (u32(vec4(b).x) == 0u) { if (u32(vec4(b).x) == 0u) {
dpdx(0.5); _ = dpdx(0.5);
} }
} }
)"; )";
@ -619,9 +619,9 @@ fn main(@builtin()" + GetParam().name +
RunTest(src, should_pass); RunTest(src, should_pass);
if (!should_pass) { if (!should_pass) {
EXPECT_EQ(error_, EXPECT_EQ(error_,
R"(test:5:5 error: 'dpdx' must only be called from uniform control flow R"(test:5:9 error: 'dpdx' must only be called from uniform control flow
dpdx(0.5); _ = dpdx(0.5);
^^^^ ^^^^^^^^^
test:4:3 note: control flow depends on possibly non-uniform value test:4:3 note: control flow depends on possibly non-uniform value
if (u32(vec4(b).x) == 0u) { if (u32(vec4(b).x) == 0u) {
@ -644,7 +644,7 @@ struct S {
@fragment @fragment
fn main(s : S) { fn main(s : S) {
if (u32(vec4(s.b).x) == 0u) { if (u32(vec4(s.b).x) == 0u) {
dpdx(0.5); _ = dpdx(0.5);
} }
} }
)"; )";
@ -653,9 +653,9 @@ fn main(s : S) {
RunTest(src, should_pass); RunTest(src, should_pass);
if (!should_pass) { if (!should_pass) {
EXPECT_EQ(error_, EXPECT_EQ(error_,
R"(test:9:5 error: 'dpdx' must only be called from uniform control flow R"(test:9:9 error: 'dpdx' must only be called from uniform control flow
dpdx(0.5); _ = dpdx(0.5);
^^^^ ^^^^^^^^^
test:8:3 note: control flow depends on possibly non-uniform value test:8:3 note: control flow depends on possibly non-uniform value
if (u32(vec4(s.b).x) == 0u) { if (u32(vec4(s.b).x) == 0u) {
@ -683,16 +683,16 @@ TEST_F(UniformityAnalysisTest, FragmentLocation) {
@fragment @fragment
fn main(@location(0) l : f32) { fn main(@location(0) l : f32) {
if (l == 0.0) { if (l == 0.0) {
dpdx(0.5); _ = dpdx(0.5);
} }
} }
)"; )";
RunTest(src, false); RunTest(src, false);
EXPECT_EQ(error_, EXPECT_EQ(error_,
R"(test:5:5 error: 'dpdx' must only be called from uniform control flow R"(test:5:9 error: 'dpdx' must only be called from uniform control flow
dpdx(0.5); _ = dpdx(0.5);
^^^^ ^^^^^^^^^
test:4:3 note: control flow depends on possibly non-uniform value test:4:3 note: control flow depends on possibly non-uniform value
if (l == 0.0) { if (l == 0.0) {
@ -713,16 +713,16 @@ struct S {
@fragment @fragment
fn main(s : S) { fn main(s : S) {
if (s.l == 0.0) { if (s.l == 0.0) {
dpdx(0.5); _ = dpdx(0.5);
} }
} }
)"; )";
RunTest(src, false); RunTest(src, false);
EXPECT_EQ(error_, EXPECT_EQ(error_,
R"(test:9:5 error: 'dpdx' must only be called from uniform control flow R"(test:9:9 error: 'dpdx' must only be called from uniform control flow
dpdx(0.5); _ = dpdx(0.5);
^^^^ ^^^^^^^^^
test:8:3 note: control flow depends on possibly non-uniform value test:8:3 note: control flow depends on possibly non-uniform value
if (s.l == 0.0) { if (s.l == 0.0) {
@ -7981,7 +7981,7 @@ TEST_F(UniformityAnalysisDiagnosticFilterTest, AttributeOnFunction_CalledByAnoth
@diagnostic(info, derivative_uniformity) @diagnostic(info, derivative_uniformity)
fn bar() { fn bar() {
dpdx(1.0); _ = dpdx(1.0);
} }
fn foo() { fn foo() {
@ -8002,7 +8002,7 @@ TEST_F(UniformityAnalysisDiagnosticFilterTest, AttributeOnFunction_RequirementOn
@diagnostic(info, derivative_uniformity) @diagnostic(info, derivative_uniformity)
fn bar(x : i32) { fn bar(x : i32) {
if (x == 0) { if (x == 0) {
dpdx(1.0); _ = dpdx(1.0);
} }
} }
@ -8022,7 +8022,7 @@ TEST_F(UniformityAnalysisDiagnosticFilterTest, AttributeOnFunction_BuiltinInChil
@group(0) @binding(0) var<storage, read_write> non_uniform : i32; @group(0) @binding(0) var<storage, read_write> non_uniform : i32;
fn bar() { fn bar() {
dpdx(1.0); _ = dpdx(1.0);
} }
@diagnostic(off, derivative_uniformity) @diagnostic(off, derivative_uniformity)
@ -8047,44 +8047,44 @@ diagnostic(info, derivative_uniformity);
fn a() { fn a() {
if (non_uniform == 42) { if (non_uniform == 42) {
dpdx(1.0); _ = dpdx(1.0);
} }
} }
@diagnostic(off, derivative_uniformity) @diagnostic(off, derivative_uniformity)
fn b() { fn b() {
if (non_uniform == 42) { if (non_uniform == 42) {
dpdx(1.0); _ = dpdx(1.0);
} }
} }
@diagnostic(info, derivative_uniformity) @diagnostic(info, derivative_uniformity)
fn c() { fn c() {
if (non_uniform == 42) { if (non_uniform == 42) {
dpdx(1.0); _ = dpdx(1.0);
} }
} }
@diagnostic(warning, derivative_uniformity) @diagnostic(warning, derivative_uniformity)
fn d() { fn d() {
if (non_uniform == 42) { if (non_uniform == 42) {
dpdx(1.0); _ = dpdx(1.0);
} }
} }
@diagnostic(error, derivative_uniformity) @diagnostic(error, derivative_uniformity)
fn e() { fn e() {
if (non_uniform == 42) { if (non_uniform == 42) {
dpdx(1.0); _ = dpdx(1.0);
} }
} }
)"; )";
RunTest(src, false); RunTest(src, false);
EXPECT_EQ(error_, EXPECT_EQ(error_,
R"(test:8:5 note: 'dpdx' must only be called from uniform control flow R"(test:8:9 note: 'dpdx' must only be called from uniform control flow
dpdx(1.0); _ = dpdx(1.0);
^^^^ ^^^^^^^^^
test:7:3 note: control flow depends on possibly non-uniform value test:7:3 note: control flow depends on possibly non-uniform value
if (non_uniform == 42) { if (non_uniform == 42) {
@ -8094,9 +8094,9 @@ test:7:7 note: reading from read_write storage buffer 'non_uniform' may result i
if (non_uniform == 42) { if (non_uniform == 42) {
^^^^^^^^^^^ ^^^^^^^^^^^
test:22:5 note: 'dpdx' must only be called from uniform control flow test:22:9 note: 'dpdx' must only be called from uniform control flow
dpdx(1.0); _ = dpdx(1.0);
^^^^ ^^^^^^^^^
test:21:3 note: control flow depends on possibly non-uniform value test:21:3 note: control flow depends on possibly non-uniform value
if (non_uniform == 42) { if (non_uniform == 42) {
@ -8106,9 +8106,9 @@ test:21:7 note: reading from read_write storage buffer 'non_uniform' may result
if (non_uniform == 42) { if (non_uniform == 42) {
^^^^^^^^^^^ ^^^^^^^^^^^
test:29:5 warning: 'dpdx' must only be called from uniform control flow test:29:9 warning: 'dpdx' must only be called from uniform control flow
dpdx(1.0); _ = dpdx(1.0);
^^^^ ^^^^^^^^^
test:28:3 note: control flow depends on possibly non-uniform value test:28:3 note: control flow depends on possibly non-uniform value
if (non_uniform == 42) { if (non_uniform == 42) {
@ -8118,9 +8118,9 @@ test:28:7 note: reading from read_write storage buffer 'non_uniform' may result
if (non_uniform == 42) { if (non_uniform == 42) {
^^^^^^^^^^^ ^^^^^^^^^^^
test:36:5 error: 'dpdx' must only be called from uniform control flow test:36:9 error: 'dpdx' must only be called from uniform control flow
dpdx(1.0); _ = dpdx(1.0);
^^^^ ^^^^^^^^^
test:35:3 note: control flow depends on possibly non-uniform value test:35:3 note: control flow depends on possibly non-uniform value
if (non_uniform == 42) { if (non_uniform == 42) {
@ -8215,7 +8215,7 @@ diagnostic(off, derivative_uniformity);
fn foo() { fn foo() {
if (non_uniform == 42) { if (non_uniform == 42) {
dpdx(1.0); _ = dpdx(1.0);
} }
} }

50
src/tint/resolver/validator.cc
View File

@ -1009,17 +1009,34 @@ bool Validator::Function(const sem::Function* func, ast::PipelineStage stage) co
auto* decl = func->Declaration(); auto* decl = func->Declaration();
for (auto* attr : decl->attributes) { for (auto* attr : decl->attributes) {
if (attr->Is<ast::WorkgroupAttribute>()) { bool ok = Switch(
attr, //
[&](const ast::WorkgroupAttribute*) {
if (decl->PipelineStage() != ast::PipelineStage::kCompute) { if (decl->PipelineStage() != ast::PipelineStage::kCompute) {
AddError("the workgroup_size attribute is only valid for compute stages", AddError("@workgroup_size is only valid for compute stages", attr->source);
return false;
}
return true;
},
[&](const ast::MustUseAttribute*) {
if (func->ReturnType()->Is<type::Void>()) {
AddError("@must_use can only be applied to functions that return a value",
attr->source); attr->source);
return false; return false;
} }
} else if (!attr->IsAnyOf<ast::DiagnosticAttribute, ast::StageAttribute, return true;
ast::MustUseAttribute, ast::InternalAttribute>()) { },
[&](Default) {
if (!attr->IsAnyOf<ast::DiagnosticAttribute, ast::StageAttribute,
ast::InternalAttribute>()) {
AddError("attribute is not valid for functions", attr->source); AddError("attribute is not valid for functions", attr->source);
return false; return false;
} }
return true;
});
if (!ok) {
return false;
}
} }
if (decl->params.Length() > 255) { if (decl->params.Length() > 255) {
@ -1467,22 +1484,39 @@ bool Validator::ContinueStatement(const sem::Statement* stmt,
} }
bool Validator::Call(const sem::Call* call, sem::Statement* current_statement) const { bool Validator::Call(const sem::Call* call, sem::Statement* current_statement) const {
if (!call->Target()->MustUse()) {
return true;
}
auto* expr = call->Declaration(); auto* expr = call->Declaration();
bool is_call_stmt = bool is_call_stmt =
current_statement && Is<ast::CallStatement>(current_statement->Declaration(), current_statement && Is<ast::CallStatement>(current_statement->Declaration(),
[&](auto* stmt) { return stmt->expr == expr; }); [&](auto* stmt) { return stmt->expr == expr; });
if (is_call_stmt) { if (is_call_stmt) {
return Switch( // Call target is annotated with @must_use, but was used as a call statement.
Switch(
call->Target(), // call->Target(), //
[&](const sem::Function* fn) {
AddError("ignoring return value of function '" +
symbols_.NameFor(fn->Declaration()->name->symbol) +
"' annotated with @must_use",
call->Declaration()->source);
sem_.NoteDeclarationSource(fn->Declaration());
},
[&](const sem::Builtin* b) {
AddError("ignoring return value of builtin '" + utils::ToString(b->Type()) + "'",
call->Declaration()->source);
},
[&](const sem::TypeConversion*) { [&](const sem::TypeConversion*) {
AddError("type conversion evaluated but not used", call->Declaration()->source); AddError("type conversion evaluated but not used", call->Declaration()->source);
return false;
}, },
[&](const sem::TypeInitializer*) { [&](const sem::TypeInitializer*) {
AddError("type initializer evaluated but not used", call->Declaration()->source); AddError("type initializer evaluated but not used", call->Declaration()->source);
return false;
}, },
[&](Default) { return true; }); [&](Default) {
AddError("return value of call not used", call->Declaration()->source);
});
return false;
} }
return true; return true;

5
src/tint/sem/builtin.cc
View File

@ -110,8 +110,9 @@ Builtin::Builtin(BuiltinType type,
utils::VectorRef<Parameter*> parameters, utils::VectorRef<Parameter*> parameters,
EvaluationStage eval_stage, EvaluationStage eval_stage,
PipelineStageSet supported_stages, PipelineStageSet supported_stages,
bool is_deprecated) bool is_deprecated,
: Base(return_type, SetOwner(std::move(parameters), this), eval_stage), bool must_use)
: Base(return_type, SetOwner(std::move(parameters), this), eval_stage, must_use),
type_(type), type_(type),
supported_stages_(supported_stages), supported_stages_(supported_stages),
is_deprecated_(is_deprecated) {} is_deprecated_(is_deprecated) {}

10
src/tint/sem/builtin.h
View File

@ -84,16 +84,16 @@ class Builtin final : public Castable<Builtin, CallTarget> {
/// @param return_type the return type for the builtin call /// @param return_type the return type for the builtin call
/// @param parameters the parameters for the builtin overload /// @param parameters the parameters for the builtin overload
/// @param eval_stage the earliest evaluation stage for a call to the builtin /// @param eval_stage the earliest evaluation stage for a call to the builtin
/// @param supported_stages the pipeline stages that this builtin can be /// @param supported_stages the pipeline stages that this builtin can be used in
/// used in /// @param is_deprecated true if the particular overload is considered deprecated
/// @param is_deprecated true if the particular overload is considered /// @param must_use true if the builtin was annotated with `@must_use`
/// deprecated
Builtin(BuiltinType type, Builtin(BuiltinType type,
const type::Type* return_type, const type::Type* return_type,
utils::VectorRef<Parameter*> parameters, utils::VectorRef<Parameter*> parameters,
EvaluationStage eval_stage, EvaluationStage eval_stage,
PipelineStageSet supported_stages, PipelineStageSet supported_stages,
bool is_deprecated); bool is_deprecated,
bool must_use);
/// Destructor /// Destructor
~Builtin() override; ~Builtin() override;

5
src/tint/sem/call_target.cc
View File

@ -25,8 +25,9 @@ namespace tint::sem {
CallTarget::CallTarget(const type::Type* return_type, CallTarget::CallTarget(const type::Type* return_type,
utils::VectorRef<const Parameter*> parameters, utils::VectorRef<const Parameter*> parameters,
EvaluationStage stage) EvaluationStage stage,
: signature_{return_type, std::move(parameters)}, stage_(stage) { bool must_use)
: signature_{return_type, std::move(parameters)}, stage_(stage), must_use_(must_use) {
TINT_ASSERT(Semantic, return_type); TINT_ASSERT(Semantic, return_type);
} }

12
src/tint/sem/call_target.h
View File

@ -67,12 +67,15 @@ struct CallTargetSignature {
class CallTarget : public Castable<CallTarget, Node> { class CallTarget : public Castable<CallTarget, Node> {
public: public:
/// Constructor /// Constructor
/// @param stage the earliest evaluation stage for a call to this target
/// @param return_type the return type of the call target /// @param return_type the return type of the call target
/// @param parameters the parameters for the call target /// @param parameters the parameters for the call target
/// @param stage the earliest evaluation stage for a call to this target
/// @param must_use the result of the call target must be used, i.e. it cannot be used as a call
/// statement.
CallTarget(const type::Type* return_type, CallTarget(const type::Type* return_type,
utils::VectorRef<const Parameter*> parameters, utils::VectorRef<const Parameter*> parameters,
EvaluationStage stage); EvaluationStage stage,
bool must_use);
/// Copy constructor /// Copy constructor
CallTarget(const CallTarget&); CallTarget(const CallTarget&);
@ -92,9 +95,14 @@ class CallTarget : public Castable<CallTarget, Node> {
/// @return the earliest evaluation stage for a call to this target /// @return the earliest evaluation stage for a call to this target
EvaluationStage Stage() const { return stage_; } EvaluationStage Stage() const { return stage_; }
/// @returns true if the result of the call target must be used, i.e. it cannot be used as a
/// call statement.
bool MustUse() const { return must_use_; }
private: private:
CallTargetSignature signature_; CallTargetSignature signature_;
EvaluationStage stage_; EvaluationStage stage_;
const bool must_use_;
}; };
} // namespace tint::sem } // namespace tint::sem

6
src/tint/sem/function.cc
View File

@ -16,6 +16,7 @@
#include "src/tint/ast/function.h" #include "src/tint/ast/function.h"
#include "src/tint/ast/identifier.h" #include "src/tint/ast/identifier.h"
#include "src/tint/ast/must_use_attribute.h"
#include "src/tint/sem/variable.h" #include "src/tint/sem/variable.h"
#include "src/tint/type/depth_texture.h" #include "src/tint/type/depth_texture.h"
#include "src/tint/type/external_texture.h" #include "src/tint/type/external_texture.h"
@ -43,7 +44,10 @@ Function::Function(const ast::Function* declaration,
type::Type* return_type, type::Type* return_type,
std::optional<uint32_t> return_location, std::optional<uint32_t> return_location,
utils::VectorRef<Parameter*> parameters) utils::VectorRef<Parameter*> parameters)
: Base(return_type, SetOwner(std::move(parameters), this), EvaluationStage::kRuntime), : Base(return_type,
SetOwner(std::move(parameters), this),
EvaluationStage::kRuntime,
ast::HasAttribute<ast::MustUseAttribute>(declaration->attributes)),
declaration_(declaration), declaration_(declaration),
workgroup_size_{1, 1, 1}, workgroup_size_{1, 1, 1},
return_location_(return_location) {} return_location_(return_location) {}

2
src/tint/sem/type_conversion.cc
View File

@ -21,7 +21,7 @@ namespace tint::sem {
TypeConversion::TypeConversion(const type::Type* type, TypeConversion::TypeConversion(const type::Type* type,
const sem::Parameter* parameter, const sem::Parameter* parameter,
EvaluationStage stage) EvaluationStage stage)
: Base(type, utils::Vector<const sem::Parameter*, 1>{parameter}, stage) {} : Base(type, utils::Vector<const sem::Parameter*, 1>{parameter}, stage, /* must_use */ true) {}
TypeConversion::~TypeConversion() = default; TypeConversion::~TypeConversion() = default;

2
src/tint/sem/type_initializer.cc
View File

@ -23,7 +23,7 @@ namespace tint::sem {
TypeInitializer::TypeInitializer(const type::Type* type, TypeInitializer::TypeInitializer(const type::Type* type,
utils::VectorRef<const Parameter*> parameters, utils::VectorRef<const Parameter*> parameters,
EvaluationStage stage) EvaluationStage stage)
: Base(type, std::move(parameters), stage) {} : Base(type, std::move(parameters), stage, /* must_use */ true) {}
TypeInitializer::~TypeInitializer() = default; TypeInitializer::~TypeInitializer() = default;

38
src/tint/transform/array_length_from_uniform_test.cc
View File

@ -514,43 +514,5 @@ struct SB {
got.data.Get<ArrayLengthFromUniform::Result>()->used_size_indices); got.data.Get<ArrayLengthFromUniform::Result>()->used_size_indices);
} }
TEST_F(ArrayLengthFromUniformTest, CallStatement) {
auto* src = R"(
struct SB {
arr : array<i32>,
}
@group(0) @binding(0) var<storage, read> a : SB;
@compute @workgroup_size(1)
fn main() {
arrayLength(&a.arr);
}
)";
auto* expect =
R"(
struct SB {
arr : array<i32>,
}
@group(0) @binding(0) var<storage, read> a : SB;
@compute @workgroup_size(1)
fn main() {
}
)";
ArrayLengthFromUniform::Config cfg({0, 30u});
cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{0, 0}, 0);
DataMap data;
data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src, data);
EXPECT_EQ(expect, str(got));
}
} // namespace } // namespace
} // namespace tint::transform } // namespace tint::transform

54
src/tint/transform/builtin_polyfill_test.cc
View File

@ -52,7 +52,7 @@ TEST_F(BuiltinPolyfillTest, ShouldRunAcosh) {
auto* src = R"( auto* src = R"(
fn f() { fn f() {
let v = 1.0; let v = 1.0;
acosh(v); _ = acosh(v);
} }
)"; )";
@ -183,7 +183,7 @@ TEST_F(BuiltinPolyfillTest, ShouldRunAsinh) {
auto* src = R"( auto* src = R"(
fn f() { fn f() {
let v = 1.0; let v = 1.0;
asinh(v); _ = asinh(v);
} }
)"; )";
@ -264,7 +264,7 @@ TEST_F(BuiltinPolyfillTest, ShouldRunAtanh) {
auto* src = R"( auto* src = R"(
fn f() { fn f() {
let v = 1.0; let v = 1.0;
atanh(v); _ = atanh(v);
} }
)"; )";
@ -666,7 +666,7 @@ TEST_F(BuiltinPolyfillTest, ShouldRunClampInteger_i32) {
auto* src = R"( auto* src = R"(
fn f() { fn f() {
let v = 1i; let v = 1i;
clamp(v, 2i, 3i); _ = clamp(v, 2i, 3i);
} }
)"; )";
@ -678,7 +678,7 @@ TEST_F(BuiltinPolyfillTest, ShouldRunClampInteger_u32) {
auto* src = R"( auto* src = R"(
fn f() { fn f() {
let v = 1u; let v = 1u;
clamp(v, 2u, 3u); _ = clamp(v, 2u, 3u);
} }
)"; )";
@ -690,7 +690,7 @@ TEST_F(BuiltinPolyfillTest, ShouldRunClampInteger_f32) {
auto* src = R"( auto* src = R"(
fn f() { fn f() {
let v = 1f; let v = 1f;
clamp(v, 2f, 3f); _ = clamp(v, 2f, 3f);
} }
)"; )";
@ -704,7 +704,7 @@ enable f16;
fn f() { fn f() {
let v = 1h; let v = 1h;
clamp(v, 2h, 3h); _ = clamp(v, 2h, 3h);
} }
)"; )";
@ -829,7 +829,7 @@ TEST_F(BuiltinPolyfillTest, ShouldRunCountLeadingZeros) {
auto* src = R"( auto* src = R"(
fn f() { fn f() {
let v = 15; let v = 15;
countLeadingZeros(v); _ = countLeadingZeros(v);
} }
)"; )";
@ -1002,7 +1002,7 @@ TEST_F(BuiltinPolyfillTest, ShouldRunCountTrailingZeros) {
auto* src = R"( auto* src = R"(
fn f() { fn f() {
let v = 15; let v = 15;
countTrailingZeros(v); _ = countTrailingZeros(v);
} }
)"; )";
@ -1175,7 +1175,7 @@ TEST_F(BuiltinPolyfillTest, ShouldRunExtractBits) {
auto* src = R"( auto* src = R"(
fn f() { fn f() {
let v = 1234i; let v = 1234i;
extractBits(v, 5u, 6u); _ = extractBits(v, 5u, 6u);
} }
)"; )";
@ -1430,7 +1430,7 @@ TEST_F(BuiltinPolyfillTest, ShouldRunFirstLeadingBit) {
auto* src = R"( auto* src = R"(
fn f() { fn f() {
let v = 15i; let v = 15i;
firstLeadingBit(v); _ = firstLeadingBit(v);
} }
)"; )";
@ -1603,7 +1603,7 @@ TEST_F(BuiltinPolyfillTest, ShouldRunFirstTrailingBit) {
auto* src = R"( auto* src = R"(
fn f() { fn f() {
let v = 15i; let v = 15i;
firstTrailingBit(v); _ = firstTrailingBit(v);
} }
)"; )";
@ -1776,7 +1776,7 @@ TEST_F(BuiltinPolyfillTest, ShouldRunInsertBits) {
auto* src = R"( auto* src = R"(
fn f() { fn f() {
let v = 1234i; let v = 1234i;
insertBits(v, 5678, 5u, 6u); _ = insertBits(v, 5678, 5u, 6u);
} }
)"; )";
@ -3059,7 +3059,7 @@ TEST_F(BuiltinPolyfillTest, ShouldRunSaturate) {
auto* src = R"( auto* src = R"(
fn f() { fn f() {
let v = 0.5f; let v = 0.5f;
saturate(v); _ = saturate(v);
} }
)"; )";
@ -3196,7 +3196,7 @@ TEST_F(BuiltinPolyfillTest, ShouldRunSign_i32) {
auto* src = R"( auto* src = R"(
fn f() { fn f() {
let v = 1i; let v = 1i;
sign(v); _ = sign(v);
} }
)"; )";
@ -3208,7 +3208,7 @@ TEST_F(BuiltinPolyfillTest, ShouldRunSign_f32) {
auto* src = R"( auto* src = R"(
fn f() { fn f() {
let v = 1f; let v = 1f;
sign(v); _ = sign(v);
} }
)"; )";
@ -3291,7 +3291,7 @@ TEST_F(BuiltinPolyfillTest, ShouldRunTextureSampleBaseClampToEdge_2d_f32) {
@group(0) @binding(1) var s : sampler; @group(0) @binding(1) var s : sampler;
fn f() { fn f() {
textureSampleBaseClampToEdge(t, s, vec2<f32>(0.5)); _ = textureSampleBaseClampToEdge(t, s, vec2<f32>(0.5));
} }
)"; )";
@ -3305,7 +3305,7 @@ TEST_F(BuiltinPolyfillTest, ShouldRunTextureSampleBaseClampToEdge_external) {
@group(0) @binding(1) var s : sampler; @group(0) @binding(1) var s : sampler;
fn f() { fn f() {
textureSampleBaseClampToEdge(t, s, vec2<f32>(0.5)); _ = textureSampleBaseClampToEdge(t, s, vec2<f32>(0.5));
} }
)"; )";
@ -3361,7 +3361,7 @@ TEST_F(BuiltinPolyfillTest, ShouldRunWorkgroupUniformLoad) {
var<workgroup> v : i32; var<workgroup> v : i32;
fn f() { fn f() {
workgroupUniformLoad(&v); _ = workgroupUniformLoad(&v);
} }
)"; )";
@ -3523,7 +3523,7 @@ TEST_F(BuiltinPolyfillTest, ShouldRunQuantizeToF16_Scalar) {
auto* src = R"( auto* src = R"(
fn f() { fn f() {
let v = 0.5; let v = 0.5;
quantizeToF16(0.5); _ = quantizeToF16(0.5);
} }
)"; )";
@ -3535,7 +3535,7 @@ TEST_F(BuiltinPolyfillTest, ShouldRunQuantizeToF16_Vector) {
auto* src = R"( auto* src = R"(
fn f() { fn f() {
let v = 0.5; let v = 0.5;
quantizeToF16(vec2(v)); _ = quantizeToF16(vec2(v));
} }
)"; )";
@ -3547,7 +3547,7 @@ TEST_F(BuiltinPolyfillTest, QuantizeToF16_Vec2) {
auto* src = R"( auto* src = R"(
fn f() { fn f() {
let v = 0.5; let v = 0.5;
quantizeToF16(vec2(v)); _ = quantizeToF16(vec2(v));
} }
)"; )";
@ -3558,7 +3558,7 @@ fn tint_quantizeToF16(v : vec2<f32>) -> vec2<f32> {
fn f() { fn f() {
let v = 0.5; let v = 0.5;
tint_quantizeToF16(vec2(v)); _ = tint_quantizeToF16(vec2(v));
} }
)"; )";
@ -3571,7 +3571,7 @@ TEST_F(BuiltinPolyfillTest, QuantizeToF16_Vec3) {
auto* src = R"( auto* src = R"(
fn f() { fn f() {
let v = 0.5; let v = 0.5;
quantizeToF16(vec3(v)); _ = quantizeToF16(vec3(v));
} }
)"; )";
@ -3582,7 +3582,7 @@ fn tint_quantizeToF16(v : vec3<f32>) -> vec3<f32> {
fn f() { fn f() {
let v = 0.5; let v = 0.5;
tint_quantizeToF16(vec3(v)); _ = tint_quantizeToF16(vec3(v));
} }
)"; )";
@ -3595,7 +3595,7 @@ TEST_F(BuiltinPolyfillTest, QuantizeToF16_Vec4) {
auto* src = R"( auto* src = R"(
fn f() { fn f() {
let v = 0.5; let v = 0.5;
quantizeToF16(vec4(v)); _ = quantizeToF16(vec4(v));
} }
)"; )";
@ -3606,7 +3606,7 @@ fn tint_quantizeToF16(v : vec4<f32>) -> vec4<f32> {
fn f() { fn f() {
let v = 0.5; let v = 0.5;
tint_quantizeToF16(vec4(v)); _ = tint_quantizeToF16(vec4(v));
} }
)"; )";

32
src/tint/transform/calculate_array_length_test.cc
View File

@ -547,37 +547,5 @@ struct SB2 {
EXPECT_EQ(expect, str(got)); EXPECT_EQ(expect, str(got));
} }
TEST_F(CalculateArrayLengthTest, CallStatement) {
auto* src = R"(
struct SB {
arr : array<i32>,
}
@group(0) @binding(0) var<storage, read> a : SB;
@compute @workgroup_size(1)
fn main() {
arrayLength(&a.arr);
}
)";
auto* expect =
R"(
struct SB {
arr : array<i32>,
}
@group(0) @binding(0) var<storage, read> a : SB;
@compute @workgroup_size(1)
fn main() {
}
)";
auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
EXPECT_EQ(expect, str(got));
}
} // namespace } // namespace
} // namespace tint::transform } // namespace tint::transform

44
src/tint/transform/multiplanar_external_texture_test.cc
View File

@ -887,7 +887,7 @@ fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
TEST_F(MultiplanarExternalTextureTest, ExternalTexturePassedAsParam) { TEST_F(MultiplanarExternalTextureTest, ExternalTexturePassedAsParam) {
auto* src = R"( auto* src = R"(
fn f(t : texture_external, s : sampler) { fn f(t : texture_external, s : sampler) {
textureSampleBaseClampToEdge(t, s, vec2<f32>(1.0, 2.0)); _ = textureSampleBaseClampToEdge(t, s, vec2<f32>(1.0, 2.0));
} }
@group(0) @binding(0) var ext_tex : texture_external; @group(0) @binding(0) var ext_tex : texture_external;
@ -955,7 +955,7 @@ fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp
} }
fn f(t : texture_2d<f32>, ext_tex_plane_1_1 : texture_2d<f32>, ext_tex_params_1 : ExternalTextureParams, s : sampler) { fn f(t : texture_2d<f32>, ext_tex_plane_1_1 : texture_2d<f32>, ext_tex_params_1 : ExternalTextureParams, s : sampler) {
textureSampleExternal(t, ext_tex_plane_1_1, s, vec2<f32>(1.0, 2.0), ext_tex_params_1); _ = textureSampleExternal(t, ext_tex_plane_1_1, s, vec2<f32>(1.0, 2.0), ext_tex_params_1);
} }
@group(0) @binding(0) var ext_tex : texture_2d<f32>; @group(0) @binding(0) var ext_tex : texture_2d<f32>;
@ -985,7 +985,7 @@ fn main() {
} }
fn f(t : texture_external, s : sampler) { fn f(t : texture_external, s : sampler) {
textureSampleBaseClampToEdge(t, s, vec2<f32>(1.0, 2.0)); _ = textureSampleBaseClampToEdge(t, s, vec2<f32>(1.0, 2.0));
} }
@group(0) @binding(0) var ext_tex : texture_external; @group(0) @binding(0) var ext_tex : texture_external;
@ -1053,7 +1053,7 @@ fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp
} }
fn f(t : texture_2d<f32>, ext_tex_plane_1_1 : texture_2d<f32>, ext_tex_params_1 : ExternalTextureParams, s : sampler) { fn f(t : texture_2d<f32>, ext_tex_plane_1_1 : texture_2d<f32>, ext_tex_params_1 : ExternalTextureParams, s : sampler) {
textureSampleExternal(t, ext_tex_plane_1_1, s, vec2<f32>(1.0, 2.0), ext_tex_params_1); _ = textureSampleExternal(t, ext_tex_plane_1_1, s, vec2<f32>(1.0, 2.0), ext_tex_params_1);
} }
@group(0) @binding(0) var ext_tex : texture_2d<f32>; @group(0) @binding(0) var ext_tex : texture_2d<f32>;
@ -1073,7 +1073,7 @@ fn f(t : texture_2d<f32>, ext_tex_plane_1_1 : texture_2d<f32>, ext_tex_params_1
TEST_F(MultiplanarExternalTextureTest, ExternalTexturePassedAsSecondParam) { TEST_F(MultiplanarExternalTextureTest, ExternalTexturePassedAsSecondParam) {
auto* src = R"( auto* src = R"(
fn f(s : sampler, t : texture_external) { fn f(s : sampler, t : texture_external) {
textureSampleBaseClampToEdge(t, s, vec2<f32>(1.0, 2.0)); _ = textureSampleBaseClampToEdge(t, s, vec2<f32>(1.0, 2.0));
} }
@group(0) @binding(0) var ext_tex : texture_external; @group(0) @binding(0) var ext_tex : texture_external;
@ -1141,7 +1141,7 @@ fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp
} }
fn f(s : sampler, t : texture_2d<f32>, ext_tex_plane_1_1 : texture_2d<f32>, ext_tex_params_1 : ExternalTextureParams) { fn f(s : sampler, t : texture_2d<f32>, ext_tex_plane_1_1 : texture_2d<f32>, ext_tex_params_1 : ExternalTextureParams) {
textureSampleExternal(t, ext_tex_plane_1_1, s, vec2<f32>(1.0, 2.0), ext_tex_params_1); _ = textureSampleExternal(t, ext_tex_plane_1_1, s, vec2<f32>(1.0, 2.0), ext_tex_params_1);
} }
@group(0) @binding(0) var ext_tex : texture_2d<f32>; @group(0) @binding(0) var ext_tex : texture_2d<f32>;
@ -1166,8 +1166,8 @@ fn main() {
TEST_F(MultiplanarExternalTextureTest, ExternalTexturePassedAsParamMultiple) { TEST_F(MultiplanarExternalTextureTest, ExternalTexturePassedAsParamMultiple) {
auto* src = R"( auto* src = R"(
fn f(t : texture_external, s : sampler, t2 : texture_external) { fn f(t : texture_external, s : sampler, t2 : texture_external) {
textureSampleBaseClampToEdge(t, s, vec2<f32>(1.0, 2.0)); _ = textureSampleBaseClampToEdge(t, s, vec2<f32>(1.0, 2.0));
textureSampleBaseClampToEdge(t2, s, vec2<f32>(1.0, 2.0)); _ = textureSampleBaseClampToEdge(t2, s, vec2<f32>(1.0, 2.0));
} }
@group(0) @binding(0) var ext_tex : texture_external; @group(0) @binding(0) var ext_tex : texture_external;
@ -1240,8 +1240,8 @@ fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp
} }
fn f(t : texture_2d<f32>, ext_tex_plane_1_2 : texture_2d<f32>, ext_tex_params_2 : ExternalTextureParams, s : sampler, t2 : texture_2d<f32>, ext_tex_plane_1_3 : texture_2d<f32>, ext_tex_params_3 : ExternalTextureParams) { fn f(t : texture_2d<f32>, ext_tex_plane_1_2 : texture_2d<f32>, ext_tex_params_2 : ExternalTextureParams, s : sampler, t2 : texture_2d<f32>, ext_tex_plane_1_3 : texture_2d<f32>, ext_tex_params_3 : ExternalTextureParams) {
textureSampleExternal(t, ext_tex_plane_1_2, s, vec2<f32>(1.0, 2.0), ext_tex_params_2); _ = textureSampleExternal(t, ext_tex_plane_1_2, s, vec2<f32>(1.0, 2.0), ext_tex_params_2);
textureSampleExternal(t2, ext_tex_plane_1_3, s, vec2<f32>(1.0, 2.0), ext_tex_params_3); _ = textureSampleExternal(t2, ext_tex_plane_1_3, s, vec2<f32>(1.0, 2.0), ext_tex_params_3);
} }
@group(0) @binding(0) var ext_tex : texture_2d<f32>; @group(0) @binding(0) var ext_tex : texture_2d<f32>;
@ -1274,8 +1274,8 @@ fn main() {
} }
fn f(t : texture_external, s : sampler, t2 : texture_external) { fn f(t : texture_external, s : sampler, t2 : texture_external) {
textureSampleBaseClampToEdge(t, s, vec2<f32>(1.0, 2.0)); _ = textureSampleBaseClampToEdge(t, s, vec2<f32>(1.0, 2.0));
textureSampleBaseClampToEdge(t2, s, vec2<f32>(1.0, 2.0)); _ = textureSampleBaseClampToEdge(t2, s, vec2<f32>(1.0, 2.0));
} }
@group(0) @binding(0) var ext_tex : texture_external; @group(0) @binding(0) var ext_tex : texture_external;
@ -1349,8 +1349,8 @@ fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp
} }
fn f(t : texture_2d<f32>, ext_tex_plane_1_2 : texture_2d<f32>, ext_tex_params_2 : ExternalTextureParams, s : sampler, t2 : texture_2d<f32>, ext_tex_plane_1_3 : texture_2d<f32>, ext_tex_params_3 : ExternalTextureParams) { fn f(t : texture_2d<f32>, ext_tex_plane_1_2 : texture_2d<f32>, ext_tex_params_2 : ExternalTextureParams, s : sampler, t2 : texture_2d<f32>, ext_tex_plane_1_3 : texture_2d<f32>, ext_tex_params_3 : ExternalTextureParams) {
textureSampleExternal(t, ext_tex_plane_1_2, s, vec2<f32>(1.0, 2.0), ext_tex_params_2); _ = textureSampleExternal(t, ext_tex_plane_1_2, s, vec2<f32>(1.0, 2.0), ext_tex_params_2);
textureSampleExternal(t2, ext_tex_plane_1_3, s, vec2<f32>(1.0, 2.0), ext_tex_params_3); _ = textureSampleExternal(t2, ext_tex_plane_1_3, s, vec2<f32>(1.0, 2.0), ext_tex_params_3);
} }
@group(0) @binding(0) var ext_tex : texture_2d<f32>; @group(0) @binding(0) var ext_tex : texture_2d<f32>;
@ -1373,7 +1373,7 @@ fn f(t : texture_2d<f32>, ext_tex_plane_1_2 : texture_2d<f32>, ext_tex_params_2
TEST_F(MultiplanarExternalTextureTest, ExternalTexturePassedAsParamNested) { TEST_F(MultiplanarExternalTextureTest, ExternalTexturePassedAsParamNested) {
auto* src = R"( auto* src = R"(
fn nested(t : texture_external, s : sampler) { fn nested(t : texture_external, s : sampler) {
textureSampleBaseClampToEdge(t, s, vec2<f32>(1.0, 2.0)); _ = textureSampleBaseClampToEdge(t, s, vec2<f32>(1.0, 2.0));
} }
fn f(t : texture_external, s : sampler) { fn f(t : texture_external, s : sampler) {
@ -1445,7 +1445,7 @@ fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp
} }
fn nested(t : texture_2d<f32>, ext_tex_plane_1_1 : texture_2d<f32>, ext_tex_params_1 : ExternalTextureParams, s : sampler) { fn nested(t : texture_2d<f32>, ext_tex_plane_1_1 : texture_2d<f32>, ext_tex_params_1 : ExternalTextureParams, s : sampler) {
textureSampleExternal(t, ext_tex_plane_1_1, s, vec2<f32>(1.0, 2.0), ext_tex_params_1); _ = textureSampleExternal(t, ext_tex_plane_1_1, s, vec2<f32>(1.0, 2.0), ext_tex_params_1);
} }
fn f(t : texture_2d<f32>, ext_tex_plane_1_2 : texture_2d<f32>, ext_tex_params_2 : ExternalTextureParams, s : sampler) { fn f(t : texture_2d<f32>, ext_tex_plane_1_2 : texture_2d<f32>, ext_tex_params_2 : ExternalTextureParams, s : sampler) {
@ -1474,7 +1474,7 @@ fn main() {
TEST_F(MultiplanarExternalTextureTest, ExternalTexturePassedAsParamNested_OutOfOrder) { TEST_F(MultiplanarExternalTextureTest, ExternalTexturePassedAsParamNested_OutOfOrder) {
auto* src = R"( auto* src = R"(
fn nested(t : texture_external, s : sampler) { fn nested(t : texture_external, s : sampler) {
textureSampleBaseClampToEdge(t, s, vec2<f32>(1.0, 2.0)); _ = textureSampleBaseClampToEdge(t, s, vec2<f32>(1.0, 2.0));
} }
fn f(t : texture_external, s : sampler) { fn f(t : texture_external, s : sampler) {
@ -1546,7 +1546,7 @@ fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp
} }
fn nested(t : texture_2d<f32>, ext_tex_plane_1_1 : texture_2d<f32>, ext_tex_params_1 : ExternalTextureParams, s : sampler) { fn nested(t : texture_2d<f32>, ext_tex_plane_1_1 : texture_2d<f32>, ext_tex_params_1 : ExternalTextureParams, s : sampler) {
textureSampleExternal(t, ext_tex_plane_1_1, s, vec2<f32>(1.0, 2.0), ext_tex_params_1); _ = textureSampleExternal(t, ext_tex_plane_1_1, s, vec2<f32>(1.0, 2.0), ext_tex_params_1);
} }
fn f(t : texture_2d<f32>, ext_tex_plane_1_2 : texture_2d<f32>, ext_tex_params_2 : ExternalTextureParams, s : sampler) { fn f(t : texture_2d<f32>, ext_tex_plane_1_2 : texture_2d<f32>, ext_tex_params_2 : ExternalTextureParams, s : sampler) {
@ -1619,7 +1619,7 @@ TEST_F(MultiplanarExternalTextureTest, ExternalTextureAlias) {
alias ET = texture_external; alias ET = texture_external;
fn f(t : ET, s : sampler) { fn f(t : ET, s : sampler) {
textureSampleBaseClampToEdge(t, s, vec2<f32>(1.0, 2.0)); _ = textureSampleBaseClampToEdge(t, s, vec2<f32>(1.0, 2.0));
} }
@group(0) @binding(0) var ext_tex : ET; @group(0) @binding(0) var ext_tex : ET;
@ -1689,7 +1689,7 @@ fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp
} }
fn f(t : texture_2d<f32>, ext_tex_plane_1_1 : texture_2d<f32>, ext_tex_params_1 : ExternalTextureParams, s : sampler) { fn f(t : texture_2d<f32>, ext_tex_plane_1_1 : texture_2d<f32>, ext_tex_params_1 : ExternalTextureParams, s : sampler) {
textureSampleExternal(t, ext_tex_plane_1_1, s, vec2<f32>(1.0, 2.0), ext_tex_params_1); _ = textureSampleExternal(t, ext_tex_plane_1_1, s, vec2<f32>(1.0, 2.0), ext_tex_params_1);
} }
@group(0) @binding(0) var ext_tex : texture_2d<f32>; @group(0) @binding(0) var ext_tex : texture_2d<f32>;
@ -1718,7 +1718,7 @@ fn main() {
} }
fn f(t : ET, s : sampler) { fn f(t : ET, s : sampler) {
textureSampleBaseClampToEdge(t, s, vec2<f32>(1.0, 2.0)); _ = textureSampleBaseClampToEdge(t, s, vec2<f32>(1.0, 2.0));
} }
@group(0) @binding(0) var ext_tex : ET; @group(0) @binding(0) var ext_tex : ET;
@ -1788,7 +1788,7 @@ fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp
} }
fn f(t : texture_2d<f32>, ext_tex_plane_1_1 : texture_2d<f32>, ext_tex_params_1 : ExternalTextureParams, s : sampler) { fn f(t : texture_2d<f32>, ext_tex_plane_1_1 : texture_2d<f32>, ext_tex_params_1 : ExternalTextureParams, s : sampler) {
textureSampleExternal(t, ext_tex_plane_1_1, s, vec2<f32>(1.0, 2.0), ext_tex_params_1); _ = textureSampleExternal(t, ext_tex_plane_1_1, s, vec2<f32>(1.0, 2.0), ext_tex_params_1);
} }
@group(0) @binding(0) var ext_tex : texture_2d<f32>; @group(0) @binding(0) var ext_tex : texture_2d<f32>;

2
src/tint/transform/remove_phonies_test.cc
View File

@ -67,8 +67,6 @@ fn f() {
_ = mat2x2<f32>(9.0, 10.0, 11.0, 12.0); _ = mat2x2<f32>(9.0, 10.0, 11.0, 12.0);
_ = atan2(1.0, 2.0); _ = atan2(1.0, 2.0);
_ = clamp(1.0, 2.0, 3.0); _ = clamp(1.0, 2.0, 3.0);
atan2(1.0, 2.0);
clamp(1.0, 2.0, 3.0);
} }
)"; )";

128
src/tint/transform/robustness_test.cc
View File

@ -792,14 +792,14 @@ fn idx_signed() {
var level_idx : i32; var level_idx : i32;
var sample_idx : i32; var sample_idx : i32;
textureLoad(tex_1d, 1i, level_idx); _ = textureLoad(tex_1d, 1i, level_idx);
textureLoad(tex_2d, vec2<i32>(1, 2), level_idx); _ = textureLoad(tex_2d, vec2<i32>(1, 2), level_idx);
textureLoad(tex_2d_arr, vec2<i32>(1, 2), array_idx, level_idx); _ = textureLoad(tex_2d_arr, vec2<i32>(1, 2), array_idx, level_idx);
textureLoad(tex_3d, vec3<i32>(1, 2, 3), level_idx); _ = textureLoad(tex_3d, vec3<i32>(1, 2, 3), level_idx);
textureLoad(tex_ms_2d, vec2<i32>(1, 2), sample_idx); _ = textureLoad(tex_ms_2d, vec2<i32>(1, 2), sample_idx);
textureLoad(tex_depth_2d, vec2<i32>(1, 2), level_idx); _ = textureLoad(tex_depth_2d, vec2<i32>(1, 2), level_idx);
textureLoad(tex_depth_2d_arr, vec2<i32>(1, 2), array_idx, level_idx); _ = textureLoad(tex_depth_2d_arr, vec2<i32>(1, 2), array_idx, level_idx);
textureLoad(tex_external, vec2<i32>(1, 2)); _ = textureLoad(tex_external, vec2<i32>(1, 2));
} }
fn idx_unsigned() { fn idx_unsigned() {
@ -807,14 +807,14 @@ fn idx_unsigned() {
var level_idx : u32; var level_idx : u32;
var sample_idx : u32; var sample_idx : u32;
textureLoad(tex_1d, 1u, level_idx); _ = textureLoad(tex_1d, 1u, level_idx);
textureLoad(tex_2d, vec2<u32>(1, 2), level_idx); _ = textureLoad(tex_2d, vec2<u32>(1, 2), level_idx);
textureLoad(tex_2d_arr, vec2<u32>(1, 2), array_idx, level_idx); _ = textureLoad(tex_2d_arr, vec2<u32>(1, 2), array_idx, level_idx);
textureLoad(tex_3d, vec3<u32>(1, 2, 3), level_idx); _ = textureLoad(tex_3d, vec3<u32>(1, 2, 3), level_idx);
textureLoad(tex_ms_2d, vec2<u32>(1, 2), sample_idx); _ = textureLoad(tex_ms_2d, vec2<u32>(1, 2), sample_idx);
textureLoad(tex_depth_2d, vec2<u32>(1, 2), level_idx); _ = textureLoad(tex_depth_2d, vec2<u32>(1, 2), level_idx);
textureLoad(tex_depth_2d_arr, vec2<u32>(1, 2), array_idx, level_idx); _ = textureLoad(tex_depth_2d_arr, vec2<u32>(1, 2), array_idx, level_idx);
textureLoad(tex_external, vec2<u32>(1, 2)); _ = textureLoad(tex_external, vec2<u32>(1, 2));
} }
)"; )";
@ -840,28 +840,28 @@ fn idx_signed() {
var array_idx : i32; var array_idx : i32;
var level_idx : i32; var level_idx : i32;
var sample_idx : i32; var sample_idx : i32;
textureLoad(tex_1d, clamp(1i, 0, i32((u32(textureDimensions(tex_1d, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_1d)) - 1))))) - 1))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_1d)) - 1)))); _ = textureLoad(tex_1d, clamp(1i, 0, i32((u32(textureDimensions(tex_1d, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_1d)) - 1))))) - 1))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_1d)) - 1))));
textureLoad(tex_2d, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_2d, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_2d)) - 1))))) - vec2(1)))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_2d)) - 1)))); _ = textureLoad(tex_2d, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_2d, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_2d)) - 1))))) - vec2(1)))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_2d)) - 1))));
textureLoad(tex_2d_arr, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_2d_arr, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_2d_arr)) - 1))))) - vec2(1)))), clamp(array_idx, 0, i32((u32(textureNumLayers(tex_2d_arr)) - 1))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_2d_arr)) - 1)))); _ = textureLoad(tex_2d_arr, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_2d_arr, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_2d_arr)) - 1))))) - vec2(1)))), clamp(array_idx, 0, i32((u32(textureNumLayers(tex_2d_arr)) - 1))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_2d_arr)) - 1))));
textureLoad(tex_3d, clamp(vec3<i32>(1, 2, 3), vec3(0), vec3<i32>((vec3<u32>(textureDimensions(tex_3d, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_3d)) - 1))))) - vec3(1)))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_3d)) - 1)))); _ = textureLoad(tex_3d, clamp(vec3<i32>(1, 2, 3), vec3(0), vec3<i32>((vec3<u32>(textureDimensions(tex_3d, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_3d)) - 1))))) - vec3(1)))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_3d)) - 1))));
textureLoad(tex_ms_2d, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_ms_2d)) - vec2(1)))), sample_idx); _ = textureLoad(tex_ms_2d, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_ms_2d)) - vec2(1)))), sample_idx);
textureLoad(tex_depth_2d, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_depth_2d, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_depth_2d)) - 1))))) - vec2(1)))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_depth_2d)) - 1)))); _ = textureLoad(tex_depth_2d, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_depth_2d, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_depth_2d)) - 1))))) - vec2(1)))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_depth_2d)) - 1))));
textureLoad(tex_depth_2d_arr, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_depth_2d_arr, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_depth_2d_arr)) - 1))))) - vec2(1)))), clamp(array_idx, 0, i32((u32(textureNumLayers(tex_depth_2d_arr)) - 1))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_depth_2d_arr)) - 1)))); _ = textureLoad(tex_depth_2d_arr, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_depth_2d_arr, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_depth_2d_arr)) - 1))))) - vec2(1)))), clamp(array_idx, 0, i32((u32(textureNumLayers(tex_depth_2d_arr)) - 1))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_depth_2d_arr)) - 1))));
textureLoad(tex_external, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_external)) - vec2(1))))); _ = textureLoad(tex_external, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_external)) - vec2(1)))));
} }
fn idx_unsigned() { fn idx_unsigned() {
var array_idx : u32; var array_idx : u32;
var level_idx : u32; var level_idx : u32;
var sample_idx : u32; var sample_idx : u32;
textureLoad(tex_1d, min(1u, (u32(textureDimensions(tex_1d, min(level_idx, (u32(textureNumLevels(tex_1d)) - 1)))) - 1)), min(level_idx, (u32(textureNumLevels(tex_1d)) - 1))); _ = textureLoad(tex_1d, min(1u, (u32(textureDimensions(tex_1d, min(level_idx, (u32(textureNumLevels(tex_1d)) - 1)))) - 1)), min(level_idx, (u32(textureNumLevels(tex_1d)) - 1)));
textureLoad(tex_2d, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_2d, min(level_idx, (u32(textureNumLevels(tex_2d)) - 1)))) - vec2(1))), min(level_idx, (u32(textureNumLevels(tex_2d)) - 1))); _ = textureLoad(tex_2d, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_2d, min(level_idx, (u32(textureNumLevels(tex_2d)) - 1)))) - vec2(1))), min(level_idx, (u32(textureNumLevels(tex_2d)) - 1)));
textureLoad(tex_2d_arr, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_2d_arr, min(level_idx, (u32(textureNumLevels(tex_2d_arr)) - 1)))) - vec2(1))), min(array_idx, (u32(textureNumLayers(tex_2d_arr)) - 1)), min(level_idx, (u32(textureNumLevels(tex_2d_arr)) - 1))); _ = textureLoad(tex_2d_arr, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_2d_arr, min(level_idx, (u32(textureNumLevels(tex_2d_arr)) - 1)))) - vec2(1))), min(array_idx, (u32(textureNumLayers(tex_2d_arr)) - 1)), min(level_idx, (u32(textureNumLevels(tex_2d_arr)) - 1)));
textureLoad(tex_3d, min(vec3<u32>(1, 2, 3), (vec3<u32>(textureDimensions(tex_3d, min(level_idx, (u32(textureNumLevels(tex_3d)) - 1)))) - vec3(1))), min(level_idx, (u32(textureNumLevels(tex_3d)) - 1))); _ = textureLoad(tex_3d, min(vec3<u32>(1, 2, 3), (vec3<u32>(textureDimensions(tex_3d, min(level_idx, (u32(textureNumLevels(tex_3d)) - 1)))) - vec3(1))), min(level_idx, (u32(textureNumLevels(tex_3d)) - 1)));
textureLoad(tex_ms_2d, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_ms_2d)) - vec2(1))), sample_idx); _ = textureLoad(tex_ms_2d, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_ms_2d)) - vec2(1))), sample_idx);
textureLoad(tex_depth_2d, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_depth_2d, min(level_idx, (u32(textureNumLevels(tex_depth_2d)) - 1)))) - vec2(1))), min(level_idx, (u32(textureNumLevels(tex_depth_2d)) - 1))); _ = textureLoad(tex_depth_2d, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_depth_2d, min(level_idx, (u32(textureNumLevels(tex_depth_2d)) - 1)))) - vec2(1))), min(level_idx, (u32(textureNumLevels(tex_depth_2d)) - 1)));
textureLoad(tex_depth_2d_arr, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_depth_2d_arr, min(level_idx, (u32(textureNumLevels(tex_depth_2d_arr)) - 1)))) - vec2(1))), min(array_idx, (u32(textureNumLayers(tex_depth_2d_arr)) - 1)), min(level_idx, (u32(textureNumLevels(tex_depth_2d_arr)) - 1))); _ = textureLoad(tex_depth_2d_arr, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_depth_2d_arr, min(level_idx, (u32(textureNumLevels(tex_depth_2d_arr)) - 1)))) - vec2(1))), min(array_idx, (u32(textureNumLayers(tex_depth_2d_arr)) - 1)), min(level_idx, (u32(textureNumLevels(tex_depth_2d_arr)) - 1)));
textureLoad(tex_external, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_external)) - vec2(1)))); _ = textureLoad(tex_external, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_external)) - vec2(1))));
} }
)"; )";
@ -878,14 +878,14 @@ fn idx_signed() {
var level_idx : i32; var level_idx : i32;
var sample_idx : i32; var sample_idx : i32;
textureLoad(tex_1d, 1i, level_idx); _ = textureLoad(tex_1d, 1i, level_idx);
textureLoad(tex_2d, vec2<i32>(1, 2), level_idx); _ = textureLoad(tex_2d, vec2<i32>(1, 2), level_idx);
textureLoad(tex_2d_arr, vec2<i32>(1, 2), array_idx, level_idx); _ = textureLoad(tex_2d_arr, vec2<i32>(1, 2), array_idx, level_idx);
textureLoad(tex_3d, vec3<i32>(1, 2, 3), level_idx); _ = textureLoad(tex_3d, vec3<i32>(1, 2, 3), level_idx);
textureLoad(tex_ms_2d, vec2<i32>(1, 2), sample_idx); _ = textureLoad(tex_ms_2d, vec2<i32>(1, 2), sample_idx);
textureLoad(tex_depth_2d, vec2<i32>(1, 2), level_idx); _ = textureLoad(tex_depth_2d, vec2<i32>(1, 2), level_idx);
textureLoad(tex_depth_2d_arr, vec2<i32>(1, 2), array_idx, level_idx); _ = textureLoad(tex_depth_2d_arr, vec2<i32>(1, 2), array_idx, level_idx);
textureLoad(tex_external, vec2<i32>(1, 2)); _ = textureLoad(tex_external, vec2<i32>(1, 2));
} }
fn idx_unsigned() { fn idx_unsigned() {
@ -893,14 +893,14 @@ fn idx_unsigned() {
var level_idx : u32; var level_idx : u32;
var sample_idx : u32; var sample_idx : u32;
textureLoad(tex_1d, 1u, level_idx); _ = textureLoad(tex_1d, 1u, level_idx);
textureLoad(tex_2d, vec2<u32>(1, 2), level_idx); _ = textureLoad(tex_2d, vec2<u32>(1, 2), level_idx);
textureLoad(tex_2d_arr, vec2<u32>(1, 2), array_idx, level_idx); _ = textureLoad(tex_2d_arr, vec2<u32>(1, 2), array_idx, level_idx);
textureLoad(tex_3d, vec3<u32>(1, 2, 3), level_idx); _ = textureLoad(tex_3d, vec3<u32>(1, 2, 3), level_idx);
textureLoad(tex_ms_2d, vec2<u32>(1, 2), sample_idx); _ = textureLoad(tex_ms_2d, vec2<u32>(1, 2), sample_idx);
textureLoad(tex_depth_2d, vec2<u32>(1, 2), level_idx); _ = textureLoad(tex_depth_2d, vec2<u32>(1, 2), level_idx);
textureLoad(tex_depth_2d_arr, vec2<u32>(1, 2), array_idx, level_idx); _ = textureLoad(tex_depth_2d_arr, vec2<u32>(1, 2), array_idx, level_idx);
textureLoad(tex_external, vec2<u32>(1, 2)); _ = textureLoad(tex_external, vec2<u32>(1, 2));
} }
@group(0) @binding(0) var tex_1d : texture_1d<f32>; @group(0) @binding(0) var tex_1d : texture_1d<f32>;
@ -919,28 +919,28 @@ fn idx_signed() {
var array_idx : i32; var array_idx : i32;
var level_idx : i32; var level_idx : i32;
var sample_idx : i32; var sample_idx : i32;
textureLoad(tex_1d, clamp(1i, 0, i32((u32(textureDimensions(tex_1d, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_1d)) - 1))))) - 1))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_1d)) - 1)))); _ = textureLoad(tex_1d, clamp(1i, 0, i32((u32(textureDimensions(tex_1d, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_1d)) - 1))))) - 1))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_1d)) - 1))));
textureLoad(tex_2d, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_2d, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_2d)) - 1))))) - vec2(1)))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_2d)) - 1)))); _ = textureLoad(tex_2d, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_2d, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_2d)) - 1))))) - vec2(1)))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_2d)) - 1))));
textureLoad(tex_2d_arr, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_2d_arr, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_2d_arr)) - 1))))) - vec2(1)))), clamp(array_idx, 0, i32((u32(textureNumLayers(tex_2d_arr)) - 1))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_2d_arr)) - 1)))); _ = textureLoad(tex_2d_arr, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_2d_arr, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_2d_arr)) - 1))))) - vec2(1)))), clamp(array_idx, 0, i32((u32(textureNumLayers(tex_2d_arr)) - 1))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_2d_arr)) - 1))));
textureLoad(tex_3d, clamp(vec3<i32>(1, 2, 3), vec3(0), vec3<i32>((vec3<u32>(textureDimensions(tex_3d, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_3d)) - 1))))) - vec3(1)))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_3d)) - 1)))); _ = textureLoad(tex_3d, clamp(vec3<i32>(1, 2, 3), vec3(0), vec3<i32>((vec3<u32>(textureDimensions(tex_3d, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_3d)) - 1))))) - vec3(1)))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_3d)) - 1))));
textureLoad(tex_ms_2d, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_ms_2d)) - vec2(1)))), sample_idx); _ = textureLoad(tex_ms_2d, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_ms_2d)) - vec2(1)))), sample_idx);
textureLoad(tex_depth_2d, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_depth_2d, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_depth_2d)) - 1))))) - vec2(1)))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_depth_2d)) - 1)))); _ = textureLoad(tex_depth_2d, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_depth_2d, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_depth_2d)) - 1))))) - vec2(1)))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_depth_2d)) - 1))));
textureLoad(tex_depth_2d_arr, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_depth_2d_arr, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_depth_2d_arr)) - 1))))) - vec2(1)))), clamp(array_idx, 0, i32((u32(textureNumLayers(tex_depth_2d_arr)) - 1))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_depth_2d_arr)) - 1)))); _ = textureLoad(tex_depth_2d_arr, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_depth_2d_arr, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_depth_2d_arr)) - 1))))) - vec2(1)))), clamp(array_idx, 0, i32((u32(textureNumLayers(tex_depth_2d_arr)) - 1))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_depth_2d_arr)) - 1))));
textureLoad(tex_external, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_external)) - vec2(1))))); _ = textureLoad(tex_external, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_external)) - vec2(1)))));
} }
fn idx_unsigned() { fn idx_unsigned() {
var array_idx : u32; var array_idx : u32;
var level_idx : u32; var level_idx : u32;
var sample_idx : u32; var sample_idx : u32;
textureLoad(tex_1d, min(1u, (u32(textureDimensions(tex_1d, min(level_idx, (u32(textureNumLevels(tex_1d)) - 1)))) - 1)), min(level_idx, (u32(textureNumLevels(tex_1d)) - 1))); _ = textureLoad(tex_1d, min(1u, (u32(textureDimensions(tex_1d, min(level_idx, (u32(textureNumLevels(tex_1d)) - 1)))) - 1)), min(level_idx, (u32(textureNumLevels(tex_1d)) - 1)));
textureLoad(tex_2d, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_2d, min(level_idx, (u32(textureNumLevels(tex_2d)) - 1)))) - vec2(1))), min(level_idx, (u32(textureNumLevels(tex_2d)) - 1))); _ = textureLoad(tex_2d, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_2d, min(level_idx, (u32(textureNumLevels(tex_2d)) - 1)))) - vec2(1))), min(level_idx, (u32(textureNumLevels(tex_2d)) - 1)));
textureLoad(tex_2d_arr, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_2d_arr, min(level_idx, (u32(textureNumLevels(tex_2d_arr)) - 1)))) - vec2(1))), min(array_idx, (u32(textureNumLayers(tex_2d_arr)) - 1)), min(level_idx, (u32(textureNumLevels(tex_2d_arr)) - 1))); _ = textureLoad(tex_2d_arr, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_2d_arr, min(level_idx, (u32(textureNumLevels(tex_2d_arr)) - 1)))) - vec2(1))), min(array_idx, (u32(textureNumLayers(tex_2d_arr)) - 1)), min(level_idx, (u32(textureNumLevels(tex_2d_arr)) - 1)));
textureLoad(tex_3d, min(vec3<u32>(1, 2, 3), (vec3<u32>(textureDimensions(tex_3d, min(level_idx, (u32(textureNumLevels(tex_3d)) - 1)))) - vec3(1))), min(level_idx, (u32(textureNumLevels(tex_3d)) - 1))); _ = textureLoad(tex_3d, min(vec3<u32>(1, 2, 3), (vec3<u32>(textureDimensions(tex_3d, min(level_idx, (u32(textureNumLevels(tex_3d)) - 1)))) - vec3(1))), min(level_idx, (u32(textureNumLevels(tex_3d)) - 1)));
textureLoad(tex_ms_2d, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_ms_2d)) - vec2(1))), sample_idx); _ = textureLoad(tex_ms_2d, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_ms_2d)) - vec2(1))), sample_idx);
textureLoad(tex_depth_2d, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_depth_2d, min(level_idx, (u32(textureNumLevels(tex_depth_2d)) - 1)))) - vec2(1))), min(level_idx, (u32(textureNumLevels(tex_depth_2d)) - 1))); _ = textureLoad(tex_depth_2d, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_depth_2d, min(level_idx, (u32(textureNumLevels(tex_depth_2d)) - 1)))) - vec2(1))), min(level_idx, (u32(textureNumLevels(tex_depth_2d)) - 1)));
textureLoad(tex_depth_2d_arr, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_depth_2d_arr, min(level_idx, (u32(textureNumLevels(tex_depth_2d_arr)) - 1)))) - vec2(1))), min(array_idx, (u32(textureNumLayers(tex_depth_2d_arr)) - 1)), min(level_idx, (u32(textureNumLevels(tex_depth_2d_arr)) - 1))); _ = textureLoad(tex_depth_2d_arr, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_depth_2d_arr, min(level_idx, (u32(textureNumLevels(tex_depth_2d_arr)) - 1)))) - vec2(1))), min(array_idx, (u32(textureNumLayers(tex_depth_2d_arr)) - 1)), min(level_idx, (u32(textureNumLevels(tex_depth_2d_arr)) - 1)));
textureLoad(tex_external, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_external)) - vec2(1)))); _ = textureLoad(tex_external, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_external)) - vec2(1))));
} }
@group(0) @binding(0) var tex_1d : texture_1d<f32>; @group(0) @binding(0) var tex_1d : texture_1d<f32>;

8
src/tint/transform/texture_1d_to_2d_test.cc
View File

@ -142,14 +142,14 @@ TEST_F(Texture1DTo2DTest, Global1DDeclAndTextureDimensionsInCallStmt) {
@group(0) @binding(0) var t : texture_1d<f32>; @group(0) @binding(0) var t : texture_1d<f32>;
fn main() { fn main() {
textureDimensions(t); _ = textureDimensions(t);
} }
)"; )";
auto* expect = R"( auto* expect = R"(
@group(0) @binding(0) var t : texture_2d<f32>; @group(0) @binding(0) var t : texture_2d<f32>;
fn main() { fn main() {
textureDimensions(t); _ = textureDimensions(t).x;
} }
)"; )";
@ -276,7 +276,7 @@ TEST_F(Texture1DTo2DTest, TextureAndNonTextureBuiltin) {
@group(0) @binding(0) var tex : texture_1d<i32>; @group(0) @binding(0) var tex : texture_1d<i32>;
fn d() { fn d() {
textureLoad(tex, 1, 0); _ = textureLoad(tex, 1, 0);
let l = sin(3.0); let l = sin(3.0);
} }
)"; )";
@ -285,7 +285,7 @@ fn d() {
@group(0) @binding(0) var tex : texture_2d<i32>; @group(0) @binding(0) var tex : texture_2d<i32>;
fn d() { fn d() {
textureLoad(tex, vec2<i32>(1, 0), 0); _ = textureLoad(tex, vec2<i32>(1, 0), 0);
let l = sin(3.0); let l = sin(3.0);
} }
)"; )";

60
src/tint/writer/glsl/generator_impl_builtin_test.cc
View File

@ -216,7 +216,7 @@ TEST_P(GlslBuiltinTest, Emit) {
ASSERT_NE(nullptr, call) << "Unhandled builtin"; ASSERT_NE(nullptr, call) << "Unhandled builtin";
Func("func", utils::Empty, ty.void_(), Func("func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(call), Assign(Phony(), call),
}, },
utils::Vector{create<ast::StageAttribute>(ast::PipelineStage::kFragment)}); utils::Vector{create<ast::StageAttribute>(ast::PipelineStage::kFragment)});
@ -345,7 +345,7 @@ TEST_F(GlslGeneratorImplTest_Builtin, Builtin_Call) {
GlobalVar("param1", ty.vec3<f32>(), builtin::AddressSpace::kPrivate); GlobalVar("param1", ty.vec3<f32>(), builtin::AddressSpace::kPrivate);
GlobalVar("param2", ty.vec3<f32>(), builtin::AddressSpace::kPrivate); GlobalVar("param2", ty.vec3<f32>(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
@ -359,7 +359,7 @@ TEST_F(GlslGeneratorImplTest_Builtin, Select_Scalar) {
GlobalVar("a", Expr(1_f), builtin::AddressSpace::kPrivate); GlobalVar("a", Expr(1_f), builtin::AddressSpace::kPrivate);
GlobalVar("b", Expr(2_f), builtin::AddressSpace::kPrivate); GlobalVar("b", Expr(2_f), builtin::AddressSpace::kPrivate);
auto* call = Call("select", "a", "b", true); auto* call = Call("select", "a", "b", true);
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
gen.increment_indent(); gen.increment_indent();
@ -372,7 +372,7 @@ TEST_F(GlslGeneratorImplTest_Builtin, Select_Vector) {
GlobalVar("a", vec2<i32>(1_i, 2_i), builtin::AddressSpace::kPrivate); GlobalVar("a", vec2<i32>(1_i, 2_i), builtin::AddressSpace::kPrivate);
GlobalVar("b", vec2<i32>(3_i, 4_i), builtin::AddressSpace::kPrivate); GlobalVar("b", vec2<i32>(3_i, 4_i), builtin::AddressSpace::kPrivate);
auto* call = Call("select", "a", "b", vec2<bool>(true, false)); auto* call = Call("select", "a", "b", vec2<bool>(true, false));
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
gen.increment_indent(); gen.increment_indent();
@ -388,7 +388,7 @@ TEST_F(GlslGeneratorImplTest_Builtin, FMA_f32) {
GlobalVar("b", ty.vec3<f32>(), builtin::AddressSpace::kPrivate); GlobalVar("b", ty.vec3<f32>(), builtin::AddressSpace::kPrivate);
GlobalVar("c", ty.vec3<f32>(), builtin::AddressSpace::kPrivate); GlobalVar("c", ty.vec3<f32>(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
@ -406,7 +406,7 @@ TEST_F(GlslGeneratorImplTest_Builtin, FMA_f16) {
GlobalVar("c", ty.vec3<f16>(), builtin::AddressSpace::kPrivate); GlobalVar("c", ty.vec3<f16>(), builtin::AddressSpace::kPrivate);
auto* call = Call("fma", "a", "b", "c"); auto* call = Call("fma", "a", "b", "c");
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
@ -1227,7 +1227,7 @@ void main() {
TEST_F(GlslGeneratorImplTest_Builtin, Pack4x8Snorm) { TEST_F(GlslGeneratorImplTest_Builtin, Pack4x8Snorm) {
auto* call = Call("pack4x8snorm", "p1"); auto* call = Call("pack4x8snorm", "p1");
GlobalVar("p1", ty.vec4<f32>(), builtin::AddressSpace::kPrivate); GlobalVar("p1", ty.vec4<f32>(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
ASSERT_TRUE(gen.Generate()) << gen.error(); ASSERT_TRUE(gen.Generate()) << gen.error();
@ -1236,7 +1236,7 @@ TEST_F(GlslGeneratorImplTest_Builtin, Pack4x8Snorm) {
vec4 p1 = vec4(0.0f, 0.0f, 0.0f, 0.0f); vec4 p1 = vec4(0.0f, 0.0f, 0.0f, 0.0f);
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void test_function() { void test_function() {
packSnorm4x8(p1); uint r = packSnorm4x8(p1);
return; return;
} }
)"); )");
@ -1245,7 +1245,7 @@ void test_function() {
TEST_F(GlslGeneratorImplTest_Builtin, Pack4x8Unorm) { TEST_F(GlslGeneratorImplTest_Builtin, Pack4x8Unorm) {
auto* call = Call("pack4x8unorm", "p1"); auto* call = Call("pack4x8unorm", "p1");
GlobalVar("p1", ty.vec4<f32>(), builtin::AddressSpace::kPrivate); GlobalVar("p1", ty.vec4<f32>(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
ASSERT_TRUE(gen.Generate()) << gen.error(); ASSERT_TRUE(gen.Generate()) << gen.error();
@ -1254,7 +1254,7 @@ TEST_F(GlslGeneratorImplTest_Builtin, Pack4x8Unorm) {
vec4 p1 = vec4(0.0f, 0.0f, 0.0f, 0.0f); vec4 p1 = vec4(0.0f, 0.0f, 0.0f, 0.0f);
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void test_function() { void test_function() {
packUnorm4x8(p1); uint r = packUnorm4x8(p1);
return; return;
} }
)"); )");
@ -1263,7 +1263,7 @@ void test_function() {
TEST_F(GlslGeneratorImplTest_Builtin, Pack2x16Snorm) { TEST_F(GlslGeneratorImplTest_Builtin, Pack2x16Snorm) {
auto* call = Call("pack2x16snorm", "p1"); auto* call = Call("pack2x16snorm", "p1");
GlobalVar("p1", ty.vec2<f32>(), builtin::AddressSpace::kPrivate); GlobalVar("p1", ty.vec2<f32>(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
ASSERT_TRUE(gen.Generate()) << gen.error(); ASSERT_TRUE(gen.Generate()) << gen.error();
@ -1272,7 +1272,7 @@ TEST_F(GlslGeneratorImplTest_Builtin, Pack2x16Snorm) {
vec2 p1 = vec2(0.0f, 0.0f); vec2 p1 = vec2(0.0f, 0.0f);
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void test_function() { void test_function() {
packSnorm2x16(p1); uint r = packSnorm2x16(p1);
return; return;
} }
)"); )");
@ -1281,7 +1281,7 @@ void test_function() {
TEST_F(GlslGeneratorImplTest_Builtin, Pack2x16Unorm) { TEST_F(GlslGeneratorImplTest_Builtin, Pack2x16Unorm) {
auto* call = Call("pack2x16unorm", "p1"); auto* call = Call("pack2x16unorm", "p1");
GlobalVar("p1", ty.vec2<f32>(), builtin::AddressSpace::kPrivate); GlobalVar("p1", ty.vec2<f32>(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
ASSERT_TRUE(gen.Generate()) << gen.error(); ASSERT_TRUE(gen.Generate()) << gen.error();
@ -1290,7 +1290,7 @@ TEST_F(GlslGeneratorImplTest_Builtin, Pack2x16Unorm) {
vec2 p1 = vec2(0.0f, 0.0f); vec2 p1 = vec2(0.0f, 0.0f);
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void test_function() { void test_function() {
packUnorm2x16(p1); uint r = packUnorm2x16(p1);
return; return;
} }
)"); )");
@ -1299,7 +1299,7 @@ void test_function() {
TEST_F(GlslGeneratorImplTest_Builtin, Pack2x16Float) { TEST_F(GlslGeneratorImplTest_Builtin, Pack2x16Float) {
auto* call = Call("pack2x16float", "p1"); auto* call = Call("pack2x16float", "p1");
GlobalVar("p1", ty.vec2<f32>(), builtin::AddressSpace::kPrivate); GlobalVar("p1", ty.vec2<f32>(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
ASSERT_TRUE(gen.Generate()) << gen.error(); ASSERT_TRUE(gen.Generate()) << gen.error();
@ -1308,7 +1308,7 @@ TEST_F(GlslGeneratorImplTest_Builtin, Pack2x16Float) {
vec2 p1 = vec2(0.0f, 0.0f); vec2 p1 = vec2(0.0f, 0.0f);
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void test_function() { void test_function() {
packHalf2x16(p1); uint r = packHalf2x16(p1);
return; return;
} }
)"); )");
@ -1317,7 +1317,7 @@ void test_function() {
TEST_F(GlslGeneratorImplTest_Builtin, Unpack4x8Snorm) { TEST_F(GlslGeneratorImplTest_Builtin, Unpack4x8Snorm) {
auto* call = Call("unpack4x8snorm", "p1"); auto* call = Call("unpack4x8snorm", "p1");
GlobalVar("p1", ty.u32(), builtin::AddressSpace::kPrivate); GlobalVar("p1", ty.u32(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
ASSERT_TRUE(gen.Generate()) << gen.error(); ASSERT_TRUE(gen.Generate()) << gen.error();
@ -1326,7 +1326,7 @@ TEST_F(GlslGeneratorImplTest_Builtin, Unpack4x8Snorm) {
uint p1 = 0u; uint p1 = 0u;
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void test_function() { void test_function() {
unpackSnorm4x8(p1); vec4 r = unpackSnorm4x8(p1);
return; return;
} }
)"); )");
@ -1335,7 +1335,7 @@ void test_function() {
TEST_F(GlslGeneratorImplTest_Builtin, Unpack4x8Unorm) { TEST_F(GlslGeneratorImplTest_Builtin, Unpack4x8Unorm) {
auto* call = Call("unpack4x8unorm", "p1"); auto* call = Call("unpack4x8unorm", "p1");
GlobalVar("p1", ty.u32(), builtin::AddressSpace::kPrivate); GlobalVar("p1", ty.u32(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
ASSERT_TRUE(gen.Generate()) << gen.error(); ASSERT_TRUE(gen.Generate()) << gen.error();
@ -1344,7 +1344,7 @@ TEST_F(GlslGeneratorImplTest_Builtin, Unpack4x8Unorm) {
uint p1 = 0u; uint p1 = 0u;
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void test_function() { void test_function() {
unpackUnorm4x8(p1); vec4 r = unpackUnorm4x8(p1);
return; return;
} }
)"); )");
@ -1353,7 +1353,7 @@ void test_function() {
TEST_F(GlslGeneratorImplTest_Builtin, Unpack2x16Snorm) { TEST_F(GlslGeneratorImplTest_Builtin, Unpack2x16Snorm) {
auto* call = Call("unpack2x16snorm", "p1"); auto* call = Call("unpack2x16snorm", "p1");
GlobalVar("p1", ty.u32(), builtin::AddressSpace::kPrivate); GlobalVar("p1", ty.u32(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
ASSERT_TRUE(gen.Generate()) << gen.error(); ASSERT_TRUE(gen.Generate()) << gen.error();
@ -1362,7 +1362,7 @@ TEST_F(GlslGeneratorImplTest_Builtin, Unpack2x16Snorm) {
uint p1 = 0u; uint p1 = 0u;
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void test_function() { void test_function() {
unpackSnorm2x16(p1); vec2 r = unpackSnorm2x16(p1);
return; return;
} }
)"); )");
@ -1371,7 +1371,7 @@ void test_function() {
TEST_F(GlslGeneratorImplTest_Builtin, Unpack2x16Unorm) { TEST_F(GlslGeneratorImplTest_Builtin, Unpack2x16Unorm) {
auto* call = Call("unpack2x16unorm", "p1"); auto* call = Call("unpack2x16unorm", "p1");
GlobalVar("p1", ty.u32(), builtin::AddressSpace::kPrivate); GlobalVar("p1", ty.u32(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
ASSERT_TRUE(gen.Generate()) << gen.error(); ASSERT_TRUE(gen.Generate()) << gen.error();
@ -1380,7 +1380,7 @@ TEST_F(GlslGeneratorImplTest_Builtin, Unpack2x16Unorm) {
uint p1 = 0u; uint p1 = 0u;
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void test_function() { void test_function() {
unpackUnorm2x16(p1); vec2 r = unpackUnorm2x16(p1);
return; return;
} }
)"); )");
@ -1389,7 +1389,7 @@ void test_function() {
TEST_F(GlslGeneratorImplTest_Builtin, Unpack2x16Float) { TEST_F(GlslGeneratorImplTest_Builtin, Unpack2x16Float) {
auto* call = Call("unpack2x16float", "p1"); auto* call = Call("unpack2x16float", "p1");
GlobalVar("p1", ty.u32(), builtin::AddressSpace::kPrivate); GlobalVar("p1", ty.u32(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
ASSERT_TRUE(gen.Generate()) << gen.error(); ASSERT_TRUE(gen.Generate()) << gen.error();
@ -1398,7 +1398,7 @@ TEST_F(GlslGeneratorImplTest_Builtin, Unpack2x16Float) {
uint p1 = 0u; uint p1 = 0u;
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void test_function() { void test_function() {
unpackHalf2x16(p1); vec2 r = unpackHalf2x16(p1);
return; return;
} }
)"); )");
@ -1452,7 +1452,7 @@ void main() {
TEST_F(GlslGeneratorImplTest_Builtin, DotI32) { TEST_F(GlslGeneratorImplTest_Builtin, DotI32) {
GlobalVar("v", ty.vec3<i32>(), builtin::AddressSpace::kPrivate); GlobalVar("v", ty.vec3<i32>(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(Call("dot", "v", "v"))); WrapInFunction(Decl(Var("r", Call("dot", "v", "v"))));
GeneratorImpl& gen = SanitizeAndBuild(); GeneratorImpl& gen = SanitizeAndBuild();
@ -1465,7 +1465,7 @@ int tint_int_dot(ivec3 a, ivec3 b) {
ivec3 v = ivec3(0, 0, 0); ivec3 v = ivec3(0, 0, 0);
void test_function() { void test_function() {
tint_int_dot(v, v); int r = tint_int_dot(v, v);
} }
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
@ -1478,7 +1478,7 @@ void main() {
TEST_F(GlslGeneratorImplTest_Builtin, DotU32) { TEST_F(GlslGeneratorImplTest_Builtin, DotU32) {
GlobalVar("v", ty.vec3<u32>(), builtin::AddressSpace::kPrivate); GlobalVar("v", ty.vec3<u32>(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(Call("dot", "v", "v"))); WrapInFunction(Decl(Var("r", Call("dot", "v", "v"))));
GeneratorImpl& gen = SanitizeAndBuild(); GeneratorImpl& gen = SanitizeAndBuild();
@ -1491,7 +1491,7 @@ uint tint_int_dot(uvec3 a, uvec3 b) {
uvec3 v = uvec3(0u, 0u, 0u); uvec3 v = uvec3(0u, 0u, 0u);
void test_function() { void test_function() {
tint_int_dot(v, v); uint r = tint_int_dot(v, v);
} }
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;

7
src/tint/writer/glsl/generator_impl_builtin_texture_test.cc
View File

@ -254,9 +254,9 @@ ExpectedResult expected_texture_overload(ast::builtin::test::ValidTextureOverloa
case ValidTextureOverload::kLoadMultisampled2dI32: case ValidTextureOverload::kLoadMultisampled2dI32:
return R"(texelFetch(Texture_1, ivec2(uvec2(1u, 2u)), int(3u));)"; return R"(texelFetch(Texture_1, ivec2(uvec2(1u, 2u)), int(3u));)";
case ValidTextureOverload::kLoadDepth2dLevelF32: case ValidTextureOverload::kLoadDepth2dLevelF32:
return R"(texelFetch(Texture_1, ivec2(1, 2), 3);)"; return R"(texelFetch(Texture_1, ivec2(1, 2), 3).x;)";
case ValidTextureOverload::kLoadDepth2dArrayLevelF32: case ValidTextureOverload::kLoadDepth2dArrayLevelF32:
return R"(texelFetch(Texture_1, ivec3(uvec3(1u, 2u, 3u)), int(4u));)"; return R"(texelFetch(Texture_1, ivec3(uvec3(1u, 2u, 3u)), int(4u)).x;)";
case ValidTextureOverload::kLoadDepthMultisampled2dF32: case ValidTextureOverload::kLoadDepthMultisampled2dF32:
return R"(texelFetch(Texture_1, ivec2(uvec2(1u, 2u)), int(3u)).x;)"; return R"(texelFetch(Texture_1, ivec2(uvec2(1u, 2u)), int(3u)).x;)";
case ValidTextureOverload::kStoreWO1dRgba32float: case ValidTextureOverload::kStoreWO1dRgba32float:
@ -281,7 +281,8 @@ TEST_P(GlslGeneratorBuiltinTextureTest, Call) {
param.BuildSamplerVariable(this); param.BuildSamplerVariable(this);
auto* call = Call(param.function, param.args(this)); auto* call = Call(param.function, param.args(this));
auto* stmt = CallStmt(call); auto* stmt = param.returns_value ? static_cast<const ast::Statement*>(Decl(Var("v", call)))
: static_cast<const ast::Statement*>(CallStmt(call));
Func("main", utils::Empty, ty.void_(), utils::Vector{stmt}, Func("main", utils::Empty, ty.void_(), utils::Vector{stmt},
utils::Vector{Stage(ast::PipelineStage::kFragment)}); utils::Vector{Stage(ast::PipelineStage::kFragment)});

8
src/tint/writer/glsl/generator_impl_type_test.cc
View File

@ -330,7 +330,7 @@ TEST_P(GlslDepthTexturesTest, Emit) {
Func("main", utils::Empty, ty.void_(), Func("main", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(Call("textureDimensions", "tex")), Decl(Var("v", Call("textureDimensions", "tex"))),
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
@ -358,7 +358,7 @@ TEST_F(GlslDepthMultisampledTexturesTest, Emit) {
Func("main", utils::Empty, ty.void_(), Func("main", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(Call("textureDimensions", "tex")), Decl(Var("v", Call("textureDimensions", "tex"))),
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
@ -402,7 +402,7 @@ TEST_P(GlslSampledTexturesTest, Emit) {
Func("main", utils::Empty, ty.void_(), Func("main", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(Call("textureDimensions", "tex")), Decl(Var("v", Call("textureDimensions", "tex"))),
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
@ -538,7 +538,7 @@ TEST_P(GlslStorageTexturesTest, Emit) {
Func("main", utils::Empty, ty.void_(), Func("main", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(Call("textureDimensions", "tex")), Decl(Var("v", Call("textureDimensions", "tex"))),
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),

48
src/tint/writer/hlsl/generator_impl_builtin_test.cc
View File

@ -215,7 +215,7 @@ TEST_P(HlslBuiltinTest, Emit) {
ASSERT_NE(nullptr, call) << "Unhandled builtin"; ASSERT_NE(nullptr, call) << "Unhandled builtin";
Func("func", utils::Empty, ty.void_(), Func("func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(call), Assign(Phony(), call),
}, },
utils::Vector{create<ast::StageAttribute>(ast::PipelineStage::kFragment)}); utils::Vector{create<ast::StageAttribute>(ast::PipelineStage::kFragment)});
@ -342,7 +342,7 @@ TEST_F(HlslGeneratorImplTest_Builtin, Builtin_Call) {
GlobalVar("param1", ty.vec3<f32>(), builtin::AddressSpace::kPrivate); GlobalVar("param1", ty.vec3<f32>(), builtin::AddressSpace::kPrivate);
GlobalVar("param2", ty.vec3<f32>(), builtin::AddressSpace::kPrivate); GlobalVar("param2", ty.vec3<f32>(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
@ -356,7 +356,7 @@ TEST_F(HlslGeneratorImplTest_Builtin, Select_Scalar) {
GlobalVar("a", Expr(1_f), builtin::AddressSpace::kPrivate); GlobalVar("a", Expr(1_f), builtin::AddressSpace::kPrivate);
GlobalVar("b", Expr(2_f), builtin::AddressSpace::kPrivate); GlobalVar("b", Expr(2_f), builtin::AddressSpace::kPrivate);
auto* call = Call("select", "a", "b", true); auto* call = Call("select", "a", "b", true);
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
gen.increment_indent(); gen.increment_indent();
@ -369,7 +369,7 @@ TEST_F(HlslGeneratorImplTest_Builtin, Select_Vector) {
GlobalVar("a", vec2<i32>(1_i, 2_i), builtin::AddressSpace::kPrivate); GlobalVar("a", vec2<i32>(1_i, 2_i), builtin::AddressSpace::kPrivate);
GlobalVar("b", vec2<i32>(3_i, 4_i), builtin::AddressSpace::kPrivate); GlobalVar("b", vec2<i32>(3_i, 4_i), builtin::AddressSpace::kPrivate);
auto* call = Call("select", "a", "b", vec2<bool>(true, false)); auto* call = Call("select", "a", "b", vec2<bool>(true, false));
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
gen.increment_indent(); gen.increment_indent();
@ -1087,7 +1087,7 @@ void test_function() {
TEST_F(HlslGeneratorImplTest_Builtin, Pack4x8Snorm) { TEST_F(HlslGeneratorImplTest_Builtin, Pack4x8Snorm) {
auto* call = Call("pack4x8snorm", "p1"); auto* call = Call("pack4x8snorm", "p1");
GlobalVar("p1", ty.vec4<f32>(), builtin::AddressSpace::kPrivate); GlobalVar("p1", ty.vec4<f32>(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
ASSERT_TRUE(gen.Generate()) << gen.error(); ASSERT_TRUE(gen.Generate()) << gen.error();
@ -1100,7 +1100,7 @@ static float4 p1 = float4(0.0f, 0.0f, 0.0f, 0.0f);
[numthreads(1, 1, 1)] [numthreads(1, 1, 1)]
void test_function() { void test_function() {
tint_pack4x8snorm(p1); uint r = tint_pack4x8snorm(p1);
return; return;
} }
)"); )");
@ -1109,7 +1109,7 @@ void test_function() {
TEST_F(HlslGeneratorImplTest_Builtin, Pack4x8Unorm) { TEST_F(HlslGeneratorImplTest_Builtin, Pack4x8Unorm) {
auto* call = Call("pack4x8unorm", "p1"); auto* call = Call("pack4x8unorm", "p1");
GlobalVar("p1", ty.vec4<f32>(), builtin::AddressSpace::kPrivate); GlobalVar("p1", ty.vec4<f32>(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
ASSERT_TRUE(gen.Generate()) << gen.error(); ASSERT_TRUE(gen.Generate()) << gen.error();
@ -1122,7 +1122,7 @@ static float4 p1 = float4(0.0f, 0.0f, 0.0f, 0.0f);
[numthreads(1, 1, 1)] [numthreads(1, 1, 1)]
void test_function() { void test_function() {
tint_pack4x8unorm(p1); uint r = tint_pack4x8unorm(p1);
return; return;
} }
)"); )");
@ -1131,7 +1131,7 @@ void test_function() {
TEST_F(HlslGeneratorImplTest_Builtin, Pack2x16Snorm) { TEST_F(HlslGeneratorImplTest_Builtin, Pack2x16Snorm) {
auto* call = Call("pack2x16snorm", "p1"); auto* call = Call("pack2x16snorm", "p1");
GlobalVar("p1", ty.vec2<f32>(), builtin::AddressSpace::kPrivate); GlobalVar("p1", ty.vec2<f32>(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
ASSERT_TRUE(gen.Generate()) << gen.error(); ASSERT_TRUE(gen.Generate()) << gen.error();
@ -1144,7 +1144,7 @@ static float2 p1 = float2(0.0f, 0.0f);
[numthreads(1, 1, 1)] [numthreads(1, 1, 1)]
void test_function() { void test_function() {
tint_pack2x16snorm(p1); uint r = tint_pack2x16snorm(p1);
return; return;
} }
)"); )");
@ -1153,7 +1153,7 @@ void test_function() {
TEST_F(HlslGeneratorImplTest_Builtin, Pack2x16Unorm) { TEST_F(HlslGeneratorImplTest_Builtin, Pack2x16Unorm) {
auto* call = Call("pack2x16unorm", "p1"); auto* call = Call("pack2x16unorm", "p1");
GlobalVar("p1", ty.vec2<f32>(), builtin::AddressSpace::kPrivate); GlobalVar("p1", ty.vec2<f32>(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
ASSERT_TRUE(gen.Generate()) << gen.error(); ASSERT_TRUE(gen.Generate()) << gen.error();
@ -1166,7 +1166,7 @@ static float2 p1 = float2(0.0f, 0.0f);
[numthreads(1, 1, 1)] [numthreads(1, 1, 1)]
void test_function() { void test_function() {
tint_pack2x16unorm(p1); uint r = tint_pack2x16unorm(p1);
return; return;
} }
)"); )");
@ -1175,7 +1175,7 @@ void test_function() {
TEST_F(HlslGeneratorImplTest_Builtin, Pack2x16Float) { TEST_F(HlslGeneratorImplTest_Builtin, Pack2x16Float) {
auto* call = Call("pack2x16float", "p1"); auto* call = Call("pack2x16float", "p1");
GlobalVar("p1", ty.vec2<f32>(), builtin::AddressSpace::kPrivate); GlobalVar("p1", ty.vec2<f32>(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
ASSERT_TRUE(gen.Generate()) << gen.error(); ASSERT_TRUE(gen.Generate()) << gen.error();
@ -1188,7 +1188,7 @@ static float2 p1 = float2(0.0f, 0.0f);
[numthreads(1, 1, 1)] [numthreads(1, 1, 1)]
void test_function() { void test_function() {
tint_pack2x16float(p1); uint r = tint_pack2x16float(p1);
return; return;
} }
)"); )");
@ -1197,7 +1197,7 @@ void test_function() {
TEST_F(HlslGeneratorImplTest_Builtin, Unpack4x8Snorm) { TEST_F(HlslGeneratorImplTest_Builtin, Unpack4x8Snorm) {
auto* call = Call("unpack4x8snorm", "p1"); auto* call = Call("unpack4x8snorm", "p1");
GlobalVar("p1", ty.u32(), builtin::AddressSpace::kPrivate); GlobalVar("p1", ty.u32(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
ASSERT_TRUE(gen.Generate()) << gen.error(); ASSERT_TRUE(gen.Generate()) << gen.error();
@ -1211,7 +1211,7 @@ static uint p1 = 0u;
[numthreads(1, 1, 1)] [numthreads(1, 1, 1)]
void test_function() { void test_function() {
tint_unpack4x8snorm(p1); float4 r = tint_unpack4x8snorm(p1);
return; return;
} }
)"); )");
@ -1220,7 +1220,7 @@ void test_function() {
TEST_F(HlslGeneratorImplTest_Builtin, Unpack4x8Unorm) { TEST_F(HlslGeneratorImplTest_Builtin, Unpack4x8Unorm) {
auto* call = Call("unpack4x8unorm", "p1"); auto* call = Call("unpack4x8unorm", "p1");
GlobalVar("p1", ty.u32(), builtin::AddressSpace::kPrivate); GlobalVar("p1", ty.u32(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
ASSERT_TRUE(gen.Generate()) << gen.error(); ASSERT_TRUE(gen.Generate()) << gen.error();
@ -1234,7 +1234,7 @@ static uint p1 = 0u;
[numthreads(1, 1, 1)] [numthreads(1, 1, 1)]
void test_function() { void test_function() {
tint_unpack4x8unorm(p1); float4 r = tint_unpack4x8unorm(p1);
return; return;
} }
)"); )");
@ -1243,7 +1243,7 @@ void test_function() {
TEST_F(HlslGeneratorImplTest_Builtin, Unpack2x16Snorm) { TEST_F(HlslGeneratorImplTest_Builtin, Unpack2x16Snorm) {
auto* call = Call("unpack2x16snorm", "p1"); auto* call = Call("unpack2x16snorm", "p1");
GlobalVar("p1", ty.u32(), builtin::AddressSpace::kPrivate); GlobalVar("p1", ty.u32(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
ASSERT_TRUE(gen.Generate()) << gen.error(); ASSERT_TRUE(gen.Generate()) << gen.error();
@ -1257,7 +1257,7 @@ static uint p1 = 0u;
[numthreads(1, 1, 1)] [numthreads(1, 1, 1)]
void test_function() { void test_function() {
tint_unpack2x16snorm(p1); float2 r = tint_unpack2x16snorm(p1);
return; return;
} }
)"); )");
@ -1266,7 +1266,7 @@ void test_function() {
TEST_F(HlslGeneratorImplTest_Builtin, Unpack2x16Unorm) { TEST_F(HlslGeneratorImplTest_Builtin, Unpack2x16Unorm) {
auto* call = Call("unpack2x16unorm", "p1"); auto* call = Call("unpack2x16unorm", "p1");
GlobalVar("p1", ty.u32(), builtin::AddressSpace::kPrivate); GlobalVar("p1", ty.u32(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
ASSERT_TRUE(gen.Generate()) << gen.error(); ASSERT_TRUE(gen.Generate()) << gen.error();
@ -1280,7 +1280,7 @@ static uint p1 = 0u;
[numthreads(1, 1, 1)] [numthreads(1, 1, 1)]
void test_function() { void test_function() {
tint_unpack2x16unorm(p1); float2 r = tint_unpack2x16unorm(p1);
return; return;
} }
)"); )");
@ -1289,7 +1289,7 @@ void test_function() {
TEST_F(HlslGeneratorImplTest_Builtin, Unpack2x16Float) { TEST_F(HlslGeneratorImplTest_Builtin, Unpack2x16Float) {
auto* call = Call("unpack2x16float", "p1"); auto* call = Call("unpack2x16float", "p1");
GlobalVar("p1", ty.u32(), builtin::AddressSpace::kPrivate); GlobalVar("p1", ty.u32(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
ASSERT_TRUE(gen.Generate()) << gen.error(); ASSERT_TRUE(gen.Generate()) << gen.error();
@ -1302,7 +1302,7 @@ static uint p1 = 0u;
[numthreads(1, 1, 1)] [numthreads(1, 1, 1)]
void test_function() { void test_function() {
tint_unpack2x16float(p1); float2 r = tint_unpack2x16float(p1);
return; return;
} }
)"); )");

3
src/tint/writer/hlsl/generator_impl_builtin_texture_test.cc
View File

@ -373,7 +373,8 @@ TEST_P(HlslGeneratorBuiltinTextureTest, Call) {
param.BuildSamplerVariable(this); param.BuildSamplerVariable(this);
auto* call = Call(param.function, param.args(this)); auto* call = Call(param.function, param.args(this));
auto* stmt = CallStmt(call); auto* stmt = param.returns_value ? static_cast<const ast::Statement*>(Decl(Var("v", call)))
: static_cast<const ast::Statement*>(CallStmt(call));
Func("main", utils::Empty, ty.void_(), utils::Vector{stmt}, Func("main", utils::Empty, ty.void_(), utils::Vector{stmt},
utils::Vector{Stage(ast::PipelineStage::kFragment)}); utils::Vector{Stage(ast::PipelineStage::kFragment)});

8
src/tint/writer/hlsl/generator_impl_type_test.cc
View File

@ -325,7 +325,7 @@ TEST_P(HlslDepthTexturesTest, Emit) {
Func("main", utils::Empty, ty.void_(), Func("main", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(Call("textureDimensions", "tex")), Decl(Var("v", Call("textureDimensions", "tex"))),
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
@ -356,7 +356,7 @@ TEST_F(HlslDepthMultisampledTexturesTest, Emit) {
Func("main", utils::Empty, ty.void_(), Func("main", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(Call("textureDimensions", "tex")), Decl(Var("v", Call("textureDimensions", "tex"))),
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
@ -400,7 +400,7 @@ TEST_P(HlslSampledTexturesTest, Emit) {
Func("main", utils::Empty, ty.void_(), Func("main", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(Call("textureDimensions", "tex")), Decl(Var("v", Call("textureDimensions", "tex"))),
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
@ -541,7 +541,7 @@ TEST_P(HlslStorageTexturesTest, Emit) {
Func("main", utils::Empty, ty.void_(), Func("main", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(Call("textureDimensions", "tex")), Decl(Var("v", Call("textureDimensions", "tex"))),
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),

10
src/tint/writer/msl/generator_impl_builtin_test.cc
View File

@ -374,7 +374,7 @@ TEST_F(MslGeneratorImplTest, Builtin_Call) {
GlobalVar("param2", ty.vec2<f32>(), builtin::AddressSpace::kPrivate); GlobalVar("param2", ty.vec2<f32>(), builtin::AddressSpace::kPrivate);
auto* call = Call("dot", "param1", "param2"); auto* call = Call("dot", "param1", "param2");
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
@ -1048,7 +1048,7 @@ kernel void test_function() {
TEST_F(MslGeneratorImplTest, Pack2x16Float) { TEST_F(MslGeneratorImplTest, Pack2x16Float) {
auto* call = Call("pack2x16float", "p1"); auto* call = Call("pack2x16float", "p1");
GlobalVar("p1", ty.vec2<f32>(), builtin::AddressSpace::kPrivate); GlobalVar("p1", ty.vec2<f32>(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
@ -1060,7 +1060,7 @@ TEST_F(MslGeneratorImplTest, Pack2x16Float) {
TEST_F(MslGeneratorImplTest, Unpack2x16Float) { TEST_F(MslGeneratorImplTest, Unpack2x16Float) {
auto* call = Call("unpack2x16float", "p1"); auto* call = Call("unpack2x16float", "p1");
GlobalVar("p1", ty.u32(), builtin::AddressSpace::kPrivate); GlobalVar("p1", ty.u32(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(call)); WrapInFunction(Decl(Var("r", call)));
GeneratorImpl& gen = Build(); GeneratorImpl& gen = Build();
@ -1071,7 +1071,7 @@ TEST_F(MslGeneratorImplTest, Unpack2x16Float) {
TEST_F(MslGeneratorImplTest, DotI32) { TEST_F(MslGeneratorImplTest, DotI32) {
GlobalVar("v", ty.vec3<i32>(), builtin::AddressSpace::kPrivate); GlobalVar("v", ty.vec3<i32>(), builtin::AddressSpace::kPrivate);
WrapInFunction(CallStmt(Call("dot", "v", "v"))); WrapInFunction(Decl(Var("r", Call("dot", "v", "v"))));
GeneratorImpl& gen = SanitizeAndBuild(); GeneratorImpl& gen = SanitizeAndBuild();
@ -1086,7 +1086,7 @@ T tint_dot3(vec<T,3> a, vec<T,3> b) {
} }
kernel void test_function() { kernel void test_function() {
thread int3 tint_symbol = 0; thread int3 tint_symbol = 0;
tint_dot3(tint_symbol, tint_symbol); int r = tint_dot3(tint_symbol, tint_symbol);
return; return;
} }

3
src/tint/writer/msl/generator_impl_builtin_texture_test.cc
View File

@ -277,7 +277,8 @@ TEST_P(MslGeneratorBuiltinTextureTest, Call) {
param.BuildSamplerVariable(this); param.BuildSamplerVariable(this);
auto* call = Call(param.function, param.args(this)); auto* call = Call(param.function, param.args(this));
auto* stmt = CallStmt(call); auto* stmt = param.returns_value ? static_cast<const ast::Statement*>(Assign(Phony(), call))
: static_cast<const ast::Statement*>(CallStmt(call));
Func("main", utils::Empty, ty.void_(), utils::Vector{stmt}, Func("main", utils::Empty, ty.void_(), utils::Vector{stmt},
utils::Vector{Stage(ast::PipelineStage::kFragment)}); utils::Vector{Stage(ast::PipelineStage::kFragment)});

325
src/tint/writer/spirv/builder_builtin_test.cc
View File

@ -52,12 +52,12 @@ TEST_F(BuiltinBuilderTest, Call_TextureSampleCompare_Twice) {
Func("f1", utils::Empty, ty.void_(), Func("f1", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(expr1), Decl(Let("l", expr1)),
}, },
utils::Empty); utils::Empty);
Func("f2", utils::Empty, ty.void_(), Func("f2", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(expr2), Decl(Let("l", expr2)),
}, },
utils::Empty); utils::Empty);
@ -102,7 +102,7 @@ TEST_F(BuiltinBuilderTest, Call_GLSLMethod_WithLoad_f32) {
auto* expr = Call("round", "ident"); auto* expr = Call("round", "ident");
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -138,7 +138,7 @@ TEST_F(BuiltinBuilderTest, Call_GLSLMethod_WithLoad_f16) {
auto* expr = Call("round", "ident"); auto* expr = Call("round", "ident");
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -177,7 +177,7 @@ TEST_P(BuiltinBoolTest, Call_Bool_Scalar) {
auto* expr = Call(param.name, "v"); auto* expr = Call(param.name, "v");
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -203,7 +203,7 @@ TEST_P(BuiltinBoolTest, Call_Bool_Vector) {
auto* expr = Call(param.name, "v"); auto* expr = Call(param.name, "v");
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -238,7 +238,7 @@ TEST_F(BuiltinBuilderTest, Call_Select) {
auto* expr = Call("select", "v3", "v3", "bool_v3"); auto* expr = Call("select", "v3", "v3", "bool_v3");
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -284,7 +284,7 @@ TEST_F(BuiltinBuilderTest, Call_ArrayLength) {
Func("a_func", utils::Empty, ty.void_(), Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(expr), Decl(Let("l", expr)),
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
@ -328,7 +328,7 @@ TEST_F(BuiltinBuilderTest, Call_ArrayLength_OtherMembersInStruct) {
Func("a_func", utils::Empty, ty.void_(), Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(expr), Decl(Let("l", expr)),
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
@ -376,7 +376,7 @@ TEST_F(BuiltinBuilderTest, Call_ArrayLength_ViaLets) {
utils::Vector{ utils::Vector{
Decl(p), Decl(p),
Decl(p2), Decl(p2),
CallStmt(expr), Decl(Let("l", expr)),
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
@ -437,7 +437,7 @@ TEST_F(BuiltinBuilderTest, Call_ArrayLength_ViaLets_WithPtrNoise) {
Decl(p), Decl(p),
Decl(p2), Decl(p2),
Decl(p3), Decl(p3),
CallStmt(expr), Decl(Let("l", expr)),
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
@ -484,7 +484,7 @@ TEST_P(Builtin_Builder_SingleParam_Float_Test, Call_Scalar_f32) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(scalar), Decl(scalar),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -525,7 +525,7 @@ TEST_P(Builtin_Builder_SingleParam_Float_Test, Call_Scalar_f16) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(scalar), Decl(scalar),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -565,7 +565,7 @@ TEST_P(Builtin_Builder_SingleParam_Float_Test, Call_Vector_f32) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(vec), Decl(vec),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -609,7 +609,7 @@ TEST_P(Builtin_Builder_SingleParam_Float_Test, Call_Vector_f16) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(vec), Decl(vec),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -675,7 +675,7 @@ TEST_F(BuiltinBuilderTest, Call_Length_Scalar_f32) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(scalar), Decl(scalar),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -712,7 +712,7 @@ TEST_F(BuiltinBuilderTest, Call_Length_Scalar_f16) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(scalar), Decl(scalar),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -747,7 +747,7 @@ TEST_F(BuiltinBuilderTest, Call_Length_Vector_f32) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(vec), Decl(vec),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -786,7 +786,7 @@ TEST_F(BuiltinBuilderTest, Call_Length_Vector_f16) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(vec), Decl(vec),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -823,7 +823,7 @@ TEST_F(BuiltinBuilderTest, Call_Normalize_f32) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(vec), Decl(vec),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -862,7 +862,7 @@ TEST_F(BuiltinBuilderTest, Call_Normalize_f16) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(vec), Decl(vec),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -901,7 +901,7 @@ TEST_P(Builtin_Builder_DualParam_Float_Test, Call_Scalar_f32) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(scalar), Decl(scalar),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -942,7 +942,7 @@ TEST_P(Builtin_Builder_DualParam_Float_Test, Call_Scalar_f16) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(scalar), Decl(scalar),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -981,7 +981,7 @@ TEST_P(Builtin_Builder_DualParam_Float_Test, Call_Vector_f32) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(vec), Decl(vec),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -1024,7 +1024,7 @@ TEST_P(Builtin_Builder_DualParam_Float_Test, Call_Vector_f16) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(vec), Decl(vec),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -1072,7 +1072,7 @@ TEST_F(BuiltinBuilderTest, Call_Reflect_Vector_f32) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(vec), Decl(vec),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -1112,7 +1112,7 @@ TEST_F(BuiltinBuilderTest, Call_Reflect_Vector_f16) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(vec), Decl(vec),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -1150,7 +1150,7 @@ TEST_F(BuiltinBuilderTest, Call_Distance_Scalar_f32) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(scalar), Decl(scalar),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -1188,7 +1188,7 @@ TEST_F(BuiltinBuilderTest, Call_Distance_Scalar_f16) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(scalar), Decl(scalar),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -1224,7 +1224,7 @@ TEST_F(BuiltinBuilderTest, Call_Distance_Vector_f32) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(vec), Decl(vec),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -1264,7 +1264,7 @@ TEST_F(BuiltinBuilderTest, Call_Distance_Vector_f16) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(vec), Decl(vec),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -1302,7 +1302,7 @@ TEST_F(BuiltinBuilderTest, Call_Cross_f32) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(vec), Decl(vec),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -1342,7 +1342,7 @@ TEST_F(BuiltinBuilderTest, Call_Cross_f16) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(vec), Decl(vec),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -1382,7 +1382,7 @@ TEST_P(Builtin_Builder_ThreeParam_Float_Test, Call_Scalar_f32) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(scalar), Decl(scalar),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -1424,7 +1424,7 @@ TEST_P(Builtin_Builder_ThreeParam_Float_Test, Call_Scalar_f16) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(scalar), Decl(scalar),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -1464,7 +1464,7 @@ TEST_P(Builtin_Builder_ThreeParam_Float_Test, Call_Vector_f32) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(vec), Decl(vec),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -1508,7 +1508,7 @@ TEST_P(Builtin_Builder_ThreeParam_Float_Test, Call_Vector_f16) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(vec), Decl(vec),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -1557,7 +1557,7 @@ TEST_F(BuiltinBuilderTest, Call_FaceForward_Vector_f32) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(vec), Decl(vec),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -1598,7 +1598,7 @@ TEST_F(BuiltinBuilderTest, Call_FaceForward_Vector_f16) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(vec), Decl(vec),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -1637,7 +1637,7 @@ TEST_F(BuiltinBuilderTest, Runtime_Call_Modf_f32) {
Func("a_func", utils::Empty, ty.void_(), Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(vec), Decl(vec),
CallStmt(expr), Decl(Let("l", expr)),
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
@ -1691,7 +1691,7 @@ TEST_F(BuiltinBuilderTest, Runtime_Call_Modf_f16) {
Func("a_func", utils::Empty, ty.void_(), Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(vec), Decl(vec),
CallStmt(expr), Decl(Let("l", expr)),
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
@ -1745,7 +1745,7 @@ TEST_F(BuiltinBuilderTest, Const_Call_Modf_f32) {
auto* expr = Call("modf", vec2<f32>(1_f, 2_f)); auto* expr = Call("modf", vec2<f32>(1_f, 2_f));
Func("a_func", utils::Empty, ty.void_(), Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(expr), Decl(Let("l", expr)),
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
@ -1756,27 +1756,27 @@ TEST_F(BuiltinBuilderTest, Const_Call_Modf_f32) {
ASSERT_TRUE(b.Build()) << b.error(); ASSERT_TRUE(b.Build()) << b.error();
auto got = DumpBuilder(b); auto got = DumpBuilder(b);
auto* expect = R"(OpCapability Shader auto* expect = R"(OpCapability Shader
%9 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450 OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %3 "a_func" OpEntryPoint Fragment %3 "a_func"
OpExecutionMode %3 OriginUpperLeft OpExecutionMode %3 OriginUpperLeft
OpName %3 "a_func" OpName %3 "a_func"
OpName %6 "__modf_result_vec2_f32" OpName %5 "__modf_result_vec2_f32"
OpMemberName %6 0 "fract" OpMemberName %5 0 "fract"
OpMemberName %6 1 "whole" OpMemberName %5 1 "whole"
OpMemberDecorate %6 0 Offset 0 OpMemberDecorate %5 0 Offset 0
OpMemberDecorate %6 1 Offset 8 OpMemberDecorate %5 1 Offset 8
%2 = OpTypeVoid %2 = OpTypeVoid
%1 = OpTypeFunction %2 %1 = OpTypeFunction %2
%8 = OpTypeFloat 32 %7 = OpTypeFloat 32
%7 = OpTypeVector %8 2 %6 = OpTypeVector %7 2
%6 = OpTypeStruct %7 %7 %5 = OpTypeStruct %6 %6
%10 = OpConstant %8 1 %8 = OpConstantNull %6
%11 = OpConstant %8 2 %9 = OpConstant %7 1
%12 = OpConstantComposite %7 %10 %11 %10 = OpConstant %7 2
%11 = OpConstantComposite %6 %9 %10
%12 = OpConstantComposite %5 %8 %11
%3 = OpFunction %2 None %1 %3 = OpFunction %2 None %1
%4 = OpLabel %4 = OpLabel
%5 = OpExtInst %6 %9 ModfStruct %12
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd
)"; )";
@ -1791,7 +1791,7 @@ TEST_F(BuiltinBuilderTest, Const_Call_Modf_f16) {
auto* expr = Call("modf", vec2<f16>(1_h, 2_h)); auto* expr = Call("modf", vec2<f16>(1_h, 2_h));
Func("a_func", utils::Empty, ty.void_(), Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(expr), Decl(Let("l", expr)),
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
@ -1806,27 +1806,27 @@ OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess OpCapability StorageBuffer16BitAccess
OpCapability StorageInputOutput16 OpCapability StorageInputOutput16
%9 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450 OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %3 "a_func" OpEntryPoint Fragment %3 "a_func"
OpExecutionMode %3 OriginUpperLeft OpExecutionMode %3 OriginUpperLeft
OpName %3 "a_func" OpName %3 "a_func"
OpName %6 "__modf_result_vec2_f16" OpName %5 "__modf_result_vec2_f16"
OpMemberName %6 0 "fract" OpMemberName %5 0 "fract"
OpMemberName %6 1 "whole" OpMemberName %5 1 "whole"
OpMemberDecorate %6 0 Offset 0 OpMemberDecorate %5 0 Offset 0
OpMemberDecorate %6 1 Offset 4 OpMemberDecorate %5 1 Offset 4
%2 = OpTypeVoid %2 = OpTypeVoid
%1 = OpTypeFunction %2 %1 = OpTypeFunction %2
%8 = OpTypeFloat 16 %7 = OpTypeFloat 16
%7 = OpTypeVector %8 2 %6 = OpTypeVector %7 2
%6 = OpTypeStruct %7 %7 %5 = OpTypeStruct %6 %6
%10 = OpConstant %8 0x1p+0 %8 = OpConstantNull %6
%11 = OpConstant %8 0x1p+1 %9 = OpConstant %7 0x1p+0
%12 = OpConstantComposite %7 %10 %11 %10 = OpConstant %7 0x1p+1
%11 = OpConstantComposite %6 %9 %10
%12 = OpConstantComposite %5 %8 %11
%3 = OpFunction %2 None %1 %3 = OpFunction %2 None %1
%4 = OpLabel %4 = OpLabel
%5 = OpExtInst %6 %9 ModfStruct %12
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd
)"; )";
@ -1841,7 +1841,7 @@ TEST_F(BuiltinBuilderTest, Runtime_Call_Frexp_f32) {
Func("a_func", utils::Empty, ty.void_(), Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(vec), Decl(vec),
CallStmt(expr), Decl(Let("l", expr)),
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
@ -1897,7 +1897,7 @@ TEST_F(BuiltinBuilderTest, Runtime_Call_Frexp_f16) {
Func("a_func", utils::Empty, ty.void_(), Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(vec), Decl(vec),
CallStmt(expr), Decl(Let("l", expr)),
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
@ -1952,7 +1952,7 @@ OpFunctionEnd
TEST_F(BuiltinBuilderTest, Const_Call_Frexp_f32) { TEST_F(BuiltinBuilderTest, Const_Call_Frexp_f32) {
Func("a_func", utils::Empty, ty.void_(), Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(Call("frexp", vec2<f32>(1_f, 2_f))), Decl(Let("l", Call("frexp", vec2<f32>(1_f, 2_f)))),
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
@ -1963,29 +1963,30 @@ TEST_F(BuiltinBuilderTest, Const_Call_Frexp_f32) {
ASSERT_TRUE(b.Build()) << b.error(); ASSERT_TRUE(b.Build()) << b.error();
auto got = DumpBuilder(b); auto got = DumpBuilder(b);
auto* expect = R"(OpCapability Shader auto* expect = R"(OpCapability Shader
%11 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450 OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %3 "a_func" OpEntryPoint Fragment %3 "a_func"
OpExecutionMode %3 OriginUpperLeft OpExecutionMode %3 OriginUpperLeft
OpName %3 "a_func" OpName %3 "a_func"
OpName %6 "__frexp_result_vec2_f32" OpName %5 "__frexp_result_vec2_f32"
OpMemberName %6 0 "fract" OpMemberName %5 0 "fract"
OpMemberName %6 1 "exp" OpMemberName %5 1 "exp"
OpMemberDecorate %6 0 Offset 0 OpMemberDecorate %5 0 Offset 0
OpMemberDecorate %6 1 Offset 8 OpMemberDecorate %5 1 Offset 8
%2 = OpTypeVoid %2 = OpTypeVoid
%1 = OpTypeFunction %2 %1 = OpTypeFunction %2
%8 = OpTypeFloat 32 %7 = OpTypeFloat 32
%7 = OpTypeVector %8 2 %6 = OpTypeVector %7 2
%10 = OpTypeInt 32 1 %9 = OpTypeInt 32 1
%9 = OpTypeVector %10 2 %8 = OpTypeVector %9 2
%6 = OpTypeStruct %7 %9 %5 = OpTypeStruct %6 %8
%12 = OpConstant %8 1 %10 = OpConstant %7 0.5
%13 = OpConstant %8 2 %11 = OpConstantComposite %6 %10 %10
%14 = OpConstantComposite %7 %12 %13 %12 = OpConstant %9 1
%13 = OpConstant %9 2
%14 = OpConstantComposite %8 %12 %13
%15 = OpConstantComposite %5 %11 %14
%3 = OpFunction %2 None %1 %3 = OpFunction %2 None %1
%4 = OpLabel %4 = OpLabel
%5 = OpExtInst %6 %11 FrexpStruct %14
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd
)"; )";
@ -1999,7 +2000,7 @@ TEST_F(BuiltinBuilderTest, Const_Call_Frexp_f16) {
Func("a_func", utils::Empty, ty.void_(), Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(Call("frexp", vec2<f16>(1_h, 2_h))), Decl(Let("l", Call("frexp", vec2<f16>(1_h, 2_h)))),
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
@ -2014,29 +2015,30 @@ OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess OpCapability StorageBuffer16BitAccess
OpCapability StorageInputOutput16 OpCapability StorageInputOutput16
%11 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450 OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %3 "a_func" OpEntryPoint Fragment %3 "a_func"
OpExecutionMode %3 OriginUpperLeft OpExecutionMode %3 OriginUpperLeft
OpName %3 "a_func" OpName %3 "a_func"
OpName %6 "__frexp_result_vec2_f16" OpName %5 "__frexp_result_vec2_f16"
OpMemberName %6 0 "fract" OpMemberName %5 0 "fract"
OpMemberName %6 1 "exp" OpMemberName %5 1 "exp"
OpMemberDecorate %6 0 Offset 0 OpMemberDecorate %5 0 Offset 0
OpMemberDecorate %6 1 Offset 8 OpMemberDecorate %5 1 Offset 8
%2 = OpTypeVoid %2 = OpTypeVoid
%1 = OpTypeFunction %2 %1 = OpTypeFunction %2
%8 = OpTypeFloat 16 %7 = OpTypeFloat 16
%7 = OpTypeVector %8 2 %6 = OpTypeVector %7 2
%10 = OpTypeInt 32 1 %9 = OpTypeInt 32 1
%9 = OpTypeVector %10 2 %8 = OpTypeVector %9 2
%6 = OpTypeStruct %7 %9 %5 = OpTypeStruct %6 %8
%12 = OpConstant %8 0x1p+0 %10 = OpConstant %7 0x1p-1
%13 = OpConstant %8 0x1p+1 %11 = OpConstantComposite %6 %10 %10
%14 = OpConstantComposite %7 %12 %13 %12 = OpConstant %9 1
%13 = OpConstant %9 2
%14 = OpConstantComposite %8 %12 %13
%15 = OpConstantComposite %5 %11 %14
%3 = OpFunction %2 None %1 %3 = OpFunction %2 None %1
%4 = OpLabel %4 = OpLabel
%5 = OpExtInst %6 %11 FrexpStruct %14
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd
)"; )";
@ -2156,7 +2158,7 @@ TEST_P(BuiltinIntTest, Call_SInt_Scalar) {
auto* expr = Call(param.name, "v"); auto* expr = Call(param.name, "v");
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -2186,7 +2188,7 @@ TEST_P(BuiltinIntTest, Call_SInt_Vector) {
auto* expr = Call(param.name, "v"); auto* expr = Call(param.name, "v");
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -2217,7 +2219,7 @@ TEST_P(BuiltinIntTest, Call_UInt_Scalar) {
auto* expr = Call(param.name, "v"); auto* expr = Call(param.name, "v");
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -2247,7 +2249,7 @@ TEST_P(BuiltinIntTest, Call_UInt_Vector) {
auto* expr = Call(param.name, "v"); auto* expr = Call(param.name, "v");
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -2284,7 +2286,7 @@ TEST_P(Builtin_Builder_SingleParam_Sint_Test, Call_Scalar) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(scalar), Decl(scalar),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -2322,7 +2324,7 @@ TEST_P(Builtin_Builder_SingleParam_Sint_Test, Call_Vector) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(vec), Decl(vec),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -2366,7 +2368,7 @@ TEST_F(Builtin_Builder_Abs_Uint_Test, Call_Scalar) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(scalar), Decl(scalar),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -2399,7 +2401,7 @@ TEST_F(Builtin_Builder_Abs_Uint_Test, Call_Vector) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(scalar), Decl(scalar),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -2436,7 +2438,7 @@ TEST_P(Builtin_Builder_DualParam_SInt_Test, Call_Scalar) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(scalar), Decl(scalar),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -2475,7 +2477,7 @@ TEST_P(Builtin_Builder_DualParam_SInt_Test, Call_Vector) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(vec), Decl(vec),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -2520,7 +2522,7 @@ TEST_P(Builtin_Builder_DualParam_UInt_Test, Call_Scalar) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(scalar), Decl(scalar),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -2559,7 +2561,7 @@ TEST_P(Builtin_Builder_DualParam_UInt_Test, Call_Vector) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(vec), Decl(vec),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -2604,7 +2606,7 @@ TEST_P(Builtin_Builder_ThreeParam_Sint_Test, Call_Scalar) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(scalar), Decl(scalar),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -2644,7 +2646,7 @@ TEST_P(Builtin_Builder_ThreeParam_Sint_Test, Call_Vector) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(vec), Decl(vec),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -2690,7 +2692,7 @@ TEST_P(Builtin_Builder_ThreeParam_Uint_Test, Call_Scalar) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(scalar), Decl(scalar),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -2730,7 +2732,7 @@ TEST_P(Builtin_Builder_ThreeParam_Uint_Test, Call_Vector) {
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Decl(vec), Decl(vec),
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -3118,7 +3120,7 @@ TEST_F(BuiltinBuilderTest, Call_Determinant_f32) {
auto* expr = Call("determinant", "var"); auto* expr = Call("determinant", "var");
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -3155,7 +3157,7 @@ TEST_F(BuiltinBuilderTest, Call_Determinant_f16) {
auto* expr = Call("determinant", "var"); auto* expr = Call("determinant", "var");
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -3190,7 +3192,7 @@ TEST_F(BuiltinBuilderTest, Call_Transpose_f32) {
auto* expr = Call("transpose", "var"); auto* expr = Call("transpose", "var");
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -3228,7 +3230,7 @@ TEST_F(BuiltinBuilderTest, Call_Transpose_f16) {
auto* expr = Call("transpose", "var"); auto* expr = Call("transpose", "var");
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -3269,7 +3271,7 @@ TEST_F(BuiltinBuilderTest, Call_Dot_F32) {
auto* expr = Call("dot", "v", "v"); auto* expr = Call("dot", "v", "v");
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -3300,7 +3302,7 @@ TEST_F(BuiltinBuilderTest, Call_Dot_F16) {
auto* expr = Call("dot", "v", "v"); auto* expr = Call("dot", "v", "v");
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -3329,7 +3331,7 @@ TEST_F(BuiltinBuilderTest, Call_Dot_U32) {
auto* expr = Call("dot", "v", "v"); auto* expr = Call("dot", "v", "v");
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -3368,7 +3370,7 @@ TEST_F(BuiltinBuilderTest, Call_Dot_I32) {
auto* expr = Call("dot", "v", "v"); auto* expr = Call("dot", "v", "v");
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
Assign(Phony(), expr), Decl(Let("l", expr)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -3414,7 +3416,7 @@ TEST_P(BuiltinDeriveTest, Call_Derivative_Scalar) {
auto* expr = Call(param.name, "v"); auto* expr = Call(param.name, "v");
auto* func = Func("func", utils::Empty, ty.void_(), auto* func = Func("func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(expr), Decl(Let("l", expr)),
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
@ -3447,7 +3449,7 @@ TEST_P(BuiltinDeriveTest, Call_Derivative_Vector) {
auto* expr = Call(param.name, "v"); auto* expr = Call(param.name, "v");
auto* func = Func("func", utils::Empty, ty.void_(), auto* func = Func("func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(expr), Decl(Let("l", expr)),
}, },
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
@ -3942,11 +3944,11 @@ TEST_P(Builtin_Builder_DataPacking_Test, Binary) {
auto param = GetParam(); auto param = GetParam();
bool pack4 = param.name == "pack4x8snorm" || param.name == "pack4x8unorm"; bool pack4 = param.name == "pack4x8snorm" || param.name == "pack4x8unorm";
auto* call = pack4 ? Call(param.name, vec4<f32>(1_f, 1_f, 1_f, 1_f)) auto* call = pack4 ? Call(param.name, vec4<f32>("one")) : Call(param.name, vec2<f32>("one"));
: Call(param.name, vec2<f32>(1_f, 1_f));
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(call), Decl(Let("one", Expr(1_f))),
Decl(Let("l", call)),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -3955,18 +3957,18 @@ TEST_P(Builtin_Builder_DataPacking_Test, Binary) {
if (pack4) { if (pack4) {
auto got = DumpBuilder(b); auto got = DumpBuilder(b);
auto expect = R"(%7 = OpExtInstImport "GLSL.std.450" auto expect = R"(%9 = OpExtInstImport "GLSL.std.450"
OpName %3 "a_func" OpName %3 "a_func"
%2 = OpTypeVoid %2 = OpTypeVoid
%1 = OpTypeFunction %2 %1 = OpTypeFunction %2
%6 = OpTypeInt 32 0 %5 = OpTypeFloat 32
%9 = OpTypeFloat 32 %6 = OpConstant %5 1
%8 = OpTypeVector %9 4 %8 = OpTypeInt 32 0
%10 = OpConstant %9 1 %10 = OpTypeVector %5 4
%11 = OpConstantComposite %8 %10 %10 %10 %10
%3 = OpFunction %2 None %1 %3 = OpFunction %2 None %1
%4 = OpLabel %4 = OpLabel
%5 = OpExtInst %6 %7 )" + %11 = OpCompositeConstruct %10 %6 %6 %6 %6
%7 = OpExtInst %8 %9 )" +
param.op + param.op +
R"( %11 R"( %11
OpReturn OpReturn
@ -3975,18 +3977,18 @@ OpFunctionEnd
EXPECT_EQ(got, expect); EXPECT_EQ(got, expect);
} else { } else {
auto got = DumpBuilder(b); auto got = DumpBuilder(b);
auto expect = R"(%7 = OpExtInstImport "GLSL.std.450" auto expect = R"(%9 = OpExtInstImport "GLSL.std.450"
OpName %3 "a_func" OpName %3 "a_func"
%2 = OpTypeVoid %2 = OpTypeVoid
%1 = OpTypeFunction %2 %1 = OpTypeFunction %2
%6 = OpTypeInt 32 0 %5 = OpTypeFloat 32
%9 = OpTypeFloat 32 %6 = OpConstant %5 1
%8 = OpTypeVector %9 2 %8 = OpTypeInt 32 0
%10 = OpConstant %9 1 %10 = OpTypeVector %5 2
%11 = OpConstantComposite %8 %10 %10
%3 = OpFunction %2 None %1 %3 = OpFunction %2 None %1
%4 = OpLabel %4 = OpLabel
%5 = OpExtInst %6 %7 )" + %11 = OpCompositeConstruct %10 %6 %6
%7 = OpExtInst %8 %9 )" +
param.op + param.op +
R"( %11 R"( %11
OpReturn OpReturn
@ -4016,7 +4018,8 @@ TEST_P(Builtin_Builder_DataUnpacking_Test, Binary) {
bool pack4 = param.name == "unpack4x8snorm" || param.name == "unpack4x8unorm"; bool pack4 = param.name == "unpack4x8snorm" || param.name == "unpack4x8unorm";
auto* func = Func("a_func", utils::Empty, ty.void_(), auto* func = Func("a_func", utils::Empty, ty.void_(),
utils::Vector{ utils::Vector{
CallStmt(Call(param.name, 1_u)), Decl(Let("one", Expr(1_u))),
Decl(Let("l", Call(param.name, "one"))),
}); });
spirv::Builder& b = Build(); spirv::Builder& b = Build();
@ -4025,38 +4028,38 @@ TEST_P(Builtin_Builder_DataUnpacking_Test, Binary) {
if (pack4) { if (pack4) {
auto got = DumpBuilder(b); auto got = DumpBuilder(b);
auto expect = R"(%8 = OpExtInstImport "GLSL.std.450" auto expect = R"(%10 = OpExtInstImport "GLSL.std.450"
OpName %3 "a_func" OpName %3 "a_func"
%2 = OpTypeVoid %2 = OpTypeVoid
%1 = OpTypeFunction %2 %1 = OpTypeFunction %2
%7 = OpTypeFloat 32 %5 = OpTypeInt 32 0
%6 = OpTypeVector %7 4 %6 = OpConstant %5 1
%9 = OpTypeInt 32 0 %9 = OpTypeFloat 32
%10 = OpConstant %9 1 %8 = OpTypeVector %9 4
%3 = OpFunction %2 None %1 %3 = OpFunction %2 None %1
%4 = OpLabel %4 = OpLabel
%5 = OpExtInst %6 %8 )" + %7 = OpExtInst %8 %10 )" +
param.op + param.op +
R"( %10 R"( %6
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd
)"; )";
EXPECT_EQ(got, expect); EXPECT_EQ(got, expect);
} else { } else {
auto got = DumpBuilder(b); auto got = DumpBuilder(b);
auto expect = R"(%8 = OpExtInstImport "GLSL.std.450" auto expect = R"(%10 = OpExtInstImport "GLSL.std.450"
OpName %3 "a_func" OpName %3 "a_func"
%2 = OpTypeVoid %2 = OpTypeVoid
%1 = OpTypeFunction %2 %1 = OpTypeFunction %2
%7 = OpTypeFloat 32 %5 = OpTypeInt 32 0
%6 = OpTypeVector %7 2 %6 = OpConstant %5 1
%9 = OpTypeInt 32 0 %9 = OpTypeFloat 32
%10 = OpConstant %9 1 %8 = OpTypeVector %9 2
%3 = OpFunction %2 None %1 %3 = OpFunction %2 None %1
%4 = OpLabel %4 = OpLabel
%5 = OpExtInst %6 %8 )" + %7 = OpExtInst %8 %10 )" +
param.op + param.op +
R"( %10 R"( %6
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd
)"; )";

10
src/tint/writer/spirv/builder_builtin_texture_test.cc
View File

@ -3717,7 +3717,8 @@ TEST_P(BuiltinTextureTest, Call) {
auto* sampler = param.BuildSamplerVariable(this); auto* sampler = param.BuildSamplerVariable(this);
auto* call = Call(param.function, param.args(this)); auto* call = Call(param.function, param.args(this));
auto* stmt = CallStmt(call); auto* stmt = param.returns_value ? static_cast<const ast::Statement*>(Assign(Phony(), call))
: static_cast<const ast::Statement*>(CallStmt(call));
Func("func", utils::Empty, ty.void_(), utils::Vector{stmt}, Func("func", utils::Empty, ty.void_(), utils::Vector{stmt},
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
@ -3745,8 +3746,9 @@ TEST_P(BuiltinTextureTest, ValidateSPIRV) {
param.BuildSamplerVariable(this); param.BuildSamplerVariable(this);
auto* call = Call(param.function, param.args(this)); auto* call = Call(param.function, param.args(this));
auto* stmt = param.returns_value ? static_cast<const ast::Statement*>(Assign(Phony(), call))
: static_cast<const ast::Statement*>(CallStmt(call));
auto* stmt = CallStmt(call);
Func("main", utils::Empty, ty.void_(), utils::Vector{stmt}, Func("main", utils::Empty, ty.void_(), utils::Vector{stmt},
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
@ -3769,7 +3771,9 @@ TEST_P(BuiltinTextureTest, OutsideFunction_IsError) {
auto* sampler = param.BuildSamplerVariable(this); auto* sampler = param.BuildSamplerVariable(this);
auto* call = Call(param.function, param.args(this)); auto* call = Call(param.function, param.args(this));
auto* stmt = CallStmt(call); auto* stmt = param.returns_value ? static_cast<const ast::Statement*>(Assign(Phony(), call))
: static_cast<const ast::Statement*>(CallStmt(call));
Func("func", utils::Empty, ty.void_(), utils::Vector{stmt}, Func("func", utils::Empty, ty.void_(), utils::Vector{stmt},
utils::Vector{ utils::Vector{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),

8
src/tint/writer/wgsl/generator_impl_function_test.cc
View File

@ -87,9 +87,9 @@ TEST_F(WgslGeneratorImplTest, Emit_Function_WithAttribute_WorkgroupSize) {
} }
TEST_F(WgslGeneratorImplTest, Emit_Function_WithAttribute_MustUse) { TEST_F(WgslGeneratorImplTest, Emit_Function_WithAttribute_MustUse) {
auto* func = Func("my_func", utils::Empty, ty.void_(), auto* func = Func("my_func", utils::Empty, ty.i32(),
utils::Vector{ utils::Vector{
Return(), Return(1_i),
}, },
utils::Vector{ utils::Vector{
MustUse(), MustUse(),
@ -101,8 +101,8 @@ TEST_F(WgslGeneratorImplTest, Emit_Function_WithAttribute_MustUse) {
ASSERT_TRUE(gen.EmitFunction(func)); ASSERT_TRUE(gen.EmitFunction(func));
EXPECT_EQ(gen.result(), R"( @must_use EXPECT_EQ(gen.result(), R"( @must_use
fn my_func() { fn my_func() -> i32 {
return; return 1i;
} }
)"); )");
} }

6
test/tint/bug/chromium/1350147.wgsl
View File

@ -1,5 +1,5 @@
fn original_clusterfuzz_code() { fn original_clusterfuzz_code() {
atan2(1,.1); _ = atan2(1,.1);
} }
fn more_tests_that_would_fail() { fn more_tests_that_would_fail() {
@ -19,7 +19,7 @@ fn more_tests_that_would_fail() {
// abstract constant value is not handled by backends. These should be removed by RemovePhonies // abstract constant value is not handled by backends. These should be removed by RemovePhonies
// transform. // transform.
{ {
atan2(1, 0.1); _ = atan2(1, 0.1);
atan2(0.1, 1); _ = atan2(0.1, 1);
} }
} }

6
test/tint/bug/chromium/1350147.wgsl.expected.wgsl
View File

@ -1,5 +1,5 @@
fn original_clusterfuzz_code() { fn original_clusterfuzz_code() {
atan2(1, 0.1); _ = atan2(1, 0.1);
} }
fn more_tests_that_would_fail() { fn more_tests_that_would_fail() {
@ -12,7 +12,7 @@ fn more_tests_that_would_fail() {
let b = (1.5 + 1); let b = (1.5 + 1);
} }
{ {
atan2(1, 0.1); _ = atan2(1, 0.1);
atan2(0.1, 1); _ = atan2(0.1, 1);
} }
} }

2
test/tint/bug/chromium/1360925.wgsl
View File

@ -3,5 +3,5 @@ var<storage> G : array<i32>;
fn n() { fn n() {
let p = &G; let p = &G;
arrayLength(p); _ = arrayLength(p);
} }

1
test/tint/bug/chromium/1360925.wgsl.expected.glsl
View File

@ -9,6 +9,5 @@ layout(binding = 0, std430) buffer G_block_ssbo {
} G; } G;
void n() { void n() {
uint(G.inner.length());
} }

4
test/tint/bug/chromium/1360925.wgsl.expected.spvasm
View File

@ -1,7 +1,7 @@
; SPIR-V ; SPIR-V
; Version: 1.3 ; Version: 1.3
; Generator: Google Tint Compiler; 0 ; Generator: Google Tint Compiler; 0
; Bound: 14 ; Bound: 12
; Schema: 0 ; Schema: 0
OpCapability Shader OpCapability Shader
OpMemoryModel Logical GLSL450 OpMemoryModel Logical GLSL450
@ -25,13 +25,11 @@
%G = OpVariable %_ptr_StorageBuffer_G_block StorageBuffer %G = OpVariable %_ptr_StorageBuffer_G_block StorageBuffer
%void = OpTypeVoid %void = OpTypeVoid
%6 = OpTypeFunction %void %6 = OpTypeFunction %void
%uint = OpTypeInt 32 0
%unused_entry_point = OpFunction %void None %6 %unused_entry_point = OpFunction %void None %6
%9 = OpLabel %9 = OpLabel
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd
%n = OpFunction %void None %6 %n = OpFunction %void None %6
%11 = OpLabel %11 = OpLabel
%12 = OpArrayLength %uint %G 0
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

2
test/tint/bug/chromium/1360925.wgsl.expected.wgsl
View File

@ -2,5 +2,5 @@
fn n() { fn n() {
let p = &(G); let p = &(G);
arrayLength(p); _ = arrayLength(p);
} }

2
test/tint/bug/chromium/1405676.wgsl
View File

@ -1,6 +1,6 @@
@group(0) @binding(0) var arg_0 : texture_1d<i32>; @group(0) @binding(0) var arg_0 : texture_1d<i32>;
fn d() { fn d() {
textureLoad(arg_0, 1, 0); _ = textureLoad(arg_0, 1, 0);
let l = sin(3.0); let l = sin(3.0);
} }

1
test/tint/bug/chromium/1405676.wgsl.expected.dxc.hlsl
View File

@ -6,6 +6,5 @@ void unused_entry_point() {
Texture1D<int4> arg_0 : register(t0, space0); Texture1D<int4> arg_0 : register(t0, space0);
void d() { void d() {
arg_0.Load(int2(1, 0));
const float l = 0.141120002f; const float l = 0.141120002f;
} }

1
test/tint/bug/chromium/1405676.wgsl.expected.fxc.hlsl
View File

@ -6,6 +6,5 @@ void unused_entry_point() {
Texture1D<int4> arg_0 : register(t0, space0); Texture1D<int4> arg_0 : register(t0, space0);
void d() { void d() {
arg_0.Load(int2(1, 0));
const float l = 0.141120002f; const float l = 0.141120002f;
} }

2
test/tint/bug/chromium/1405676.wgsl.expected.glsl
View File

@ -4,9 +4,7 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void unused_entry_point() { void unused_entry_point() {
return; return;
} }
uniform highp isampler2D arg_0_1;
void d() { void d() {
texelFetch(arg_0_1, ivec2(1, 0), 0);
float l = 0.141120002f; float l = 0.141120002f;
} }

3
test/tint/bug/chromium/1405676.wgsl.expected.msl
View File

@ -1,8 +1,7 @@
#include <metal_stdlib> #include <metal_stdlib>
using namespace metal; using namespace metal;
void d(texture1d<int, access::sample> tint_symbol) { void d() {
tint_symbol.read(uint(1), 0);
float const l = 0.141120002f; float const l = 0.141120002f;
} }

7
test/tint/bug/chromium/1405676.wgsl.expected.spvasm
View File

@ -1,7 +1,7 @@
; SPIR-V ; SPIR-V
; Version: 1.3 ; Version: 1.3
; Generator: Google Tint Compiler; 0 ; Generator: Google Tint Compiler; 0
; Bound: 18 ; Bound: 13
; Schema: 0 ; Schema: 0
OpCapability Shader OpCapability Shader
OpCapability Sampled1D OpCapability Sampled1D
@ -19,9 +19,6 @@
%arg_0 = OpVariable %_ptr_UniformConstant_3 UniformConstant %arg_0 = OpVariable %_ptr_UniformConstant_3 UniformConstant
%void = OpTypeVoid %void = OpTypeVoid
%5 = OpTypeFunction %void %5 = OpTypeFunction %void
%v4int = OpTypeVector %int 4
%int_1 = OpConstant %int 1
%15 = OpConstantNull %int
%float = OpTypeFloat 32 %float = OpTypeFloat 32
%float_0_141120002 = OpConstant %float 0.141120002 %float_0_141120002 = OpConstant %float 0.141120002
%unused_entry_point = OpFunction %void None %5 %unused_entry_point = OpFunction %void None %5
@ -30,7 +27,5 @@
OpFunctionEnd OpFunctionEnd
%d = OpFunction %void None %5 %d = OpFunction %void None %5
%10 = OpLabel %10 = OpLabel
%13 = OpLoad %3 %arg_0
%11 = OpImageFetch %v4int %13 %int_1 Lod %15
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

2
test/tint/bug/chromium/1405676.wgsl.expected.wgsl
View File

@ -1,6 +1,6 @@
@group(0) @binding(0) var arg_0 : texture_1d<i32>; @group(0) @binding(0) var arg_0 : texture_1d<i32>;
fn d() { fn d() {
textureLoad(arg_0, 1, 0); _ = textureLoad(arg_0, 1, 0);
let l = sin(3.0); let l = sin(3.0);
} }