tint/intrinsics: Texture queries now return unsigned integer / vectors

To match the spec.

Also add a bunch of missing texture test cases to
src/tint/ast/builtin_texture_helper_test.cc. Fix all the tests that were
broken because these were not being exercised.

Fixed: tint:1526
Change-Id: I207b51d307bbdc054b595e0e0e0fd3330607e171
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/106681
Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
Auto-Submit: Ben Clayton <bclayton@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
This commit is contained in:
Ben Clayton 2022-10-26 18:36:44 +00:00 committed by Dawn LUCI CQ
parent 980145bc16
commit 13f089095f
5443 changed files with 116933 additions and 115307 deletions

View File

@ -2,6 +2,10 @@
## Changes for M109
### Breaking changes
* `textureDimensions()`, `textureNumLayers()` and `textureNumLevels()` now return unsigned integers / vectors. [tint:1526](crbug.com/tint/1526)
### New features
* Uniformity analysis failures are warnings again [tint:1728](crbug.com/tint/1728)

View File

@ -310,23 +310,23 @@ fn IsEqualTo(pixel : vec4<f32>, expected : vec4<f32>) -> bool {
auto texelType = "vec4<" + componentFmt + ">";
std::string sliceCount;
std::string textureStore;
std::string textureSize = "textureDimensions(storageImage0).xy";
std::string textureSize = "vec2<i32>(textureDimensions(storageImage0).xy)";
switch (dimension) {
case wgpu::TextureViewDimension::e1D:
sliceCount = "1";
textureStore = "textureStore(storageImage0, x, expected)";
textureSize = "vec2<i32>(textureDimensions(storageImage0), 1)";
textureSize = "vec2<i32>(i32(textureDimensions(storageImage0)), 1)";
break;
case wgpu::TextureViewDimension::e2D:
sliceCount = "1";
textureStore = "textureStore(storageImage0, vec2<i32>(x, y), expected)";
break;
case wgpu::TextureViewDimension::e2DArray:
sliceCount = "textureNumLayers(storageImage0)";
sliceCount = "i32(textureNumLayers(storageImage0))";
textureStore = "textureStore(storageImage0, vec2<i32>(x, y), slice, expected)";
break;
case wgpu::TextureViewDimension::e3D:
sliceCount = "textureDimensions(storageImage0).z";
sliceCount = "i32(textureDimensions(storageImage0).z)";
textureStore = "textureStore(storageImage0, vec3<i32>(x, y, slice), expected)";
break;
default:

View File

@ -75,7 +75,7 @@ class SubresourceTrackingPerf : public DawnPerfTestWithParams<SubresourceTrackin
pipelineDesc.cFragment.module = utils::CreateShaderModule(device, R"(
@group(0) @binding(0) var materials : texture_2d<f32>;
@fragment fn main() -> @location(0) vec4<f32> {
let foo : vec2<i32> = textureDimensions(materials);
_ = materials;
return vec4<f32>(1.0, 0.0, 0.0, 1.0);
}
)");

View File

@ -183,7 +183,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
return {
{
ValidTextureOverload::kDimensions1d,
"textureDimensions(t : texture_1d<f32>) -> i32",
"textureDimensions(t : texture_1d<f32>) -> u32",
TextureKind::kRegular,
ast::SamplerKind::kSampler,
ast::TextureDimension::k1d,
@ -193,7 +193,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kDimensions2d,
"textureDimensions(t : texture_2d<f32>) -> vec2<i32>",
"textureDimensions(t : texture_2d<f32>) -> vec2<u32>",
TextureKind::kRegular,
ast::SamplerKind::kSampler,
ast::TextureDimension::k2d,
@ -204,7 +204,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
{
ValidTextureOverload::kDimensions2dLevel,
"textureDimensions(t : texture_2d<f32>,\n"
" level : i32) -> vec2<i32>",
" level : i32) -> vec2<u32>",
TextureKind::kRegular,
ast::SamplerKind::kSampler,
ast::TextureDimension::k2d,
@ -214,7 +214,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kDimensions2dArray,
"textureDimensions(t : texture_2d_array<f32>) -> vec2<i32>",
"textureDimensions(t : texture_2d_array<f32>) -> vec2<u32>",
TextureKind::kRegular,
ast::SamplerKind::kSampler,
ast::TextureDimension::k2dArray,
@ -225,7 +225,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
{
ValidTextureOverload::kDimensions2dArrayLevel,
"textureDimensions(t : texture_2d_array<f32>,\n"
" level : i32) -> vec2<i32>",
" level : i32) -> vec2<u32>",
TextureKind::kRegular,
ast::SamplerKind::kSampler,
ast::TextureDimension::k2dArray,
@ -235,7 +235,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kDimensions3d,
"textureDimensions(t : texture_3d<f32>) -> vec3<i32>",
"textureDimensions(t : texture_3d<f32>) -> vec3<u32>",
TextureKind::kRegular,
ast::SamplerKind::kSampler,
ast::TextureDimension::k3d,
@ -246,7 +246,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
{
ValidTextureOverload::kDimensions3dLevel,
"textureDimensions(t : texture_3d<f32>,\n"
" level : i32) -> vec3<i32>",
" level : i32) -> vec3<u32>",
TextureKind::kRegular,
ast::SamplerKind::kSampler,
ast::TextureDimension::k3d,
@ -256,7 +256,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kDimensionsCube,
"textureDimensions(t : texture_cube<f32>) -> vec2<i32>",
"textureDimensions(t : texture_cube<f32>) -> vec2<u32>",
TextureKind::kRegular,
ast::SamplerKind::kSampler,
ast::TextureDimension::kCube,
@ -267,7 +267,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
{
ValidTextureOverload::kDimensionsCubeLevel,
"textureDimensions(t : texture_cube<f32>,\n"
" level : i32) -> vec2<i32>",
" level : i32) -> vec2<u32>",
TextureKind::kRegular,
ast::SamplerKind::kSampler,
ast::TextureDimension::kCube,
@ -277,7 +277,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kDimensionsCubeArray,
"textureDimensions(t : texture_cube_array<f32>) -> vec2<i32>",
"textureDimensions(t : texture_cube_array<f32>) -> vec2<u32>",
TextureKind::kRegular,
ast::SamplerKind::kSampler,
ast::TextureDimension::kCubeArray,
@ -288,7 +288,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
{
ValidTextureOverload::kDimensionsCubeArrayLevel,
"textureDimensions(t : texture_cube_array<f32>,\n"
" level : i32) -> vec2<i32>",
" level : i32) -> vec2<u32>",
TextureKind::kRegular,
ast::SamplerKind::kSampler,
ast::TextureDimension::kCubeArray,
@ -298,7 +298,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kDimensionsMultisampled2d,
"textureDimensions(t : texture_multisampled_2d<f32>)-> vec2<i32>",
"textureDimensions(t : texture_multisampled_2d<f32>)-> vec2<u32>",
TextureKind::kMultisampled,
ast::SamplerKind::kSampler,
ast::TextureDimension::k2d,
@ -308,7 +308,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kDimensionsDepth2d,
"textureDimensions(t : texture_depth_2d) -> vec2<i32>",
"textureDimensions(t : texture_depth_2d) -> vec2<u32>",
TextureKind::kDepth,
ast::SamplerKind::kSampler,
ast::TextureDimension::k2d,
@ -319,7 +319,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
{
ValidTextureOverload::kDimensionsDepth2dLevel,
"textureDimensions(t : texture_depth_2d,\n"
" level : i32) -> vec2<i32>",
" level : i32) -> vec2<u32>",
TextureKind::kDepth,
ast::SamplerKind::kSampler,
ast::TextureDimension::k2d,
@ -329,7 +329,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kDimensionsDepth2dArray,
"textureDimensions(t : texture_depth_2d_array) -> vec2<i32>",
"textureDimensions(t : texture_depth_2d_array) -> vec2<u32>",
TextureKind::kDepth,
ast::SamplerKind::kSampler,
ast::TextureDimension::k2dArray,
@ -340,7 +340,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
{
ValidTextureOverload::kDimensionsDepth2dArrayLevel,
"textureDimensions(t : texture_depth_2d_array,\n"
" level : i32) -> vec2<i32>",
" level : i32) -> vec2<u32>",
TextureKind::kDepth,
ast::SamplerKind::kSampler,
ast::TextureDimension::k2dArray,
@ -350,7 +350,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kDimensionsDepthCube,
"textureDimensions(t : texture_depth_cube) -> vec2<i32>",
"textureDimensions(t : texture_depth_cube) -> vec2<u32>",
TextureKind::kDepth,
ast::SamplerKind::kSampler,
ast::TextureDimension::kCube,
@ -361,7 +361,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
{
ValidTextureOverload::kDimensionsDepthCubeLevel,
"textureDimensions(t : texture_depth_cube,\n"
" level : i32) -> vec2<i32>",
" level : i32) -> vec2<u32>",
TextureKind::kDepth,
ast::SamplerKind::kSampler,
ast::TextureDimension::kCube,
@ -371,7 +371,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kDimensionsDepthCubeArray,
"textureDimensions(t : texture_depth_cube_array) -> vec2<i32>",
"textureDimensions(t : texture_depth_cube_array) -> vec2<u32>",
TextureKind::kDepth,
ast::SamplerKind::kSampler,
ast::TextureDimension::kCubeArray,
@ -382,7 +382,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
{
ValidTextureOverload::kDimensionsDepthCubeArrayLevel,
"textureDimensions(t : texture_depth_cube_array,\n"
" level : i32) -> vec2<i32>",
" level : i32) -> vec2<u32>",
TextureKind::kDepth,
ast::SamplerKind::kSampler,
ast::TextureDimension::kCubeArray,
@ -392,7 +392,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kDimensionsDepthMultisampled2d,
"textureDimensions(t : texture_depth_multisampled_2d) -> vec2<i32>",
"textureDimensions(t : texture_depth_multisampled_2d) -> vec2<u32>",
TextureKind::kDepthMultisampled,
ast::SamplerKind::kSampler,
ast::TextureDimension::k2d,
@ -402,7 +402,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kDimensionsStorageWO1d,
"textureDimensions(t : texture_storage_1d<rgba32float>) -> i32",
"textureDimensions(t : texture_storage_1d<rgba32float>) -> u32",
ast::Access::kWrite,
ast::TexelFormat::kRgba32Float,
ast::TextureDimension::k1d,
@ -412,8 +412,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kDimensionsStorageWO2d,
"textureDimensions(t : texture_storage_2d<rgba32float>) -> "
"vec2<i32>",
"textureDimensions(t : texture_storage_2d<rgba32float>) -> vec2<u32>",
ast::Access::kWrite,
ast::TexelFormat::kRgba32Float,
ast::TextureDimension::k2d,
@ -423,8 +422,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kDimensionsStorageWO2dArray,
"textureDimensions(t : texture_storage_2d_array<rgba32float>) -> "
"vec2<i32>",
"textureDimensions(t : texture_storage_2d_array<rgba32float>) -> vec2<u32>",
ast::Access::kWrite,
ast::TexelFormat::kRgba32Float,
ast::TextureDimension::k2dArray,
@ -434,8 +432,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kDimensionsStorageWO3d,
"textureDimensions(t : texture_storage_3d<rgba32float>) -> "
"vec3<i32>",
"textureDimensions(t : texture_storage_3d<rgba32float>) -> vec3<u32>",
ast::Access::kWrite,
ast::TexelFormat::kRgba32Float,
ast::TextureDimension::k3d,
@ -788,7 +785,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kNumLayers2dArray,
"textureNumLayers(t : texture_2d_array<f32>) -> i32",
"textureNumLayers(t : texture_2d_array<f32>) -> u32",
TextureKind::kRegular,
ast::SamplerKind::kSampler,
ast::TextureDimension::k2dArray,
@ -798,7 +795,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kNumLayersCubeArray,
"textureNumLayers(t : texture_cube_array<f32>) -> i32",
"textureNumLayers(t : texture_cube_array<f32>) -> u32",
TextureKind::kRegular,
ast::SamplerKind::kSampler,
ast::TextureDimension::kCubeArray,
@ -808,7 +805,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kNumLayersDepth2dArray,
"textureNumLayers(t : texture_depth_2d_array) -> i32",
"textureNumLayers(t : texture_depth_2d_array) -> u32",
TextureKind::kDepth,
ast::SamplerKind::kSampler,
ast::TextureDimension::k2dArray,
@ -818,7 +815,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kNumLayersDepthCubeArray,
"textureNumLayers(t : texture_depth_cube_array) -> i32",
"textureNumLayers(t : texture_depth_cube_array) -> u32",
TextureKind::kDepth,
ast::SamplerKind::kSampler,
ast::TextureDimension::kCubeArray,
@ -828,7 +825,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kNumLayersStorageWO2dArray,
"textureNumLayers(t : texture_storage_2d_array<rgba32float>) -> i32",
"textureNumLayers(t : texture_storage_2d_array<rgba32float>) -> u32",
ast::Access::kWrite,
ast::TexelFormat::kRgba32Float,
ast::TextureDimension::k2dArray,
@ -838,7 +835,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kNumLevels2d,
"textureNumLevels(t : texture_2d<f32>) -> i32",
"textureNumLevels(t : texture_2d<f32>) -> u32",
TextureKind::kRegular,
ast::SamplerKind::kSampler,
ast::TextureDimension::k2d,
@ -848,7 +845,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kNumLevels2dArray,
"textureNumLevels(t : texture_2d_array<f32>) -> i32",
"textureNumLevels(t : texture_2d_array<f32>) -> u32",
TextureKind::kRegular,
ast::SamplerKind::kSampler,
ast::TextureDimension::k2dArray,
@ -858,7 +855,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kNumLevels3d,
"textureNumLevels(t : texture_3d<f32>) -> i32",
"textureNumLevels(t : texture_3d<f32>) -> u32",
TextureKind::kRegular,
ast::SamplerKind::kSampler,
ast::TextureDimension::k3d,
@ -868,7 +865,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kNumLevelsCube,
"textureNumLevels(t : texture_cube<f32>) -> i32",
"textureNumLevels(t : texture_cube<f32>) -> u32",
TextureKind::kRegular,
ast::SamplerKind::kSampler,
ast::TextureDimension::kCube,
@ -878,7 +875,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kNumLevelsCubeArray,
"textureNumLevels(t : texture_cube_array<f32>) -> i32",
"textureNumLevels(t : texture_cube_array<f32>) -> u32",
TextureKind::kRegular,
ast::SamplerKind::kSampler,
ast::TextureDimension::kCubeArray,
@ -888,7 +885,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kNumLevelsDepth2d,
"textureNumLevels(t : texture_depth_2d) -> i32",
"textureNumLevels(t : texture_depth_2d) -> u32",
TextureKind::kDepth,
ast::SamplerKind::kSampler,
ast::TextureDimension::k2d,
@ -898,7 +895,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kNumLevelsDepth2dArray,
"textureNumLevels(t : texture_depth_2d_array) -> i32",
"textureNumLevels(t : texture_depth_2d_array) -> u32",
TextureKind::kDepth,
ast::SamplerKind::kSampler,
ast::TextureDimension::k2dArray,
@ -908,7 +905,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kNumLevelsDepthCube,
"textureNumLevels(t : texture_depth_cube) -> i32",
"textureNumLevels(t : texture_depth_cube) -> u32",
TextureKind::kDepth,
ast::SamplerKind::kSampler,
ast::TextureDimension::kCube,
@ -918,7 +915,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kNumLevelsDepthCubeArray,
"textureNumLevels(t : texture_depth_cube_array) -> i32",
"textureNumLevels(t : texture_depth_cube_array) -> u32",
TextureKind::kDepth,
ast::SamplerKind::kSampler,
ast::TextureDimension::kCubeArray,
@ -928,7 +925,7 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
},
{
ValidTextureOverload::kNumSamplesMultisampled2d,
"textureNumSamples(t : texture_multisampled_2d<f32>) -> i32",
"textureNumSamples(t : texture_multisampled_2d<f32>) -> u32",
TextureKind::kMultisampled,
ast::SamplerKind::kSampler,
ast::TextureDimension::k2d,
@ -936,6 +933,16 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
"textureNumSamples",
[](ProgramBuilder* b) { return b->ExprList("texture"); },
},
{
ValidTextureOverload::kNumSamplesDepthMultisampled2d,
"textureNumSamples(t : texture_depth_multisampled_2d<f32>) -> u32",
TextureKind::kMultisampled,
ast::SamplerKind::kComparisonSampler,
ast::TextureDimension::k2d,
TextureDataType::kF32,
"textureNumSamples",
[](ProgramBuilder* b) { return b->ExprList("texture"); },
},
{
ValidTextureOverload::kSample1dF32,
"textureSample(t : texture_1d<f32>,\n"
@ -1918,6 +1925,124 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
5_f); // depth_ref
},
},
{
ValidTextureOverload::kSampleCompareLevelDepth2dF32,
"textureSampleCompareLevel(t : texture_depth_2d,\n"
" s : sampler_comparison,\n"
" coords : vec2<f32>,\n"
" depth_ref : f32) -> f32",
TextureKind::kDepth,
ast::SamplerKind::kComparisonSampler,
ast::TextureDimension::k2d,
TextureDataType::kF32,
"textureSampleCompareLevel",
[](ProgramBuilder* b) {
return b->ExprList("texture", // t
"sampler", // s
b->vec2<f32>(1_f, 2_f), // coords
3_f); // depth_ref
},
},
{
ValidTextureOverload::kSampleCompareLevelDepth2dOffsetF32,
"textureSampleCompareLevel(t : texture_depth_2d,\n"
" s : sampler_comparison,\n"
" coords : vec2<f32>,\n"
" depth_ref : f32,\n"
" offset : vec2<i32>) -> f32",
TextureKind::kDepth,
ast::SamplerKind::kComparisonSampler,
ast::TextureDimension::k2d,
TextureDataType::kF32,
"textureSampleCompareLevel",
[](ProgramBuilder* b) {
return b->ExprList("texture", // t
"sampler", // s
b->vec2<f32>(1_f, 2_f), // coords
3_f, // depth_ref
b->vec2<i32>(4_i, 5_i)); // offset
},
},
{
ValidTextureOverload::kSampleCompareLevelDepth2dArrayF32,
"textureSampleCompareLevel(t : texture_depth_2d_array,\n"
" s : sampler_comparison,\n"
" coords : vec2<f32>,\n"
" array_index : i32,\n"
" depth_ref : f32) -> f32",
TextureKind::kDepth,
ast::SamplerKind::kComparisonSampler,
ast::TextureDimension::k2dArray,
TextureDataType::kF32,
"textureSampleCompareLevel",
[](ProgramBuilder* b) {
return b->ExprList("texture", // t
"sampler", // s
b->vec2<f32>(1_f, 2_f), // coords
3_i, // array_index
4_f); // depth_ref
},
},
{
ValidTextureOverload::kSampleCompareLevelDepth2dArrayOffsetF32,
"textureSampleCompareLevel(t : texture_depth_2d_array,\n"
" s : sampler_comparison,\n"
" coords : vec2<f32>,\n"
" array_index : i32,\n"
" depth_ref : f32,\n"
" offset : vec2<i32>) -> f32",
TextureKind::kDepth,
ast::SamplerKind::kComparisonSampler,
ast::TextureDimension::k2dArray,
TextureDataType::kF32,
"textureSampleCompareLevel",
[](ProgramBuilder* b) {
return b->ExprList("texture", // t
"sampler", // s
b->vec2<f32>(1_f, 2_f), // coords
3_i, // array_index
4_f, // depth_ref
b->vec2<i32>(5_i, 6_i)); // offset
},
},
{
ValidTextureOverload::kSampleCompareLevelDepthCubeF32,
"textureSampleCompareLevel(t : texture_depth_cube,\n"
" s : sampler_comparison,\n"
" coords : vec3<f32>,\n"
" depth_ref : f32) -> f32",
TextureKind::kDepth,
ast::SamplerKind::kComparisonSampler,
ast::TextureDimension::kCube,
TextureDataType::kF32,
"textureSampleCompareLevel",
[](ProgramBuilder* b) {
return b->ExprList("texture", // t
"sampler", // s
b->vec3<f32>(1_f, 2_f, 3_f), // coords
4_f); // depth_ref
},
},
{
ValidTextureOverload::kSampleCompareLevelDepthCubeArrayF32,
"textureSampleCompareLevel(t : texture_depth_cube_array,\n"
" s : sampler_comparison,\n"
" coords : vec3<f32>,\n"
" array_index : i32,\n"
" depth_ref : f32) -> f32",
TextureKind::kDepth,
ast::SamplerKind::kComparisonSampler,
ast::TextureDimension::kCubeArray,
TextureDataType::kF32,
"textureSampleCompareLevel",
[](ProgramBuilder* b) {
return b->ExprList("texture", // t
"sampler", // s
b->vec3<f32>(1_f, 2_f, 3_f), // coords
4_i, // array_index
5_f); // depth_ref
},
},
{
ValidTextureOverload::kLoad1dLevelF32,
"textureLoad(t : texture_1d<f32>,\n"
@ -2181,6 +2306,21 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
4_u); // level
},
},
{
ValidTextureOverload::kLoadDepthMultisampled2dF32,
"textureLoad(t : texture_depth_multisampled_2d,\n"
" coords : vec2<u32>,\n"
" sample_index : u32) -> f32",
TextureKind::kDepthMultisampled,
ast::TextureDimension::k2d,
TextureDataType::kF32,
"textureLoad",
[](ProgramBuilder* b) {
return b->ExprList("texture", // t
b->vec2<u32>(1_u, 2_u), // coords
3_u); // sample_index
},
},
{
ValidTextureOverload::kStoreWO1dRgba32float,
"textureStore(t : texture_storage_1d<rgba32float>,\n"

View File

@ -1618,7 +1618,7 @@ TEST_F(InspectorGetResourceBindingsTest, Simple) {
auto* st_type = MakeStorageTextureTypes(ast::TextureDimension::k2d, ast::TexelFormat::kR32Uint);
AddStorageTexture("st_var", st_type, 4, 0);
MakeStorageTextureBodyFunction("st_func", "st_var", ty.vec2<i32>(), utils::Empty);
MakeStorageTextureBodyFunction("st_func", "st_var", ty.vec2<u32>(), utils::Empty);
MakeCallerBodyFunction("ep_func",
utils::Vector{
@ -2789,14 +2789,14 @@ TEST_P(InspectorGetStorageTextureResourceBindingsTestWithParam, Simple) {
const ast::Type* dim_type = nullptr;
switch (dim) {
case ast::TextureDimension::k1d:
dim_type = ty.i32();
dim_type = ty.u32();
break;
case ast::TextureDimension::k2d:
case ast::TextureDimension::k2dArray:
dim_type = ty.vec2<i32>();
dim_type = ty.vec2<u32>();
break;
case ast::TextureDimension::k3d:
dim_type = ty.vec3<i32>();
dim_type = ty.vec3<u32>();
break;
default:
break;

View File

@ -554,33 +554,33 @@ fn unpack4x8snorm(u32) -> vec4<f32>
fn unpack4x8unorm(u32) -> vec4<f32>
@stage("compute") fn workgroupBarrier()
fn textureDimensions<T: fiu32>(texture: texture_1d<T>) -> i32
fn textureDimensions<T: fiu32, C: iu32>(texture: texture_1d<T>, level: C) -> i32
fn textureDimensions<T: fiu32>(texture: texture_2d<T>) -> vec2<i32>
fn textureDimensions<T: fiu32, C: iu32>(texture: texture_2d<T>, level: C) -> vec2<i32>
fn textureDimensions<T: fiu32>(texture: texture_2d_array<T>) -> vec2<i32>
fn textureDimensions<T: fiu32, C: iu32>(texture: texture_2d_array<T>, level: C) -> vec2<i32>
fn textureDimensions<T: fiu32>(texture: texture_3d<T>) -> vec3<i32>
fn textureDimensions<T: fiu32, C: iu32>(texture: texture_3d<T>, level: C) -> vec3<i32>
fn textureDimensions<T: fiu32>(texture: texture_cube<T>) -> vec2<i32>
fn textureDimensions<T: fiu32, C: iu32>(texture: texture_cube<T>, level: C) -> vec2<i32>
fn textureDimensions<T: fiu32>(texture: texture_cube_array<T>) -> vec2<i32>
fn textureDimensions<T: fiu32, C: iu32>(texture: texture_cube_array<T>, level: C) -> vec2<i32>
fn textureDimensions<T: fiu32>(texture: texture_multisampled_2d<T>) -> vec2<i32>
fn textureDimensions(texture: texture_depth_2d) -> vec2<i32>
fn textureDimensions<C: iu32>(texture: texture_depth_2d, level: C) -> vec2<i32>
fn textureDimensions(texture: texture_depth_2d_array) -> vec2<i32>
fn textureDimensions<C: iu32>(texture: texture_depth_2d_array, level: C) -> vec2<i32>
fn textureDimensions(texture: texture_depth_cube) -> vec2<i32>
fn textureDimensions<C: iu32>(texture: texture_depth_cube, level: C) -> vec2<i32>
fn textureDimensions(texture: texture_depth_cube_array) -> vec2<i32>
fn textureDimensions<C: iu32>(texture: texture_depth_cube_array, level: C) -> vec2<i32>
fn textureDimensions(texture: texture_depth_multisampled_2d) -> vec2<i32>
fn textureDimensions<F: texel_format, A: write>(texture: texture_storage_1d<F, A>) -> i32
fn textureDimensions<F: texel_format, A: write>(texture: texture_storage_2d<F, A>) -> vec2<i32>
fn textureDimensions<F: texel_format, A: write>(texture: texture_storage_2d_array<F, A>) -> vec2<i32>
fn textureDimensions<F: texel_format, A: write>(texture: texture_storage_3d<F, A>) -> vec3<i32>
fn textureDimensions(texture: texture_external) -> vec2<i32>
fn textureDimensions<T: fiu32>(texture: texture_1d<T>) -> u32
fn textureDimensions<T: fiu32, C: iu32>(texture: texture_1d<T>, level: C) -> u32
fn textureDimensions<T: fiu32>(texture: texture_2d<T>) -> vec2<u32>
fn textureDimensions<T: fiu32, C: iu32>(texture: texture_2d<T>, level: C) -> vec2<u32>
fn textureDimensions<T: fiu32>(texture: texture_2d_array<T>) -> vec2<u32>
fn textureDimensions<T: fiu32, C: iu32>(texture: texture_2d_array<T>, level: C) -> vec2<u32>
fn textureDimensions<T: fiu32>(texture: texture_3d<T>) -> vec3<u32>
fn textureDimensions<T: fiu32, C: iu32>(texture: texture_3d<T>, level: C) -> vec3<u32>
fn textureDimensions<T: fiu32>(texture: texture_cube<T>) -> vec2<u32>
fn textureDimensions<T: fiu32, C: iu32>(texture: texture_cube<T>, level: C) -> vec2<u32>
fn textureDimensions<T: fiu32>(texture: texture_cube_array<T>) -> vec2<u32>
fn textureDimensions<T: fiu32, C: iu32>(texture: texture_cube_array<T>, level: C) -> vec2<u32>
fn textureDimensions<T: fiu32>(texture: texture_multisampled_2d<T>) -> vec2<u32>
fn textureDimensions(texture: texture_depth_2d) -> vec2<u32>
fn textureDimensions<C: iu32>(texture: texture_depth_2d, level: C) -> vec2<u32>
fn textureDimensions(texture: texture_depth_2d_array) -> vec2<u32>
fn textureDimensions<C: iu32>(texture: texture_depth_2d_array, level: C) -> vec2<u32>
fn textureDimensions(texture: texture_depth_cube) -> vec2<u32>
fn textureDimensions<C: iu32>(texture: texture_depth_cube, level: C) -> vec2<u32>
fn textureDimensions(texture: texture_depth_cube_array) -> vec2<u32>
fn textureDimensions<C: iu32>(texture: texture_depth_cube_array, level: C) -> vec2<u32>
fn textureDimensions(texture: texture_depth_multisampled_2d) -> vec2<u32>
fn textureDimensions<F: texel_format, A: write>(texture: texture_storage_1d<F, A>) -> u32
fn textureDimensions<F: texel_format, A: write>(texture: texture_storage_2d<F, A>) -> vec2<u32>
fn textureDimensions<F: texel_format, A: write>(texture: texture_storage_2d_array<F, A>) -> vec2<u32>
fn textureDimensions<F: texel_format, A: write>(texture: texture_storage_3d<F, A>) -> vec3<u32>
fn textureDimensions(texture: texture_external) -> vec2<u32>
fn textureGather<T: fiu32, C: iu32>(@const component: C, texture: texture_2d<T>, sampler: sampler, coords: vec2<f32>) -> vec4<T>
fn textureGather<T: fiu32, C: iu32>(@const component: C, texture: texture_2d<T>, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>) -> vec4<T>
fn textureGather<T: fiu32, C: iu32>(@const component: C, texture: texture_2d_array<T>, sampler: sampler, coords: vec2<f32>, array_index: C) -> vec4<T>
@ -599,23 +599,23 @@ fn textureGatherCompare<C: iu32>(texture: texture_depth_2d_array, sampler: sampl
fn textureGatherCompare<C: iu32>(texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: C, depth_ref: f32, @const offset: vec2<i32>) -> vec4<f32>
fn textureGatherCompare(texture: texture_depth_cube, sampler: sampler_comparison, coords: vec3<f32>, depth_ref: f32) -> vec4<f32>
fn textureGatherCompare<C: iu32>(texture: texture_depth_cube_array, sampler: sampler_comparison, coords: vec3<f32>, array_index: C, depth_ref: f32) -> vec4<f32>
fn textureNumLayers<T: fiu32>(texture: texture_2d_array<T>) -> i32
fn textureNumLayers<T: fiu32>(texture: texture_cube_array<T>) -> i32
fn textureNumLayers(texture: texture_depth_2d_array) -> i32
fn textureNumLayers(texture: texture_depth_cube_array) -> i32
fn textureNumLayers<F: texel_format, A: write>(texture: texture_storage_2d_array<F, A>) -> i32
fn textureNumLevels<T: fiu32>(texture: texture_1d<T>) -> i32
fn textureNumLevels<T: fiu32>(texture: texture_2d<T>) -> i32
fn textureNumLevels<T: fiu32>(texture: texture_2d_array<T>) -> i32
fn textureNumLevels<T: fiu32>(texture: texture_3d<T>) -> i32
fn textureNumLevels<T: fiu32>(texture: texture_cube<T>) -> i32
fn textureNumLevels<T: fiu32>(texture: texture_cube_array<T>) -> i32
fn textureNumLevels(texture: texture_depth_2d) -> i32
fn textureNumLevels(texture: texture_depth_2d_array) -> i32
fn textureNumLevels(texture: texture_depth_cube) -> i32
fn textureNumLevels(texture: texture_depth_cube_array) -> i32
fn textureNumSamples<T: fiu32>(texture: texture_multisampled_2d<T>) -> i32
fn textureNumSamples(texture: texture_depth_multisampled_2d) -> i32
fn textureNumLayers<T: fiu32>(texture: texture_2d_array<T>) -> u32
fn textureNumLayers<T: fiu32>(texture: texture_cube_array<T>) -> u32
fn textureNumLayers(texture: texture_depth_2d_array) -> u32
fn textureNumLayers(texture: texture_depth_cube_array) -> u32
fn textureNumLayers<F: texel_format, A: write>(texture: texture_storage_2d_array<F, A>) -> u32
fn textureNumLevels<T: fiu32>(texture: texture_1d<T>) -> u32
fn textureNumLevels<T: fiu32>(texture: texture_2d<T>) -> u32
fn textureNumLevels<T: fiu32>(texture: texture_2d_array<T>) -> u32
fn textureNumLevels<T: fiu32>(texture: texture_3d<T>) -> u32
fn textureNumLevels<T: fiu32>(texture: texture_cube<T>) -> u32
fn textureNumLevels<T: fiu32>(texture: texture_cube_array<T>) -> u32
fn textureNumLevels(texture: texture_depth_2d) -> u32
fn textureNumLevels(texture: texture_depth_2d_array) -> u32
fn textureNumLevels(texture: texture_depth_cube) -> u32
fn textureNumLevels(texture: texture_depth_cube_array) -> u32
fn textureNumSamples<T: fiu32>(texture: texture_multisampled_2d<T>) -> u32
fn textureNumSamples(texture: texture_depth_multisampled_2d) -> u32
@stage("fragment") fn textureSample(texture: texture_1d<f32>, sampler: sampler, coords: f32) -> vec4<f32>
@stage("fragment") fn textureSample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>) -> vec4<f32>
@stage("fragment") fn textureSample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>) -> vec4<f32>

View File

@ -5709,7 +5709,7 @@ bool FunctionEmitter::EmitImageQuery(const spvtools::opt::Instruction& inst) {
Source{}, builder_.Symbols().Register("textureDimensions"));
ExpressionList dims_args{GetImageExpression(inst)};
if (opcode == SpvOpImageQuerySizeLod) {
dims_args.Push(ToI32(MakeOperand(inst, 1)).expr);
dims_args.Push(MakeOperand(inst, 1).expr);
}
const ast::Expression* dims_call =
create<ast::CallExpression>(Source{}, dims_ident, dims_args);
@ -5724,18 +5724,29 @@ bool FunctionEmitter::EmitImageQuery(const spvtools::opt::Instruction& inst) {
if (ast::IsTextureArray(dims)) {
auto* layers_ident = create<ast::IdentifierExpression>(
Source{}, builder_.Symbols().Register("textureNumLayers"));
exprs.Push(create<ast::CallExpression>(Source{}, layers_ident,
utils::Vector{GetImageExpression(inst)}));
auto num_layers = create<ast::CallExpression>(
Source{}, layers_ident, utils::Vector{GetImageExpression(inst)});
exprs.Push(num_layers);
}
auto* result_type = parser_impl_.ConvertType(inst.type_id());
auto* unsigned_type = ty_.AsUnsigned(result_type);
TypedExpression expr = {
result_type,
builder_.Construct(Source{}, result_type->Build(builder_), std::move(exprs))};
unsigned_type,
// If `exprs` holds multiple expressions, then these are the calls to
// textureDimensions() and textureNumLayers(), and these need to be placed into a
// vector initializer - otherwise, just emit the single expression to omit an
// unnecessary cast.
(exprs.Length() > 1)
? builder_.Construct(Source{}, unsigned_type->Build(builder_), std::move(exprs))
: exprs[0],
};
expr = ToSignedIfUnsigned(expr);
return EmitConstDefOrWriteToHoistedVar(inst, expr);
}
case SpvOpImageQueryLod:
return Fail() << "WGSL does not support querying the level of detail of "
"an image: "
return Fail() << "WGSL does not support querying the level of detail of an image: "
<< inst.PrettyPrint();
case SpvOpImageQueryLevels:
case SpvOpImageQuerySamples: {
@ -5746,9 +5757,10 @@ bool FunctionEmitter::EmitImageQuery(const spvtools::opt::Instruction& inst) {
const ast::Expression* ast_expr = create<ast::CallExpression>(
Source{}, levels_ident, utils::Vector{GetImageExpression(inst)});
auto* result_type = parser_impl_.ConvertType(inst.type_id());
// The SPIR-V result type must be integer scalar. The WGSL bulitin
// returns i32. If they aren't the same then convert the result.
if (!result_type->Is<I32>()) {
// The SPIR-V result type must be integer scalar.
// The WGSL bulitin returns u32.
// If they aren't the same then convert the result.
if (!result_type->Is<U32>()) {
ast_expr = builder_.Construct(Source{}, result_type->Build(builder_),
utils::Vector{ast_expr});
}

View File

@ -956,17 +956,15 @@ class FunctionEmitter {
ExpressionList MakeCoordinateOperandsForImageAccess(
const spvtools::opt::Instruction& image_access);
/// Returns the given value as an I32. If it's already an I32 then this
/// return the given value. Otherwise, wrap the value in a TypeInitializer
/// expression.
/// Returns the given value as an i32. If it's already an i32 then simply returns @p value.
/// Otherwise, wrap the value in a TypeInitializer expression.
/// @param value the value to pass through or convert
/// @returns the value as an I32 value.
/// @returns the value as an i32 value.
TypedExpression ToI32(TypedExpression value);
/// Returns the given value as a signed integer type of the same shape
/// if the value is unsigned scalar or vector, by wrapping the value
/// with a TypeInitializer expression. Returns the value itself if the
/// value otherwise.
/// Returns the given value as a signed integer type of the same shape if the value is unsigned
/// scalar or vector, by wrapping the value with a TypeInitializer expression. Returns the
/// value itself if the value was already signed.
/// @param value the value to pass through or convert
/// @returns the value itself, or converted to signed integral
TypedExpression ToSignedIfUnsigned(TypedExpression value);

View File

@ -2842,7 +2842,7 @@ INSTANTIATE_TEST_SUITE_P(
"%99 = OpImageQuerySize %v3int %im \n"
"%98 = OpImageRead %v4float %im %vi123\n",
R"(@group(2) @binding(1) var x_20 : texture_2d_array<f32>;)",
R"(let x_99 : vec3<i32> = vec3<i32>(textureDimensions(x_20), textureNumLayers(x_20));)"}
R"(let x_99 : vec3<i32> = vec3<i32>(vec3<u32>(textureDimensions(x_20), textureNumLayers(x_20)));)"}
// 3D array storage image doesn't exist.
// Multisampled array
@ -2898,7 +2898,7 @@ INSTANTIATE_TEST_SUITE_P(
// 2D array
{"%float 2D 0 1 0 1 Unknown", "%99 = OpImageQuerySizeLod %v3int %im %i1\n",
R"(@group(2) @binding(1) var x_20 : texture_2d_array<f32>;)",
R"(let x_99 : vec3<i32> = vec3<i32>(textureDimensions(x_20, i1), textureNumLayers(x_20));)"},
R"(let x_99 : vec3<i32> = vec3<i32>(vec3<u32>(textureDimensions(x_20, i1), textureNumLayers(x_20)));)"},
// There is no 3D array
@ -2909,12 +2909,12 @@ INSTANTIATE_TEST_SUITE_P(
// https://github.com/gpuweb/gpuweb/issues/1345
{"%float Cube 0 1 0 1 Unknown", "%99 = OpImageQuerySizeLod %v3int %im %i1\n",
R"(@group(2) @binding(1) var x_20 : texture_cube_array<f32>;)",
R"(let x_99 : vec3<i32> = vec3<i32>(textureDimensions(x_20, i1).xy, textureNumLayers(x_20));)"},
R"(let x_99 : vec3<i32> = vec3<i32>(vec3<u32>(textureDimensions(x_20, i1).xy, textureNumLayers(x_20)));)"},
// Depth 2D array
{"%float 2D 1 1 0 1 Unknown", "%99 = OpImageQuerySizeLod %v3int %im %i1\n",
R"(@group(2) @binding(1) var x_20 : texture_depth_2d_array;)",
R"(let x_99 : vec3<i32> = vec3<i32>(textureDimensions(x_20, i1), textureNumLayers(x_20));)"},
R"(let x_99 : vec3<i32> = vec3<i32>(vec3<u32>(textureDimensions(x_20, i1), textureNumLayers(x_20)));)"},
// Depth Cube Array
//
@ -2923,19 +2923,17 @@ INSTANTIATE_TEST_SUITE_P(
// https://github.com/gpuweb/gpuweb/issues/1345
{"%float Cube 1 1 0 1 Unknown", "%99 = OpImageQuerySizeLod %v3int %im %i1\n",
R"(@group(2) @binding(1) var x_20 : texture_depth_cube_array;)",
R"(let x_99 : vec3<i32> = vec3<i32>(textureDimensions(x_20, i1).xy, textureNumLayers(x_20));)"}}));
R"(let x_99 : vec3<i32> = vec3<i32>(vec3<u32>(textureDimensions(x_20, i1).xy, textureNumLayers(x_20)));)"}}));
INSTANTIATE_TEST_SUITE_P(
// When the level-of-detail value is given as an unsigned
// integer, we must convert it before using it as an argument
// to textureDimensions.
// textureDimensions accepts both signed and unsigned the level-of-detail values.
ImageQuerySizeLod_NonArrayed_SignedResult_UnsignedLevel,
SpvParserHandleTest_SampledImageAccessTest,
::testing::ValuesIn(std::vector<ImageAccessCase>{
{"%float 1D 0 0 0 1 Unknown", "%99 = OpImageQuerySizeLod %int %im %u1\n",
R"(@group(2) @binding(1) var x_20 : texture_1d<f32>;)",
R"(let x_99 : i32 = i32(textureDimensions(x_20, i32(u1)));)"}}));
R"(let x_99 : i32 = i32(textureDimensions(x_20, u1));)"}}));
INSTANTIATE_TEST_SUITE_P(
// When SPIR-V wants the result type to be unsigned, we have to
@ -2948,7 +2946,7 @@ INSTANTIATE_TEST_SUITE_P(
{"%float 1D 0 0 0 1 Unknown", "%99 = OpImageQuerySizeLod %uint %im %i1\n",
R"(@group(2) @binding(1) var x_20 : texture_1d<f32>;)",
R"(let x_99 : u32 = u32(textureDimensions(x_20, i1));)"}}));
R"(let x_99 : i32 = i32(textureDimensions(x_20, i1));)"}}));
INSTANTIATE_TEST_SUITE_P(ImageQueryLevels_SignedResult,
SpvParserHandleTest_SampledImageAccessTest,
@ -2961,47 +2959,47 @@ INSTANTIATE_TEST_SUITE_P(ImageQueryLevels_SignedResult,
// 2D
{"%float 2D 0 0 0 1 Unknown", "%99 = OpImageQueryLevels %int %im\n",
R"(@group(2) @binding(1) var x_20 : texture_2d<f32>;)",
R"(let x_99 : i32 = textureNumLevels(x_20);)"},
R"(let x_99 : i32 = i32(textureNumLevels(x_20));)"},
// 2D array
{"%float 2D 0 1 0 1 Unknown", "%99 = OpImageQueryLevels %int %im\n",
R"(@group(2) @binding(1) var x_20 : texture_2d_array<f32>;)",
R"(let x_99 : i32 = textureNumLevels(x_20);)"},
R"(let x_99 : i32 = i32(textureNumLevels(x_20));)"},
// 3D
{"%float 3D 0 0 0 1 Unknown", "%99 = OpImageQueryLevels %int %im\n",
R"(@group(2) @binding(1) var x_20 : texture_3d<f32>;)",
R"(let x_99 : i32 = textureNumLevels(x_20);)"},
R"(let x_99 : i32 = i32(textureNumLevels(x_20));)"},
// Cube
{"%float Cube 0 0 0 1 Unknown", "%99 = OpImageQueryLevels %int %im\n",
R"(@group(2) @binding(1) var x_20 : texture_cube<f32>;)",
R"(let x_99 : i32 = textureNumLevels(x_20);)"},
R"(let x_99 : i32 = i32(textureNumLevels(x_20));)"},
// Cube array
{"%float Cube 0 1 0 1 Unknown", "%99 = OpImageQueryLevels %int %im\n",
R"(@group(2) @binding(1) var x_20 : texture_cube_array<f32>;)",
R"(let x_99 : i32 = textureNumLevels(x_20);)"},
R"(let x_99 : i32 = i32(textureNumLevels(x_20));)"},
// depth 2d
{"%float 2D 1 0 0 1 Unknown", "%99 = OpImageQueryLevels %int %im\n",
R"(@group(2) @binding(1) var x_20 : texture_depth_2d;)",
R"(let x_99 : i32 = textureNumLevels(x_20);)"},
R"(let x_99 : i32 = i32(textureNumLevels(x_20));)"},
// depth 2d array
{"%float 2D 1 1 0 1 Unknown", "%99 = OpImageQueryLevels %int %im\n",
R"(@group(2) @binding(1) var x_20 : texture_depth_2d_array;)",
R"(let x_99 : i32 = textureNumLevels(x_20);)"},
R"(let x_99 : i32 = i32(textureNumLevels(x_20));)"},
// depth cube
{"%float Cube 1 0 0 1 Unknown", "%99 = OpImageQueryLevels %int %im\n",
R"(@group(2) @binding(1) var x_20 : texture_depth_cube;)",
R"(let x_99 : i32 = textureNumLevels(x_20);)"},
R"(let x_99 : i32 = i32(textureNumLevels(x_20));)"},
// depth cube array
{"%float Cube 1 1 0 1 Unknown", "%99 = OpImageQueryLevels %int %im\n",
R"(@group(2) @binding(1) var x_20 : texture_depth_cube_array;)",
R"(let x_99 : i32 = textureNumLevels(x_20);)"}}));
R"(let x_99 : i32 = i32(textureNumLevels(x_20));)"}}));
INSTANTIATE_TEST_SUITE_P(
// Spot check that a type conversion is inserted when SPIR-V asks for
@ -3011,7 +3009,7 @@ INSTANTIATE_TEST_SUITE_P(
::testing::ValuesIn(std::vector<ImageAccessCase>{
{"%float 2D 0 0 0 1 Unknown", "%99 = OpImageQueryLevels %uint %im\n",
R"(@group(2) @binding(1) var x_20 : texture_2d<f32>;)",
R"(let x_99 : u32 = u32(textureNumLevels(x_20));)"}}));
R"(let x_99 : u32 = textureNumLevels(x_20);)"}}));
INSTANTIATE_TEST_SUITE_P(ImageQuerySamples_SignedResult,
SpvParserHandleTest_SampledImageAccessTest,
@ -3019,21 +3017,21 @@ INSTANTIATE_TEST_SUITE_P(ImageQuerySamples_SignedResult,
// Multsample 2D
{"%float 2D 0 0 1 1 Unknown", "%99 = OpImageQuerySamples %int %im\n",
R"(@group(2) @binding(1) var x_20 : texture_multisampled_2d<f32>;)",
R"(let x_99 : i32 = textureNumSamples(x_20);)"} // namespace
R"(let x_99 : i32 = i32(textureNumSamples(x_20));)"}
// Multisample 2D array
// Not in WebGPU
}));
INSTANTIATE_TEST_SUITE_P(
// Translation must inject a type coersion from signed to unsigned.
// Translation must inject a type coersion from unsigned to signed.
ImageQuerySamples_UnsignedResult,
SpvParserHandleTest_SampledImageAccessTest,
::testing::ValuesIn(std::vector<ImageAccessCase>{
// Multsample 2D
// Multisample 2D
{"%float 2D 0 0 1 1 Unknown", "%99 = OpImageQuerySamples %uint %im\n",
R"(@group(2) @binding(1) var x_20 : texture_multisampled_2d<f32>;)",
R"(let x_99 : u32 = u32(textureNumSamples(x_20));)"}
R"(let x_99 : u32 = textureNumSamples(x_20);)"}
// Multisample 2D array
// Not in WebGPU

View File

@ -444,6 +444,20 @@ const spirv::I32* TypeManager::I32() {
return state->i32_;
}
const Type* TypeManager::AsUnsigned(const Type* ty) {
return Switch(
ty, //
[&](const spirv::I32*) { return U32(); }, //
[&](const spirv::U32*) { return ty; }, //
[&](const spirv::Vector* vec) {
return Switch(
vec->type, //
[&](const spirv::I32*) { return Vector(U32(), vec->size); }, //
[&](const spirv::U32*) { return ty; } //
);
});
}
const spirv::Pointer* TypeManager::Pointer(const Type* el,
ast::AddressSpace address_space,
ast::Access access) {

View File

@ -539,6 +539,11 @@ class TypeManager {
const spirv::F32* F32();
/// @return a I32 type. Repeated calls will return the same pointer.
const spirv::I32* I32();
/// @param ty the input type.
/// @returns the equivalent unsigned integer scalar or vector if @p ty is a scalar or vector,
/// otherwise nullptr.
const Type* AsUnsigned(const Type* ty);
/// @param ty the store type
/// @param address_space the pointer address space
/// @param access the declared access mode

View File

@ -2416,17 +2416,17 @@ static const char* expected_texture_overload(ast::builtin::test::ValidTextureOve
case ValidTextureOverload::kSampleCompareDepthCubeArrayF32:
return R"(textureSampleCompare(texture, sampler, coords, array_index, depth_ref))";
case ValidTextureOverload::kSampleCompareLevelDepth2dF32:
return R"(textureSampleCompare(texture, sampler, coords, depth_ref))";
return R"(textureSampleCompareLevel(texture, sampler, coords, depth_ref))";
case ValidTextureOverload::kSampleCompareLevelDepth2dOffsetF32:
return R"(textureSampleCompare(texture, sampler, coords, depth_ref, offset))";
return R"(textureSampleCompareLevel(texture, sampler, coords, depth_ref, offset))";
case ValidTextureOverload::kSampleCompareLevelDepth2dArrayF32:
return R"(textureSampleCompare(texture, sampler, coords, array_index, depth_ref))";
return R"(textureSampleCompareLevel(texture, sampler, coords, array_index, depth_ref))";
case ValidTextureOverload::kSampleCompareLevelDepth2dArrayOffsetF32:
return R"(textureSampleCompare(texture, sampler, coords, array_index, depth_ref, offset))";
return R"(textureSampleCompareLevel(texture, sampler, coords, array_index, depth_ref, offset))";
case ValidTextureOverload::kSampleCompareLevelDepthCubeF32:
return R"(textureSampleCompare(texture, sampler, coords, depth_ref))";
return R"(textureSampleCompareLevel(texture, sampler, coords, depth_ref))";
case ValidTextureOverload::kSampleCompareLevelDepthCubeArrayF32:
return R"(textureSampleCompare(texture, sampler, coords, array_index, depth_ref))";
return R"(textureSampleCompareLevel(texture, sampler, coords, array_index, depth_ref))";
case ValidTextureOverload::kLoad1dLevelF32:
case ValidTextureOverload::kLoad1dLevelU32:
case ValidTextureOverload::kLoad1dLevelI32:
@ -2478,7 +2478,7 @@ TEST_P(ResolverBuiltinTest_Texture, Call) {
default:
FAIL() << "invalid texture dimensions: " << param.texture_dimension;
case ast::TextureDimension::k1d:
EXPECT_TRUE(TypeOf(call)->Is<sem::I32>());
EXPECT_TRUE(TypeOf(call)->Is<sem::U32>());
break;
case ast::TextureDimension::k2d:
case ast::TextureDimension::k2dArray:
@ -2487,23 +2487,23 @@ TEST_P(ResolverBuiltinTest_Texture, Call) {
auto* vec = As<sem::Vector>(TypeOf(call));
ASSERT_NE(vec, nullptr);
EXPECT_EQ(vec->Width(), 2u);
EXPECT_TRUE(vec->type()->Is<sem::I32>());
EXPECT_TRUE(vec->type()->Is<sem::U32>());
break;
}
case ast::TextureDimension::k3d: {
auto* vec = As<sem::Vector>(TypeOf(call));
ASSERT_NE(vec, nullptr);
EXPECT_EQ(vec->Width(), 3u);
EXPECT_TRUE(vec->type()->Is<sem::I32>());
EXPECT_TRUE(vec->type()->Is<sem::U32>());
break;
}
}
} else if (std::string(param.function) == "textureNumLayers") {
EXPECT_TRUE(TypeOf(call)->Is<sem::I32>());
EXPECT_TRUE(TypeOf(call)->Is<sem::U32>());
} else if (std::string(param.function) == "textureNumLevels") {
EXPECT_TRUE(TypeOf(call)->Is<sem::I32>());
EXPECT_TRUE(TypeOf(call)->Is<sem::U32>());
} else if (std::string(param.function) == "textureNumSamples") {
EXPECT_TRUE(TypeOf(call)->Is<sem::I32>());
EXPECT_TRUE(TypeOf(call)->Is<sem::U32>());
} else if (std::string(param.function) == "textureStore") {
EXPECT_TRUE(TypeOf(call)->Is<sem::Void>());
} else if (std::string(param.function) == "textureGather") {

View File

@ -8276,7 +8276,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[0],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[839],
/* return matcher indices */ &kMatcherIndices[34],
/* return matcher indices */ &kMatcherIndices[35],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -8288,7 +8288,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[0],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[623],
/* return matcher indices */ &kMatcherIndices[34],
/* return matcher indices */ &kMatcherIndices[35],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -8300,7 +8300,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[0],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[838],
/* return matcher indices */ &kMatcherIndices[126],
/* return matcher indices */ &kMatcherIndices[124],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -8312,7 +8312,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[0],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[627],
/* return matcher indices */ &kMatcherIndices[126],
/* return matcher indices */ &kMatcherIndices[124],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -8324,7 +8324,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[0],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[837],
/* return matcher indices */ &kMatcherIndices[126],
/* return matcher indices */ &kMatcherIndices[124],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -8336,7 +8336,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[0],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[631],
/* return matcher indices */ &kMatcherIndices[126],
/* return matcher indices */ &kMatcherIndices[124],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -8348,7 +8348,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[0],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[836],
/* return matcher indices */ &kMatcherIndices[108],
/* return matcher indices */ &kMatcherIndices[114],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -8360,7 +8360,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[0],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[635],
/* return matcher indices */ &kMatcherIndices[108],
/* return matcher indices */ &kMatcherIndices[114],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -8372,7 +8372,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[0],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[835],
/* return matcher indices */ &kMatcherIndices[126],
/* return matcher indices */ &kMatcherIndices[124],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -8384,7 +8384,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[0],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[639],
/* return matcher indices */ &kMatcherIndices[126],
/* return matcher indices */ &kMatcherIndices[124],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -8396,7 +8396,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[0],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[834],
/* return matcher indices */ &kMatcherIndices[126],
/* return matcher indices */ &kMatcherIndices[124],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -8408,7 +8408,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[0],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[643],
/* return matcher indices */ &kMatcherIndices[126],
/* return matcher indices */ &kMatcherIndices[124],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -8420,7 +8420,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[0],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[833],
/* return matcher indices */ &kMatcherIndices[126],
/* return matcher indices */ &kMatcherIndices[124],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -8432,7 +8432,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[28],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[832],
/* return matcher indices */ &kMatcherIndices[126],
/* return matcher indices */ &kMatcherIndices[124],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -8444,7 +8444,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[1],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[687],
/* return matcher indices */ &kMatcherIndices[126],
/* return matcher indices */ &kMatcherIndices[124],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -8456,7 +8456,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[28],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[831],
/* return matcher indices */ &kMatcherIndices[126],
/* return matcher indices */ &kMatcherIndices[124],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -8468,7 +8468,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[1],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[653],
/* return matcher indices */ &kMatcherIndices[126],
/* return matcher indices */ &kMatcherIndices[124],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -8480,7 +8480,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[28],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[830],
/* return matcher indices */ &kMatcherIndices[126],
/* return matcher indices */ &kMatcherIndices[124],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -8492,7 +8492,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[1],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[657],
/* return matcher indices */ &kMatcherIndices[126],
/* return matcher indices */ &kMatcherIndices[124],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -8504,7 +8504,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[28],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[829],
/* return matcher indices */ &kMatcherIndices[126],
/* return matcher indices */ &kMatcherIndices[124],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -8516,7 +8516,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[1],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[605],
/* return matcher indices */ &kMatcherIndices[126],
/* return matcher indices */ &kMatcherIndices[124],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -8528,7 +8528,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[28],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[828],
/* return matcher indices */ &kMatcherIndices[126],
/* return matcher indices */ &kMatcherIndices[124],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -8540,7 +8540,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[28],
/* template numbers */ &kTemplateNumbers[3],
/* parameters */ &kParameters[827],
/* return matcher indices */ &kMatcherIndices[34],
/* return matcher indices */ &kMatcherIndices[35],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -8552,7 +8552,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[28],
/* template numbers */ &kTemplateNumbers[3],
/* parameters */ &kParameters[826],
/* return matcher indices */ &kMatcherIndices[126],
/* return matcher indices */ &kMatcherIndices[124],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -8564,7 +8564,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[28],
/* template numbers */ &kTemplateNumbers[3],
/* parameters */ &kParameters[825],
/* return matcher indices */ &kMatcherIndices[126],
/* return matcher indices */ &kMatcherIndices[124],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -8576,7 +8576,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[28],
/* template numbers */ &kTemplateNumbers[3],
/* parameters */ &kParameters[824],
/* return matcher indices */ &kMatcherIndices[108],
/* return matcher indices */ &kMatcherIndices[114],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -8588,7 +8588,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[28],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[823],
/* return matcher indices */ &kMatcherIndices[126],
/* return matcher indices */ &kMatcherIndices[124],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -9560,7 +9560,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[0],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[817],
/* return matcher indices */ &kMatcherIndices[34],
/* return matcher indices */ &kMatcherIndices[35],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -9572,7 +9572,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[0],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[816],
/* return matcher indices */ &kMatcherIndices[34],
/* return matcher indices */ &kMatcherIndices[35],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -9584,7 +9584,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[0],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[815],
/* return matcher indices */ &kMatcherIndices[34],
/* return matcher indices */ &kMatcherIndices[35],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -9596,7 +9596,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[0],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[814],
/* return matcher indices */ &kMatcherIndices[34],
/* return matcher indices */ &kMatcherIndices[35],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -9608,7 +9608,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[0],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[813],
/* return matcher indices */ &kMatcherIndices[34],
/* return matcher indices */ &kMatcherIndices[35],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -9620,7 +9620,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[0],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[812],
/* return matcher indices */ &kMatcherIndices[34],
/* return matcher indices */ &kMatcherIndices[35],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -9632,7 +9632,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[28],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[811],
/* return matcher indices */ &kMatcherIndices[34],
/* return matcher indices */ &kMatcherIndices[35],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -9644,7 +9644,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[28],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[810],
/* return matcher indices */ &kMatcherIndices[34],
/* return matcher indices */ &kMatcherIndices[35],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -9656,7 +9656,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[28],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[809],
/* return matcher indices */ &kMatcherIndices[34],
/* return matcher indices */ &kMatcherIndices[35],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -9668,7 +9668,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[28],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[808],
/* return matcher indices */ &kMatcherIndices[34],
/* return matcher indices */ &kMatcherIndices[35],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -11180,7 +11180,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[0],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[822],
/* return matcher indices */ &kMatcherIndices[34],
/* return matcher indices */ &kMatcherIndices[35],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -11192,7 +11192,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[0],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[821],
/* return matcher indices */ &kMatcherIndices[34],
/* return matcher indices */ &kMatcherIndices[35],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -11204,7 +11204,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[28],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[820],
/* return matcher indices */ &kMatcherIndices[34],
/* return matcher indices */ &kMatcherIndices[35],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -11216,7 +11216,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[28],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[819],
/* return matcher indices */ &kMatcherIndices[34],
/* return matcher indices */ &kMatcherIndices[35],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -11228,7 +11228,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[28],
/* template numbers */ &kTemplateNumbers[3],
/* parameters */ &kParameters[818],
/* return matcher indices */ &kMatcherIndices[34],
/* return matcher indices */ &kMatcherIndices[35],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -12932,7 +12932,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[0],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[988],
/* return matcher indices */ &kMatcherIndices[34],
/* return matcher indices */ &kMatcherIndices[35],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -12944,7 +12944,7 @@ constexpr OverloadInfo kOverloads[] = {
/* template types */ &kTemplateTypes[28],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[807],
/* return matcher indices */ &kMatcherIndices[34],
/* return matcher indices */ &kMatcherIndices[35],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* const eval */ nullptr,
},
@ -14499,33 +14499,33 @@ constexpr IntrinsicInfo kBuiltins[] = {
},
{
/* [85] */
/* fn textureDimensions<T : fiu32>(texture: texture_1d<T>) -> i32 */
/* fn textureDimensions<T : fiu32, C : iu32>(texture: texture_1d<T>, level: C) -> i32 */
/* fn textureDimensions<T : fiu32>(texture: texture_2d<T>) -> vec2<i32> */
/* fn textureDimensions<T : fiu32, C : iu32>(texture: texture_2d<T>, level: C) -> vec2<i32> */
/* fn textureDimensions<T : fiu32>(texture: texture_2d_array<T>) -> vec2<i32> */
/* fn textureDimensions<T : fiu32, C : iu32>(texture: texture_2d_array<T>, level: C) -> vec2<i32> */
/* fn textureDimensions<T : fiu32>(texture: texture_3d<T>) -> vec3<i32> */
/* fn textureDimensions<T : fiu32, C : iu32>(texture: texture_3d<T>, level: C) -> vec3<i32> */
/* fn textureDimensions<T : fiu32>(texture: texture_cube<T>) -> vec2<i32> */
/* fn textureDimensions<T : fiu32, C : iu32>(texture: texture_cube<T>, level: C) -> vec2<i32> */
/* fn textureDimensions<T : fiu32>(texture: texture_cube_array<T>) -> vec2<i32> */
/* fn textureDimensions<T : fiu32, C : iu32>(texture: texture_cube_array<T>, level: C) -> vec2<i32> */
/* fn textureDimensions<T : fiu32>(texture: texture_multisampled_2d<T>) -> vec2<i32> */
/* fn textureDimensions(texture: texture_depth_2d) -> vec2<i32> */
/* fn textureDimensions<C : iu32>(texture: texture_depth_2d, level: C) -> vec2<i32> */
/* fn textureDimensions(texture: texture_depth_2d_array) -> vec2<i32> */
/* fn textureDimensions<C : iu32>(texture: texture_depth_2d_array, level: C) -> vec2<i32> */
/* fn textureDimensions(texture: texture_depth_cube) -> vec2<i32> */
/* fn textureDimensions<C : iu32>(texture: texture_depth_cube, level: C) -> vec2<i32> */
/* fn textureDimensions(texture: texture_depth_cube_array) -> vec2<i32> */
/* fn textureDimensions<C : iu32>(texture: texture_depth_cube_array, level: C) -> vec2<i32> */
/* fn textureDimensions(texture: texture_depth_multisampled_2d) -> vec2<i32> */
/* fn textureDimensions<F : texel_format, A : write>(texture: texture_storage_1d<F, A>) -> i32 */
/* fn textureDimensions<F : texel_format, A : write>(texture: texture_storage_2d<F, A>) -> vec2<i32> */
/* fn textureDimensions<F : texel_format, A : write>(texture: texture_storage_2d_array<F, A>) -> vec2<i32> */
/* fn textureDimensions<F : texel_format, A : write>(texture: texture_storage_3d<F, A>) -> vec3<i32> */
/* fn textureDimensions(texture: texture_external) -> vec2<i32> */
/* fn textureDimensions<T : fiu32>(texture: texture_1d<T>) -> u32 */
/* fn textureDimensions<T : fiu32, C : iu32>(texture: texture_1d<T>, level: C) -> u32 */
/* fn textureDimensions<T : fiu32>(texture: texture_2d<T>) -> vec2<u32> */
/* fn textureDimensions<T : fiu32, C : iu32>(texture: texture_2d<T>, level: C) -> vec2<u32> */
/* fn textureDimensions<T : fiu32>(texture: texture_2d_array<T>) -> vec2<u32> */
/* fn textureDimensions<T : fiu32, C : iu32>(texture: texture_2d_array<T>, level: C) -> vec2<u32> */
/* fn textureDimensions<T : fiu32>(texture: texture_3d<T>) -> vec3<u32> */
/* fn textureDimensions<T : fiu32, C : iu32>(texture: texture_3d<T>, level: C) -> vec3<u32> */
/* fn textureDimensions<T : fiu32>(texture: texture_cube<T>) -> vec2<u32> */
/* fn textureDimensions<T : fiu32, C : iu32>(texture: texture_cube<T>, level: C) -> vec2<u32> */
/* fn textureDimensions<T : fiu32>(texture: texture_cube_array<T>) -> vec2<u32> */
/* fn textureDimensions<T : fiu32, C : iu32>(texture: texture_cube_array<T>, level: C) -> vec2<u32> */
/* fn textureDimensions<T : fiu32>(texture: texture_multisampled_2d<T>) -> vec2<u32> */
/* fn textureDimensions(texture: texture_depth_2d) -> vec2<u32> */
/* fn textureDimensions<C : iu32>(texture: texture_depth_2d, level: C) -> vec2<u32> */
/* fn textureDimensions(texture: texture_depth_2d_array) -> vec2<u32> */
/* fn textureDimensions<C : iu32>(texture: texture_depth_2d_array, level: C) -> vec2<u32> */
/* fn textureDimensions(texture: texture_depth_cube) -> vec2<u32> */
/* fn textureDimensions<C : iu32>(texture: texture_depth_cube, level: C) -> vec2<u32> */
/* fn textureDimensions(texture: texture_depth_cube_array) -> vec2<u32> */
/* fn textureDimensions<C : iu32>(texture: texture_depth_cube_array, level: C) -> vec2<u32> */
/* fn textureDimensions(texture: texture_depth_multisampled_2d) -> vec2<u32> */
/* fn textureDimensions<F : texel_format, A : write>(texture: texture_storage_1d<F, A>) -> u32 */
/* fn textureDimensions<F : texel_format, A : write>(texture: texture_storage_2d<F, A>) -> vec2<u32> */
/* fn textureDimensions<F : texel_format, A : write>(texture: texture_storage_2d_array<F, A>) -> vec2<u32> */
/* fn textureDimensions<F : texel_format, A : write>(texture: texture_storage_3d<F, A>) -> vec3<u32> */
/* fn textureDimensions(texture: texture_external) -> vec2<u32> */
/* num overloads */ 27,
/* overloads */ &kOverloads[0],
},
@ -14559,33 +14559,33 @@ constexpr IntrinsicInfo kBuiltins[] = {
},
{
/* [88] */
/* fn textureNumLayers<T : fiu32>(texture: texture_2d_array<T>) -> i32 */
/* fn textureNumLayers<T : fiu32>(texture: texture_cube_array<T>) -> i32 */
/* fn textureNumLayers(texture: texture_depth_2d_array) -> i32 */
/* fn textureNumLayers(texture: texture_depth_cube_array) -> i32 */
/* fn textureNumLayers<F : texel_format, A : write>(texture: texture_storage_2d_array<F, A>) -> i32 */
/* fn textureNumLayers<T : fiu32>(texture: texture_2d_array<T>) -> u32 */
/* fn textureNumLayers<T : fiu32>(texture: texture_cube_array<T>) -> u32 */
/* fn textureNumLayers(texture: texture_depth_2d_array) -> u32 */
/* fn textureNumLayers(texture: texture_depth_cube_array) -> u32 */
/* fn textureNumLayers<F : texel_format, A : write>(texture: texture_storage_2d_array<F, A>) -> u32 */
/* num overloads */ 5,
/* overloads */ &kOverloads[242],
},
{
/* [89] */
/* fn textureNumLevels<T : fiu32>(texture: texture_1d<T>) -> i32 */
/* fn textureNumLevels<T : fiu32>(texture: texture_2d<T>) -> i32 */
/* fn textureNumLevels<T : fiu32>(texture: texture_2d_array<T>) -> i32 */
/* fn textureNumLevels<T : fiu32>(texture: texture_3d<T>) -> i32 */
/* fn textureNumLevels<T : fiu32>(texture: texture_cube<T>) -> i32 */
/* fn textureNumLevels<T : fiu32>(texture: texture_cube_array<T>) -> i32 */
/* fn textureNumLevels(texture: texture_depth_2d) -> i32 */
/* fn textureNumLevels(texture: texture_depth_2d_array) -> i32 */
/* fn textureNumLevels(texture: texture_depth_cube) -> i32 */
/* fn textureNumLevels(texture: texture_depth_cube_array) -> i32 */
/* fn textureNumLevels<T : fiu32>(texture: texture_1d<T>) -> u32 */
/* fn textureNumLevels<T : fiu32>(texture: texture_2d<T>) -> u32 */
/* fn textureNumLevels<T : fiu32>(texture: texture_2d_array<T>) -> u32 */
/* fn textureNumLevels<T : fiu32>(texture: texture_3d<T>) -> u32 */
/* fn textureNumLevels<T : fiu32>(texture: texture_cube<T>) -> u32 */
/* fn textureNumLevels<T : fiu32>(texture: texture_cube_array<T>) -> u32 */
/* fn textureNumLevels(texture: texture_depth_2d) -> u32 */
/* fn textureNumLevels(texture: texture_depth_2d_array) -> u32 */
/* fn textureNumLevels(texture: texture_depth_cube) -> u32 */
/* fn textureNumLevels(texture: texture_depth_cube_array) -> u32 */
/* num overloads */ 10,
/* overloads */ &kOverloads[107],
},
{
/* [90] */
/* fn textureNumSamples<T : fiu32>(texture: texture_multisampled_2d<T>) -> i32 */
/* fn textureNumSamples(texture: texture_depth_multisampled_2d) -> i32 */
/* fn textureNumSamples<T : fiu32>(texture: texture_multisampled_2d<T>) -> u32 */
/* fn textureNumSamples(texture: texture_depth_multisampled_2d) -> u32 */
/* num overloads */ 2,
/* overloads */ &kOverloads[388],
},

View File

@ -596,33 +596,33 @@ TEST_F(IntrinsicTableTest, OverloadOrderByNumberOfParameters) {
R"(error: no matching call to textureDimensions(bool, bool)
27 candidate functions:
textureDimensions(texture: texture_1d<T>, level: C) -> i32 where: T is f32, i32 or u32, C is i32 or u32
textureDimensions(texture: texture_2d<T>, level: C) -> vec2<i32> where: T is f32, i32 or u32, C is i32 or u32
textureDimensions(texture: texture_2d_array<T>, level: C) -> vec2<i32> where: T is f32, i32 or u32, C is i32 or u32
textureDimensions(texture: texture_3d<T>, level: C) -> vec3<i32> where: T is f32, i32 or u32, C is i32 or u32
textureDimensions(texture: texture_cube<T>, level: C) -> vec2<i32> where: T is f32, i32 or u32, C is i32 or u32
textureDimensions(texture: texture_cube_array<T>, level: C) -> vec2<i32> where: T is f32, i32 or u32, C is i32 or u32
textureDimensions(texture: texture_depth_2d, level: C) -> vec2<i32> where: C is i32 or u32
textureDimensions(texture: texture_depth_2d_array, level: C) -> vec2<i32> where: C is i32 or u32
textureDimensions(texture: texture_depth_cube, level: C) -> vec2<i32> where: C is i32 or u32
textureDimensions(texture: texture_depth_cube_array, level: C) -> vec2<i32> where: C is i32 or u32
textureDimensions(texture: texture_1d<T>) -> i32 where: T is f32, i32 or u32
textureDimensions(texture: texture_2d<T>) -> vec2<i32> where: T is f32, i32 or u32
textureDimensions(texture: texture_2d_array<T>) -> vec2<i32> where: T is f32, i32 or u32
textureDimensions(texture: texture_3d<T>) -> vec3<i32> where: T is f32, i32 or u32
textureDimensions(texture: texture_cube<T>) -> vec2<i32> where: T is f32, i32 or u32
textureDimensions(texture: texture_cube_array<T>) -> vec2<i32> where: T is f32, i32 or u32
textureDimensions(texture: texture_multisampled_2d<T>) -> vec2<i32> where: T is f32, i32 or u32
textureDimensions(texture: texture_depth_2d) -> vec2<i32>
textureDimensions(texture: texture_depth_2d_array) -> vec2<i32>
textureDimensions(texture: texture_depth_cube) -> vec2<i32>
textureDimensions(texture: texture_depth_cube_array) -> vec2<i32>
textureDimensions(texture: texture_depth_multisampled_2d) -> vec2<i32>
textureDimensions(texture: texture_storage_1d<F, A>) -> i32 where: A is write
textureDimensions(texture: texture_storage_2d<F, A>) -> vec2<i32> where: A is write
textureDimensions(texture: texture_storage_2d_array<F, A>) -> vec2<i32> where: A is write
textureDimensions(texture: texture_storage_3d<F, A>) -> vec3<i32> where: A is write
textureDimensions(texture: texture_external) -> vec2<i32>
textureDimensions(texture: texture_1d<T>, level: C) -> u32 where: T is f32, i32 or u32, C is i32 or u32
textureDimensions(texture: texture_2d<T>, level: C) -> vec2<u32> where: T is f32, i32 or u32, C is i32 or u32
textureDimensions(texture: texture_2d_array<T>, level: C) -> vec2<u32> where: T is f32, i32 or u32, C is i32 or u32
textureDimensions(texture: texture_3d<T>, level: C) -> vec3<u32> where: T is f32, i32 or u32, C is i32 or u32
textureDimensions(texture: texture_cube<T>, level: C) -> vec2<u32> where: T is f32, i32 or u32, C is i32 or u32
textureDimensions(texture: texture_cube_array<T>, level: C) -> vec2<u32> where: T is f32, i32 or u32, C is i32 or u32
textureDimensions(texture: texture_depth_2d, level: C) -> vec2<u32> where: C is i32 or u32
textureDimensions(texture: texture_depth_2d_array, level: C) -> vec2<u32> where: C is i32 or u32
textureDimensions(texture: texture_depth_cube, level: C) -> vec2<u32> where: C is i32 or u32
textureDimensions(texture: texture_depth_cube_array, level: C) -> vec2<u32> where: C is i32 or u32
textureDimensions(texture: texture_1d<T>) -> u32 where: T is f32, i32 or u32
textureDimensions(texture: texture_2d<T>) -> vec2<u32> where: T is f32, i32 or u32
textureDimensions(texture: texture_2d_array<T>) -> vec2<u32> where: T is f32, i32 or u32
textureDimensions(texture: texture_3d<T>) -> vec3<u32> where: T is f32, i32 or u32
textureDimensions(texture: texture_cube<T>) -> vec2<u32> where: T is f32, i32 or u32
textureDimensions(texture: texture_cube_array<T>) -> vec2<u32> where: T is f32, i32 or u32
textureDimensions(texture: texture_multisampled_2d<T>) -> vec2<u32> where: T is f32, i32 or u32
textureDimensions(texture: texture_depth_2d) -> vec2<u32>
textureDimensions(texture: texture_depth_2d_array) -> vec2<u32>
textureDimensions(texture: texture_depth_cube) -> vec2<u32>
textureDimensions(texture: texture_depth_cube_array) -> vec2<u32>
textureDimensions(texture: texture_depth_multisampled_2d) -> vec2<u32>
textureDimensions(texture: texture_storage_1d<F, A>) -> u32 where: A is write
textureDimensions(texture: texture_storage_2d<F, A>) -> vec2<u32> where: A is write
textureDimensions(texture: texture_storage_2d_array<F, A>) -> vec2<u32> where: A is write
textureDimensions(texture: texture_storage_3d<F, A>) -> vec3<u32> where: A is write
textureDimensions(texture: texture_external) -> vec2<u32>
)");
}
@ -635,33 +635,33 @@ TEST_F(IntrinsicTableTest, OverloadOrderByMatchingParameter) {
R"(error: no matching call to textureDimensions(texture_depth_2d, bool)
27 candidate functions:
textureDimensions(texture: texture_depth_2d, level: C) -> vec2<i32> where: C is i32 or u32
textureDimensions(texture: texture_1d<T>, level: C) -> i32 where: T is f32, i32 or u32, C is i32 or u32
textureDimensions(texture: texture_2d<T>, level: C) -> vec2<i32> where: T is f32, i32 or u32, C is i32 or u32
textureDimensions(texture: texture_2d_array<T>, level: C) -> vec2<i32> where: T is f32, i32 or u32, C is i32 or u32
textureDimensions(texture: texture_3d<T>, level: C) -> vec3<i32> where: T is f32, i32 or u32, C is i32 or u32
textureDimensions(texture: texture_cube<T>, level: C) -> vec2<i32> where: T is f32, i32 or u32, C is i32 or u32
textureDimensions(texture: texture_cube_array<T>, level: C) -> vec2<i32> where: T is f32, i32 or u32, C is i32 or u32
textureDimensions(texture: texture_depth_2d_array, level: C) -> vec2<i32> where: C is i32 or u32
textureDimensions(texture: texture_depth_cube, level: C) -> vec2<i32> where: C is i32 or u32
textureDimensions(texture: texture_depth_cube_array, level: C) -> vec2<i32> where: C is i32 or u32
textureDimensions(texture: texture_depth_2d) -> vec2<i32>
textureDimensions(texture: texture_1d<T>) -> i32 where: T is f32, i32 or u32
textureDimensions(texture: texture_2d<T>) -> vec2<i32> where: T is f32, i32 or u32
textureDimensions(texture: texture_2d_array<T>) -> vec2<i32> where: T is f32, i32 or u32
textureDimensions(texture: texture_3d<T>) -> vec3<i32> where: T is f32, i32 or u32
textureDimensions(texture: texture_cube<T>) -> vec2<i32> where: T is f32, i32 or u32
textureDimensions(texture: texture_cube_array<T>) -> vec2<i32> where: T is f32, i32 or u32
textureDimensions(texture: texture_multisampled_2d<T>) -> vec2<i32> where: T is f32, i32 or u32
textureDimensions(texture: texture_depth_2d_array) -> vec2<i32>
textureDimensions(texture: texture_depth_cube) -> vec2<i32>
textureDimensions(texture: texture_depth_cube_array) -> vec2<i32>
textureDimensions(texture: texture_depth_multisampled_2d) -> vec2<i32>
textureDimensions(texture: texture_storage_1d<F, A>) -> i32 where: A is write
textureDimensions(texture: texture_storage_2d<F, A>) -> vec2<i32> where: A is write
textureDimensions(texture: texture_storage_2d_array<F, A>) -> vec2<i32> where: A is write
textureDimensions(texture: texture_storage_3d<F, A>) -> vec3<i32> where: A is write
textureDimensions(texture: texture_external) -> vec2<i32>
textureDimensions(texture: texture_depth_2d, level: C) -> vec2<u32> where: C is i32 or u32
textureDimensions(texture: texture_1d<T>, level: C) -> u32 where: T is f32, i32 or u32, C is i32 or u32
textureDimensions(texture: texture_2d<T>, level: C) -> vec2<u32> where: T is f32, i32 or u32, C is i32 or u32
textureDimensions(texture: texture_2d_array<T>, level: C) -> vec2<u32> where: T is f32, i32 or u32, C is i32 or u32
textureDimensions(texture: texture_3d<T>, level: C) -> vec3<u32> where: T is f32, i32 or u32, C is i32 or u32
textureDimensions(texture: texture_cube<T>, level: C) -> vec2<u32> where: T is f32, i32 or u32, C is i32 or u32
textureDimensions(texture: texture_cube_array<T>, level: C) -> vec2<u32> where: T is f32, i32 or u32, C is i32 or u32
textureDimensions(texture: texture_depth_2d_array, level: C) -> vec2<u32> where: C is i32 or u32
textureDimensions(texture: texture_depth_cube, level: C) -> vec2<u32> where: C is i32 or u32
textureDimensions(texture: texture_depth_cube_array, level: C) -> vec2<u32> where: C is i32 or u32
textureDimensions(texture: texture_depth_2d) -> vec2<u32>
textureDimensions(texture: texture_1d<T>) -> u32 where: T is f32, i32 or u32
textureDimensions(texture: texture_2d<T>) -> vec2<u32> where: T is f32, i32 or u32
textureDimensions(texture: texture_2d_array<T>) -> vec2<u32> where: T is f32, i32 or u32
textureDimensions(texture: texture_3d<T>) -> vec3<u32> where: T is f32, i32 or u32
textureDimensions(texture: texture_cube<T>) -> vec2<u32> where: T is f32, i32 or u32
textureDimensions(texture: texture_cube_array<T>) -> vec2<u32> where: T is f32, i32 or u32
textureDimensions(texture: texture_multisampled_2d<T>) -> vec2<u32> where: T is f32, i32 or u32
textureDimensions(texture: texture_depth_2d_array) -> vec2<u32>
textureDimensions(texture: texture_depth_cube) -> vec2<u32>
textureDimensions(texture: texture_depth_cube_array) -> vec2<u32>
textureDimensions(texture: texture_depth_multisampled_2d) -> vec2<u32>
textureDimensions(texture: texture_storage_1d<F, A>) -> u32 where: A is write
textureDimensions(texture: texture_storage_2d<F, A>) -> vec2<u32> where: A is write
textureDimensions(texture: texture_storage_2d_array<F, A>) -> vec2<u32> where: A is write
textureDimensions(texture: texture_storage_3d<F, A>) -> vec3<u32> where: A is write
textureDimensions(texture: texture_external) -> vec2<u32>
)");
}

View File

@ -138,7 +138,7 @@ TEST_F(MultiplanarExternalTextureTest, Dimensions) {
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
var dim : vec2<i32>;
var dim : vec2<u32>;
dim = textureDimensions(ext_tex);
return vec4<f32>(0.0, 0.0, 0.0, 0.0);
}
@ -173,7 +173,7 @@ struct ExternalTextureParams {
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
var dim : vec2<i32>;
var dim : vec2<u32>;
dim = textureDimensions(ext_tex);
return vec4<f32>(0.0, 0.0, 0.0, 0.0);
}
@ -191,7 +191,7 @@ TEST_F(MultiplanarExternalTextureTest, Dimensions_OutOfOrder) {
auto* src = R"(
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
var dim : vec2<i32>;
var dim : vec2<u32>;
dim = textureDimensions(ext_tex);
return vec4<f32>(0.0, 0.0, 0.0, 0.0);
}
@ -226,7 +226,7 @@ struct ExternalTextureParams {
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
var dim : vec2<i32>;
var dim : vec2<u32>;
dim = textureDimensions(ext_tex);
return vec4<f32>(0.0, 0.0, 0.0, 0.0);
}
@ -2006,7 +2006,7 @@ fn main() {
// even if there's no external texture declared at module scope.
TEST_F(MultiplanarExternalTextureTest, ExternalTexturePassedAsParamWithoutGlobalDecl) {
auto* src = R"(
fn f(ext_tex : texture_external) -> vec2<i32> {
fn f(ext_tex : texture_external) -> vec2<u32> {
return textureDimensions(ext_tex);
}
)";
@ -2032,7 +2032,7 @@ struct ExternalTextureParams {
gamutConversionMatrix : mat3x3<f32>,
}
fn f(ext_tex : texture_2d<f32>, ext_tex_plane_1 : texture_2d<f32>, ext_tex_params : ExternalTextureParams) -> vec2<i32> {
fn f(ext_tex : texture_2d<f32>, ext_tex_plane_1 : texture_2d<f32>, ext_tex_params : ExternalTextureParams) -> vec2<u32> {
return textureDimensions(ext_tex);
}
)";

View File

@ -1345,6 +1345,16 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& out,
}
};
auto emit_unsigned_int_type = [&](const sem::Type* ty) {
uint32_t width = 0;
sem::Type::ElementOf(ty, &width);
if (width > 1) {
out << "uvec" << width;
} else {
out << "uint";
}
};
auto emit_expr_as_signed = [&](const ast::Expression* e) {
auto* ty = TypeOf(e)->UnwrapRef();
if (!ty->is_unsigned_scalar_or_vector()) {
@ -1357,6 +1367,12 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& out,
switch (builtin->Type()) {
case sem::BuiltinType::kTextureDimensions: {
// textureDimensions() returns an unsigned scalar / vector in WGSL.
// textureSize() / imageSize() returns a signed scalar / vector in GLSL.
// Cast.
emit_unsigned_int_type(call->Type());
ScopedParen sp(out);
if (texture_type->Is<sem::StorageTexture>()) {
out << "imageSize(";
} else {
@ -1390,6 +1406,12 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& out,
return true;
}
case sem::BuiltinType::kTextureNumLayers: {
// textureNumLayers() returns an unsigned scalar in WGSL.
// textureSize() / imageSize() returns a signed scalar / vector in GLSL.
// Cast.
out << "uint";
ScopedParen sp(out);
if (texture_type->Is<sem::StorageTexture>()) {
out << "imageSize(";
} else {
@ -1418,6 +1440,12 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& out,
return true;
}
case sem::BuiltinType::kTextureNumLevels: {
// textureNumLevels() returns an unsigned scalar in WGSL.
// textureQueryLevels() returns a signed scalar in GLSL.
// Cast.
out << "uint";
ScopedParen sp(out);
out << "textureQueryLevels(";
if (!EmitExpression(out, texture)) {
return false;
@ -1426,6 +1454,12 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& out,
return true;
}
case sem::BuiltinType::kTextureNumSamples: {
// textureNumSamples() returns an unsigned scalar in WGSL.
// textureSamples() returns a signed scalar in GLSL.
// Cast.
out << "uint";
ScopedParen sp(out);
out << "textureSamples(";
if (!EmitExpression(out, texture)) {
return false;

View File

@ -219,15 +219,15 @@ ExpectedResult expected_texture_overload(ast::builtin::test::ValidTextureOverloa
case ValidTextureOverload::kSampleCompareDepthCubeArrayF32:
return R"(texture(tint_symbol_sampler, vec4(1.0f, 2.0f, 3.0f, float(4)), 5.0f);)";
case ValidTextureOverload::kSampleCompareLevelDepth2dF32:
return R"(yyytexture(tint_symbol_sampler, vec2(1.0f, 2.0f), 3.0f);)";
return R"(texture(tint_symbol_sampler, vec3(1.0f, 2.0f, 3.0f));)";
case ValidTextureOverload::kSampleCompareLevelDepth2dOffsetF32:
return R"(yyytextureOffset(tint_symbol_sampler, vec2(1.0f, 2.0f), 3.0f, ivec2(4, 5));)";
return R"(textureOffset(tint_symbol_sampler, vec3(1.0f, 2.0f, 3.0f), ivec2(4, 5));)";
case ValidTextureOverload::kSampleCompareLevelDepth2dArrayF32:
return R"(texture(tint_symbol_sampler, vec4(1.0f, 2.0f, float(4)), 3.0f);)";
return R"(texture(tint_symbol_sampler, vec4(1.0f, 2.0f, float(3), 4.0f));)";
case ValidTextureOverload::kSampleCompareLevelDepth2dArrayOffsetF32:
return R"(textureOffset(tint_symbol_sampler, vec3(1.0f, 2.0f, float(4)), 3.0f, ivec2(5, 6));)";
return R"(textureOffset(tint_symbol_sampler, vec4(1.0f, 2.0f, float(3), 4.0f), ivec2(5, 6));)";
case ValidTextureOverload::kSampleCompareLevelDepthCubeF32:
return R"(texture(tint_symbol_sampler, vec3(1.0f, 2.0f, 3.0f), 4.0f);)";
return R"(texture(tint_symbol_sampler, vec4(1.0f, 2.0f, 3.0f, 4.0f));)";
case ValidTextureOverload::kSampleCompareLevelDepthCubeArrayF32:
return R"(texture(tint_symbol_sampler, vec4(1.0f, 2.0f, 3.0f, float(4)), 5.0f);)";
case ValidTextureOverload::kLoad1dLevelF32:
@ -248,7 +248,6 @@ ExpectedResult expected_texture_overload(ast::builtin::test::ValidTextureOverloa
case ValidTextureOverload::kLoad2dArrayLevelI32:
case ValidTextureOverload::kLoad3dLevelI32:
return R"(texelFetch(tint_symbol_2, ivec3(uvec3(1u, 2u, 3u)), int(4u));)";
case ValidTextureOverload::kLoadDepthMultisampled2dF32:
case ValidTextureOverload::kLoadMultisampled2dF32:
case ValidTextureOverload::kLoadMultisampled2dU32:
return R"(texelFetch(tint_symbol_2, ivec2(1, 2), 3);)";
@ -258,6 +257,8 @@ ExpectedResult expected_texture_overload(ast::builtin::test::ValidTextureOverloa
return R"(texelFetch(tint_symbol_2, ivec2(1, 2), 3);)";
case ValidTextureOverload::kLoadDepth2dArrayLevelF32:
return R"(texelFetch(tint_symbol_2, ivec3(uvec3(1u, 2u, 3u)), int(4u));)";
case ValidTextureOverload::kLoadDepthMultisampled2dF32:
return R"(texelFetch(tint_symbol_2, ivec2(uvec2(1u, 2u)), int(3u)).x;)";
case ValidTextureOverload::kStoreWO1dRgba32float:
return R"(imageStore(tint_symbol, 1, vec4(2.0f, 3.0f, 4.0f, 5.0f));)";
case ValidTextureOverload::kStoreWO2dRgba32float:

View File

@ -315,9 +315,9 @@ ExpectedResult expected_texture_overload(ast::builtin::test::ValidTextureOverloa
case ValidTextureOverload::kSampleCompareLevelDepth2dOffsetF32:
return R"(tint_symbol.SampleCmpLevelZero(tint_symbol_1, float2(1.0f, 2.0f), 3.0f, int2(4, 5));)";
case ValidTextureOverload::kSampleCompareLevelDepth2dArrayF32:
return R"(tint_symbol.SampleCmpLevelZero(tint_symbol_1, float3(1.0f, 2.0f, float(4)), 3.0f);)";
return R"(tint_symbol.SampleCmpLevelZero(tint_symbol_1, float3(1.0f, 2.0f, float(3)), 4.0f);)";
case ValidTextureOverload::kSampleCompareLevelDepth2dArrayOffsetF32:
return R"(tint_symbol.SampleCmpLevelZero(tint_symbol_1, float3(1.0f, 2.0f, float(4)), 3.0f, int2(5, 6));)";
return R"(tint_symbol.SampleCmpLevelZero(tint_symbol_1, float3(1.0f, 2.0f, float(3)), 4.0f, int2(5, 6));)";
case ValidTextureOverload::kSampleCompareLevelDepthCubeF32:
return R"(tint_symbol.SampleCmpLevelZero(tint_symbol_1, float3(1.0f, 2.0f, 3.0f), 4.0f);)";
case ValidTextureOverload::kSampleCompareLevelDepthCubeArrayF32:
@ -340,7 +340,6 @@ ExpectedResult expected_texture_overload(ast::builtin::test::ValidTextureOverloa
case ValidTextureOverload::kLoad2dArrayLevelI32:
case ValidTextureOverload::kLoad3dLevelI32:
return R"(tint_symbol.Load(uint4(1u, 2u, 3u, 4u));)";
case ValidTextureOverload::kLoadDepthMultisampled2dF32:
case ValidTextureOverload::kLoadMultisampled2dF32:
case ValidTextureOverload::kLoadMultisampled2dU32:
return R"(tint_symbol.Load(int2(1, 2), 3);)";
@ -350,6 +349,8 @@ ExpectedResult expected_texture_overload(ast::builtin::test::ValidTextureOverloa
return R"(tint_symbol.Load(int3(1, 2, 3)).x;)";
case ValidTextureOverload::kLoadDepth2dArrayLevelF32:
return R"(tint_symbol.Load(uint4(1u, 2u, 3u, 4u)).x;)";
case ValidTextureOverload::kLoadDepthMultisampled2dF32:
return R"(tint_symbol.Load(uint3(1u, 2u, uint(0)), 3u).x;)";
case ValidTextureOverload::kStoreWO1dRgba32float:
return R"(tint_symbol[1] = float4(2.0f, 3.0f, 4.0f, 5.0f);)";
case ValidTextureOverload::kStoreWO2dRgba32float:

View File

@ -1092,9 +1092,7 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& out,
};
if (dims.size() == 1) {
out << "int(";
get_dim(dims[0]);
out << ")";
} else {
EmitType(out, TypeOf(expr)->UnwrapRef(), "");
out << "(";
@ -1109,27 +1107,24 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& out,
return true;
}
case sem::BuiltinType::kTextureNumLayers: {
out << "int(";
if (!texture_expr()) {
return false;
}
out << ".get_array_size())";
out << ".get_array_size()";
return true;
}
case sem::BuiltinType::kTextureNumLevels: {
out << "int(";
if (!texture_expr()) {
return false;
}
out << ".get_num_mip_levels())";
out << ".get_num_mip_levels()";
return true;
}
case sem::BuiltinType::kTextureNumSamples: {
out << "int(";
if (!texture_expr()) {
return false;
}
out << ".get_num_samples())";
out << ".get_num_samples()";
return true;
}
default:

View File

@ -24,7 +24,7 @@ std::string expected_texture_overload(ast::builtin::test::ValidTextureOverload o
switch (overload) {
case ValidTextureOverload::kDimensions1d:
case ValidTextureOverload::kDimensionsStorageWO1d:
return R"(int(texture.get_width(0)))";
return R"(texture.get_width(0))";
case ValidTextureOverload::kDimensions2d:
case ValidTextureOverload::kDimensions2dArray:
case ValidTextureOverload::kDimensionsCube:
@ -37,10 +37,10 @@ std::string expected_texture_overload(ast::builtin::test::ValidTextureOverload o
case ValidTextureOverload::kDimensionsDepthMultisampled2d:
case ValidTextureOverload::kDimensionsStorageWO2d:
case ValidTextureOverload::kDimensionsStorageWO2dArray:
return R"(int2(texture.get_width(), texture.get_height()))";
return R"(uint2(texture.get_width(), texture.get_height()))";
case ValidTextureOverload::kDimensions3d:
case ValidTextureOverload::kDimensionsStorageWO3d:
return R"(int3(texture.get_width(), texture.get_height(), texture.get_depth()))";
return R"(uint3(texture.get_width(), texture.get_height(), texture.get_depth()))";
case ValidTextureOverload::kDimensions2dLevel:
case ValidTextureOverload::kDimensionsCubeLevel:
case ValidTextureOverload::kDimensionsCubeArrayLevel:
@ -49,9 +49,9 @@ std::string expected_texture_overload(ast::builtin::test::ValidTextureOverload o
case ValidTextureOverload::kDimensionsDepth2dArrayLevel:
case ValidTextureOverload::kDimensionsDepthCubeLevel:
case ValidTextureOverload::kDimensionsDepthCubeArrayLevel:
return R"(int2(texture.get_width(1), texture.get_height(1)))";
return R"(uint2(texture.get_width(1), texture.get_height(1)))";
case ValidTextureOverload::kDimensions3dLevel:
return R"(int3(texture.get_width(1), texture.get_height(1), texture.get_depth(1)))";
return R"(uint3(texture.get_width(1), texture.get_height(1), texture.get_depth(1)))";
case ValidTextureOverload::kGather2dF32:
return R"(texture.gather(sampler, float2(1.0f, 2.0f), int2(0), component::x))";
case ValidTextureOverload::kGather2dOffsetF32:
@ -93,7 +93,7 @@ std::string expected_texture_overload(ast::builtin::test::ValidTextureOverload o
case ValidTextureOverload::kNumLayersDepth2dArray:
case ValidTextureOverload::kNumLayersDepthCubeArray:
case ValidTextureOverload::kNumLayersStorageWO2dArray:
return R"(int(texture.get_array_size()))";
return R"(texture.get_array_size())";
case ValidTextureOverload::kNumLevels2d:
case ValidTextureOverload::kNumLevels2dArray:
case ValidTextureOverload::kNumLevels3d:
@ -103,10 +103,10 @@ std::string expected_texture_overload(ast::builtin::test::ValidTextureOverload o
case ValidTextureOverload::kNumLevelsDepth2dArray:
case ValidTextureOverload::kNumLevelsDepthCube:
case ValidTextureOverload::kNumLevelsDepthCubeArray:
return R"(int(texture.get_num_mip_levels()))";
return R"(texture.get_num_mip_levels())";
case ValidTextureOverload::kNumSamplesDepthMultisampled2d:
case ValidTextureOverload::kNumSamplesMultisampled2d:
return R"(int(texture.get_num_samples()))";
return R"(texture.get_num_samples())";
case ValidTextureOverload::kSample1dF32:
return R"(texture.sample(sampler, 1.0f))";
case ValidTextureOverload::kSample2dF32:
@ -210,17 +210,17 @@ std::string expected_texture_overload(ast::builtin::test::ValidTextureOverload o
case ValidTextureOverload::kSampleCompareDepthCubeArrayF32:
return R"(texture.sample_compare(sampler, float3(1.0f, 2.0f, 3.0f), 4, 5.0f))";
case ValidTextureOverload::kSampleCompareLevelDepth2dF32:
return R"(texture.sample_compare(sampler, float2(1.0f, 2.0f), 3.0f))";
return R"(texture.sample_compare(sampler, float2(1.0f, 2.0f), 3.0f, level(0)))";
case ValidTextureOverload::kSampleCompareLevelDepth2dOffsetF32:
return R"(texture.sample_compare(sampler, float2(1.0f, 2.0f), 3.0f, int2(4, 5)))";
return R"(texture.sample_compare(sampler, float2(1.0f, 2.0f), 3.0f, level(0), int2(4, 5)))";
case ValidTextureOverload::kSampleCompareLevelDepth2dArrayF32:
return R"(texture.sample_compare(sampler, float2(1.0f, 2.0f), 4, 3.0f))";
return R"(texture.sample_compare(sampler, float2(1.0f, 2.0f), 3, 4.0f, level(0)))";
case ValidTextureOverload::kSampleCompareLevelDepth2dArrayOffsetF32:
return R"(texture.sample_compare(sampler, float2(1.0f, 2.0f), 4, 3.0f, int2(5, 6)))";
return R"(texture.sample_compare(sampler, float2(1.0f, 2.0f), 3, 4.0f, level(0), int2(5, 6)))";
case ValidTextureOverload::kSampleCompareLevelDepthCubeF32:
return R"(texture.sample_compare(sampler, float3(1.0f, 2.0f, 3.0f), 4.0f))";
return R"(texture.sample_compare(sampler, float3(1.0f, 2.0f, 3.0f), 4.0f, level(0)))";
case ValidTextureOverload::kSampleCompareLevelDepthCubeArrayF32:
return R"(texture.sample_compare(sampler, float3(1.0f, 2.0f, 3.0f), 4, 5.0f))";
return R"(texture.sample_compare(sampler, float3(1.0f, 2.0f, 3.0f), 4, 5.0f, level(0)))";
case ValidTextureOverload::kLoad1dLevelF32:
return R"(texture.read(uint(1u), 0))";
case ValidTextureOverload::kLoad1dLevelU32:
@ -250,10 +250,11 @@ std::string expected_texture_overload(ast::builtin::test::ValidTextureOverload o
case ValidTextureOverload::kLoadMultisampled2dI32:
return R"(texture.read(uint2(uint2(1u, 2u)), 3u))";
case ValidTextureOverload::kLoadDepth2dLevelF32:
case ValidTextureOverload::kLoadDepthMultisampled2dF32:
return R"(texture.read(uint2(int2(1, 2)), 3))";
case ValidTextureOverload::kLoadDepth2dArrayLevelF32:
return R"(texture.read(uint2(uint2(1u, 2u)), 3u, 4u))";
case ValidTextureOverload::kLoadDepthMultisampled2dF32:
return R"(texture.read(uint2(uint2(1u, 2u)), 3u))";
case ValidTextureOverload::kStoreWO1dRgba32float:
return R"(texture.write(float4(2.0f, 3.0f, 4.0f, 5.0f), uint(1)))";
case ValidTextureOverload::kStoreWO2dRgba32float:

View File

@ -42,12 +42,13 @@ expected_texture_overload_spirv expected_texture_overload(
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeInt 32 1
%11 = OpConstant %9 0
%9 = OpTypeInt 32 0
%11 = OpTypeInt 32 1
%12 = OpConstant %11 0
)",
R"(
%10 = OpLoad %3 %1
%8 = OpImageQuerySizeLod %9 %10 %11
%8 = OpImageQuerySizeLod %9 %10 %12
)",
R"(
OpCapability Sampled1D
@ -63,13 +64,14 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%10 = OpTypeInt 32 1
%10 = OpTypeInt 32 0
%9 = OpTypeVector %10 2
%12 = OpConstant %10 0
%12 = OpTypeInt 32 1
%13 = OpConstant %12 0
)",
R"(
%11 = OpLoad %3 %1
%8 = OpImageQuerySizeLod %9 %11 %12
%8 = OpImageQuerySizeLod %9 %11 %13
)",
R"(
OpCapability ImageQuery
@ -84,13 +86,14 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%10 = OpTypeInt 32 1
%10 = OpTypeInt 32 0
%9 = OpTypeVector %10 2
%12 = OpConstant %10 1
%12 = OpTypeInt 32 1
%13 = OpConstant %12 1
)",
R"(
%11 = OpLoad %3 %1
%8 = OpImageQuerySizeLod %9 %11 %12
%8 = OpImageQuerySizeLod %9 %11 %13
)",
R"(
OpCapability ImageQuery
@ -105,14 +108,15 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%10 = OpTypeInt 32 1
%10 = OpTypeInt 32 0
%9 = OpTypeVector %10 2
%12 = OpTypeVector %10 3
%14 = OpConstant %10 0
%14 = OpTypeInt 32 1
%15 = OpConstant %14 0
)",
R"(
%13 = OpLoad %3 %1
%11 = OpImageQuerySizeLod %12 %13 %14
%11 = OpImageQuerySizeLod %12 %13 %15
%8 = OpVectorShuffle %9 %11 %11 0 1
)",
R"(
@ -128,14 +132,15 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%10 = OpTypeInt 32 1
%10 = OpTypeInt 32 0
%9 = OpTypeVector %10 2
%12 = OpTypeVector %10 3
%14 = OpConstant %10 1
%14 = OpTypeInt 32 1
%15 = OpConstant %14 1
)",
R"(
%13 = OpLoad %3 %1
%11 = OpImageQuerySizeLod %12 %13 %14
%11 = OpImageQuerySizeLod %12 %13 %15
%8 = OpVectorShuffle %9 %11 %11 0 1
)",
R"(
@ -151,13 +156,14 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%10 = OpTypeInt 32 1
%10 = OpTypeInt 32 0
%9 = OpTypeVector %10 3
%12 = OpConstant %10 0
%12 = OpTypeInt 32 1
%13 = OpConstant %12 0
)",
R"(
%11 = OpLoad %3 %1
%8 = OpImageQuerySizeLod %9 %11 %12
%8 = OpImageQuerySizeLod %9 %11 %13
)",
R"(
OpCapability ImageQuery
@ -172,13 +178,14 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%10 = OpTypeInt 32 1
%10 = OpTypeInt 32 0
%9 = OpTypeVector %10 3
%12 = OpConstant %10 1
%12 = OpTypeInt 32 1
%13 = OpConstant %12 1
)",
R"(
%11 = OpLoad %3 %1
%8 = OpImageQuerySizeLod %9 %11 %12
%8 = OpImageQuerySizeLod %9 %11 %13
)",
R"(
OpCapability ImageQuery
@ -193,13 +200,14 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%10 = OpTypeInt 32 1
%10 = OpTypeInt 32 0
%9 = OpTypeVector %10 2
%12 = OpConstant %10 0
%12 = OpTypeInt 32 1
%13 = OpConstant %12 0
)",
R"(
%11 = OpLoad %3 %1
%8 = OpImageQuerySizeLod %9 %11 %12
%8 = OpImageQuerySizeLod %9 %11 %13
)",
R"(
OpCapability ImageQuery
@ -214,13 +222,14 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%10 = OpTypeInt 32 1
%10 = OpTypeInt 32 0
%9 = OpTypeVector %10 2
%12 = OpConstant %10 1
%12 = OpTypeInt 32 1
%13 = OpConstant %12 1
)",
R"(
%11 = OpLoad %3 %1
%8 = OpImageQuerySizeLod %9 %11 %12
%8 = OpImageQuerySizeLod %9 %11 %13
)",
R"(
OpCapability ImageQuery
@ -235,14 +244,15 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%10 = OpTypeInt 32 1
%10 = OpTypeInt 32 0
%9 = OpTypeVector %10 2
%12 = OpTypeVector %10 3
%14 = OpConstant %10 0
%14 = OpTypeInt 32 1
%15 = OpConstant %14 0
)",
R"(
%13 = OpLoad %3 %1
%11 = OpImageQuerySizeLod %12 %13 %14
%11 = OpImageQuerySizeLod %12 %13 %15
%8 = OpVectorShuffle %9 %11 %11 0 1
)",
R"(
@ -259,14 +269,15 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%10 = OpTypeInt 32 1
%10 = OpTypeInt 32 0
%9 = OpTypeVector %10 2
%12 = OpTypeVector %10 3
%14 = OpConstant %10 1
%14 = OpTypeInt 32 1
%15 = OpConstant %14 1
)",
R"(
%13 = OpLoad %3 %1
%11 = OpImageQuerySizeLod %12 %13 %14
%11 = OpImageQuerySizeLod %12 %13 %15
%8 = OpVectorShuffle %9 %11 %11 0 1
)",
R"(
@ -283,7 +294,7 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%10 = OpTypeInt 32 1
%10 = OpTypeInt 32 0
%9 = OpTypeVector %10 2
)",
R"(
@ -303,13 +314,14 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%10 = OpTypeInt 32 1
%10 = OpTypeInt 32 0
%9 = OpTypeVector %10 2
%12 = OpConstant %10 0
%12 = OpTypeInt 32 1
%13 = OpConstant %12 0
)",
R"(
%11 = OpLoad %3 %1
%8 = OpImageQuerySizeLod %9 %11 %12
%8 = OpImageQuerySizeLod %9 %11 %13
)",
R"(
OpCapability ImageQuery
@ -324,13 +336,14 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%10 = OpTypeInt 32 1
%10 = OpTypeInt 32 0
%9 = OpTypeVector %10 2
%12 = OpConstant %10 1
%12 = OpTypeInt 32 1
%13 = OpConstant %12 1
)",
R"(
%11 = OpLoad %3 %1
%8 = OpImageQuerySizeLod %9 %11 %12
%8 = OpImageQuerySizeLod %9 %11 %13
)",
R"(
OpCapability ImageQuery
@ -345,14 +358,15 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%10 = OpTypeInt 32 1
%10 = OpTypeInt 32 0
%9 = OpTypeVector %10 2
%12 = OpTypeVector %10 3
%14 = OpConstant %10 0
%14 = OpTypeInt 32 1
%15 = OpConstant %14 0
)",
R"(
%13 = OpLoad %3 %1
%11 = OpImageQuerySizeLod %12 %13 %14
%11 = OpImageQuerySizeLod %12 %13 %15
%8 = OpVectorShuffle %9 %11 %11 0 1
)",
R"(
@ -368,14 +382,15 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%10 = OpTypeInt 32 1
%10 = OpTypeInt 32 0
%9 = OpTypeVector %10 2
%12 = OpTypeVector %10 3
%14 = OpConstant %10 1
%14 = OpTypeInt 32 1
%15 = OpConstant %14 1
)",
R"(
%13 = OpLoad %3 %1
%11 = OpImageQuerySizeLod %12 %13 %14
%11 = OpImageQuerySizeLod %12 %13 %15
%8 = OpVectorShuffle %9 %11 %11 0 1
)",
R"(
@ -391,13 +406,14 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%10 = OpTypeInt 32 1
%10 = OpTypeInt 32 0
%9 = OpTypeVector %10 2
%12 = OpConstant %10 0
%12 = OpTypeInt 32 1
%13 = OpConstant %12 0
)",
R"(
%11 = OpLoad %3 %1
%8 = OpImageQuerySizeLod %9 %11 %12
%8 = OpImageQuerySizeLod %9 %11 %13
)",
R"(
OpCapability ImageQuery
@ -412,13 +428,14 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%10 = OpTypeInt 32 1
%10 = OpTypeInt 32 0
%9 = OpTypeVector %10 2
%12 = OpConstant %10 1
%12 = OpTypeInt 32 1
%13 = OpConstant %12 1
)",
R"(
%11 = OpLoad %3 %1
%8 = OpImageQuerySizeLod %9 %11 %12
%8 = OpImageQuerySizeLod %9 %11 %13
)",
R"(
OpCapability ImageQuery
@ -433,14 +450,15 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%10 = OpTypeInt 32 1
%10 = OpTypeInt 32 0
%9 = OpTypeVector %10 2
%12 = OpTypeVector %10 3
%14 = OpConstant %10 0
%14 = OpTypeInt 32 1
%15 = OpConstant %14 0
)",
R"(
%13 = OpLoad %3 %1
%11 = OpImageQuerySizeLod %12 %13 %14
%11 = OpImageQuerySizeLod %12 %13 %15
%8 = OpVectorShuffle %9 %11 %11 0 1
)",
R"(
@ -457,14 +475,15 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%10 = OpTypeInt 32 1
%10 = OpTypeInt 32 0
%9 = OpTypeVector %10 2
%12 = OpTypeVector %10 3
%14 = OpConstant %10 1
%14 = OpTypeInt 32 1
%15 = OpConstant %14 1
)",
R"(
%13 = OpLoad %3 %1
%11 = OpImageQuerySizeLod %12 %13 %14
%11 = OpImageQuerySizeLod %12 %13 %15
%8 = OpVectorShuffle %9 %11 %11 0 1
)",
R"(
@ -481,7 +500,7 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%10 = OpTypeInt 32 1
%10 = OpTypeInt 32 0
%9 = OpTypeVector %10 2
)",
R"(
@ -501,7 +520,7 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeInt 32 1
%9 = OpTypeInt 32 0
)",
R"(
%10 = OpLoad %3 %1
@ -521,7 +540,7 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%10 = OpTypeInt 32 1
%10 = OpTypeInt 32 0
%9 = OpTypeVector %10 2
)",
R"(
@ -541,7 +560,7 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%10 = OpTypeInt 32 1
%10 = OpTypeInt 32 0
%9 = OpTypeVector %10 2
%12 = OpTypeVector %10 3
)",
@ -563,7 +582,7 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%10 = OpTypeInt 32 1
%10 = OpTypeInt 32 0
%9 = OpTypeVector %10 3
)",
R"(
@ -1100,13 +1119,14 @@ OpCapability SampledCubeArray
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeInt 32 1
%9 = OpTypeInt 32 0
%11 = OpTypeVector %9 3
%13 = OpConstant %9 0
%13 = OpTypeInt 32 1
%14 = OpConstant %13 0
)",
R"(
%12 = OpLoad %3 %1
%10 = OpImageQuerySizeLod %11 %12 %13
%10 = OpImageQuerySizeLod %11 %12 %14
%8 = OpCompositeExtract %9 %10 2
)",
R"(
@ -1121,13 +1141,14 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeInt 32 1
%9 = OpTypeInt 32 0
%11 = OpTypeVector %9 3
%13 = OpConstant %9 0
%13 = OpTypeInt 32 1
%14 = OpConstant %13 0
)",
R"(
%12 = OpLoad %3 %1
%10 = OpImageQuerySizeLod %11 %12 %13
%10 = OpImageQuerySizeLod %11 %12 %14
%8 = OpCompositeExtract %9 %10 2
)",
R"(
@ -1143,13 +1164,14 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeInt 32 1
%9 = OpTypeInt 32 0
%11 = OpTypeVector %9 3
%13 = OpConstant %9 0
%13 = OpTypeInt 32 1
%14 = OpConstant %13 0
)",
R"(
%12 = OpLoad %3 %1
%10 = OpImageQuerySizeLod %11 %12 %13
%10 = OpImageQuerySizeLod %11 %12 %14
%8 = OpCompositeExtract %9 %10 2
)",
R"(
@ -1164,13 +1186,14 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeInt 32 1
%9 = OpTypeInt 32 0
%11 = OpTypeVector %9 3
%13 = OpConstant %9 0
%13 = OpTypeInt 32 1
%14 = OpConstant %13 0
)",
R"(
%12 = OpLoad %3 %1
%10 = OpImageQuerySizeLod %11 %12 %13
%10 = OpImageQuerySizeLod %11 %12 %14
%8 = OpCompositeExtract %9 %10 2
)",
R"(
@ -1186,7 +1209,7 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeInt 32 1
%9 = OpTypeInt 32 0
%11 = OpTypeVector %9 3
)",
R"(
@ -1206,7 +1229,7 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeInt 32 1
%9 = OpTypeInt 32 0
)",
R"(
%10 = OpLoad %3 %1
@ -1224,7 +1247,7 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeInt 32 1
%9 = OpTypeInt 32 0
)",
R"(
%10 = OpLoad %3 %1
@ -1242,7 +1265,7 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeInt 32 1
%9 = OpTypeInt 32 0
)",
R"(
%10 = OpLoad %3 %1
@ -1260,7 +1283,7 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeInt 32 1
%9 = OpTypeInt 32 0
)",
R"(
%10 = OpLoad %3 %1
@ -1278,7 +1301,7 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeInt 32 1
%9 = OpTypeInt 32 0
)",
R"(
%10 = OpLoad %3 %1
@ -1297,7 +1320,7 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeInt 32 1
%9 = OpTypeInt 32 0
)",
R"(
%10 = OpLoad %3 %1
@ -1315,7 +1338,7 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeInt 32 1
%9 = OpTypeInt 32 0
)",
R"(
%10 = OpLoad %3 %1
@ -1333,7 +1356,7 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeInt 32 1
%9 = OpTypeInt 32 0
)",
R"(
%10 = OpLoad %3 %1
@ -1351,7 +1374,7 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeInt 32 1
%9 = OpTypeInt 32 0
)",
R"(
%10 = OpLoad %3 %1
@ -1370,7 +1393,7 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeInt 32 1
%9 = OpTypeInt 32 0
)",
R"(
%10 = OpLoad %3 %1
@ -1388,7 +1411,7 @@ OpCapability ImageQuery
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeInt 32 1
%9 = OpTypeInt 32 0
)",
R"(
%10 = OpLoad %3 %1
@ -3019,8 +3042,8 @@ OpCapability SampledCubeArray
%14 = OpConstant %4 1
%15 = OpConstant %4 2
%17 = OpTypeInt 32 1
%18 = OpConstant %17 4
%20 = OpConstant %4 3
%18 = OpConstant %17 3
%20 = OpConstant %4 4
%21 = OpConstant %4 0
)",
R"(
@ -3048,8 +3071,8 @@ OpCapability SampledCubeArray
%14 = OpConstant %4 1
%15 = OpConstant %4 2
%17 = OpTypeInt 32 1
%18 = OpConstant %17 4
%20 = OpConstant %4 3
%18 = OpConstant %17 3
%20 = OpConstant %4 4
%21 = OpConstant %4 0
%22 = OpTypeVector %17 2
%23 = OpConstant %17 5
@ -3538,24 +3561,23 @@ OpCapability Sampled1D
return {
R"(
%4 = OpTypeFloat 32
%3 = OpTypeImage %4 2D 0 1 0 1 Unknown
%3 = OpTypeImage %4 2D 0 0 1 1 Unknown
%2 = OpTypePointer UniformConstant %3
%1 = OpVariable %2 UniformConstant
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%10 = OpTypeVector %4 4
%13 = OpTypeInt 32 1
%12 = OpTypeVector %13 3
%13 = OpTypeInt 32 0
%12 = OpTypeVector %13 2
%14 = OpConstant %13 1
%15 = OpConstant %13 2
%16 = OpConstant %13 3
%17 = OpConstantComposite %12 %14 %15 %16
%18 = OpConstant %13 4
%16 = OpConstantComposite %12 %14 %15
%17 = OpConstant %13 3
)",
R"(
%11 = OpLoad %3 %1
%9 = OpImageFetch %10 %11 %17 Sample %18
%9 = OpImageFetch %10 %11 %16 Sample %17
%8 = OpCompositeExtract %4 %9 0
)",
R"(

View File

@ -98,7 +98,7 @@ fn simulate(@builtin(global_invocation_id) GlobalInvocationID : vec3<u32>) {
if (particle.lifetime < 0.0) {
// Use the probability map to find where the particle should be spawned.
// Starting with the 1x1 mip level.
var coord = vec2<i32>(0, 0);
var coord = vec2<u32>(0, 0);
for (var level = textureNumLevels(texture) - 1; level > 0; level = level - 1) {
// Load the probability value from the mip-level
// Generate a random number and using the probabilty values, pick the
@ -112,8 +112,8 @@ fn simulate(@builtin(global_invocation_id) GlobalInvocationID : vec3<u32>) {
let value = vec4<f32>(rand());
let mask = (value >= vec4<f32>(0.0, probabilites.xyz)) & (value < probabilites);
coord = coord * 2;
coord.x = coord.x + select(0, 1, any(mask.yw)); // x y
coord.y = coord.y + select(0, 1, any(mask.zw)); // z w
coord.x += select(0u, 1u, any(mask.yw)); // x y
coord.y += select(0u, 1u, any(mask.zw)); // z w
}
let uv = vec2<f32>(coord) / vec2<f32>(textureDimensions(texture));
particle.position = vec3<f32>((uv - 0.5) * 3.0 * vec2<f32>(1.0, -1.0), 0.0);

View File

@ -177,21 +177,21 @@ void simulate(uvec3 GlobalInvocationID) {
particle.lifetime = (particle.lifetime - sim_params.deltaTime);
particle.color.a = smoothstep(0.0f, 0.5f, particle.lifetime);
if ((particle.lifetime < 0.0f)) {
ivec2 coord = ivec2(0);
uvec2 coord = uvec2(0u);
{
for(int level = (textureQueryLevels(tint_symbol_6) - 1); (level > 0); level = (level - 1)) {
vec4 probabilites = texelFetch(tint_symbol_6, coord, level);
for(uint level = (uint(textureQueryLevels(tint_symbol_6)) - 1u); (level > 0u); level = (level - 1u)) {
vec4 probabilites = texelFetch(tint_symbol_6, ivec2(coord), int(level));
float tint_symbol_5 = rand();
vec4 value = vec4(tint_symbol_5);
bvec4 mask = bvec4(uvec4(greaterThanEqual(value, vec4(0.0f, probabilites.xyz))) & uvec4(lessThan(value, probabilites)));
coord = (coord * 2);
coord.x = (coord.x + (any(mask.yw) ? 1 : 0));
coord.y = (coord.y + (any(mask.zw) ? 1 : 0));
coord = (coord * 2u);
coord.x = (coord.x + (any(mask.yw) ? 1u : 0u));
coord.y = (coord.y + (any(mask.zw) ? 1u : 0u));
}
}
vec2 uv = (vec2(coord) / vec2(textureSize(tint_symbol_6, 0)));
vec2 uv = (vec2(coord) / vec2(uvec2(textureSize(tint_symbol_6, 0))));
particle.position = vec3((((uv - 0.5f) * 3.0f) * vec2(1.0f, -1.0f)), 0.0f);
particle.color = texelFetch(tint_symbol_6, coord, 0);
particle.color = texelFetch(tint_symbol_6, ivec2(coord), int(0u));
float tint_symbol_1 = rand();
particle.velocity.x = ((tint_symbol_1 - 0.5f) * 0.100000001f);
float tint_symbol_2 = rand();
@ -323,7 +323,7 @@ layout(binding = 5, std430) buffer Buffer_ssbo_1 {
layout(rgba8) uniform highp writeonly image2D tex_out;
void export_level(uvec3 coord) {
if (all(lessThan(coord.xy, uvec2(imageSize(tex_out))))) {
if (all(lessThan(coord.xy, uvec2(uvec2(imageSize(tex_out)))))) {
uint dst_offset = (coord.x + (coord.y * ubo.width));
uint src_offset = ((coord.x * 2u) + ((coord.y * 2u) * ubo.width));
float a = buf_in.weights[(src_offset + 0u)];

View File

@ -20,9 +20,9 @@ fn ConvertToFp16FloatValue(fp32 : f32) -> u32 {
@compute @workgroup_size(1, 1, 1)
fn main(@builtin(global_invocation_id) GlobalInvocationID : vec3<u32>) {
var size : vec2<i32> = textureDimensions(src);
var dstTexCoord : vec2<i32> = vec2<i32>(GlobalInvocationID.xy);
var srcTexCoord : vec2<i32> = dstTexCoord;
var size = textureDimensions(src);
var dstTexCoord = GlobalInvocationID.xy;
var srcTexCoord = dstTexCoord;
if (uniforms.dstTextureFlipY == 1u) {
srcTexCoord.y = size.y - dstTexCoord.y - 1;
}

View File

@ -20,14 +20,14 @@ struct tint_symbol_2 {
void main_inner(uint3 GlobalInvocationID) {
int2 tint_tmp;
src.GetDimensions(tint_tmp.x, tint_tmp.y);
int2 size = tint_tmp;
int2 dstTexCoord = int2(GlobalInvocationID.xy);
int2 srcTexCoord = dstTexCoord;
uint2 size = tint_tmp;
uint2 dstTexCoord = GlobalInvocationID.xy;
uint2 srcTexCoord = dstTexCoord;
if ((uniforms[0].x == 1u)) {
srcTexCoord.y = ((size.y - dstTexCoord.y) - 1);
srcTexCoord.y = ((size.y - dstTexCoord.y) - 1u);
}
float4 srcColor = src.Load(int3(srcTexCoord, 0));
float4 dstColor = tint_symbol.Load(int3(dstTexCoord, 0));
float4 srcColor = src.Load(uint3(srcTexCoord, 0u));
float4 dstColor = tint_symbol.Load(uint3(dstTexCoord, 0u));
bool success = true;
uint4 srcColorBits = uint4(0u, 0u, 0u, 0u);
uint4 dstColorBits = uint4(dstColor);

View File

@ -20,14 +20,14 @@ struct tint_symbol_2 {
void main_inner(uint3 GlobalInvocationID) {
int2 tint_tmp;
src.GetDimensions(tint_tmp.x, tint_tmp.y);
int2 size = tint_tmp;
int2 dstTexCoord = int2(GlobalInvocationID.xy);
int2 srcTexCoord = dstTexCoord;
uint2 size = tint_tmp;
uint2 dstTexCoord = GlobalInvocationID.xy;
uint2 srcTexCoord = dstTexCoord;
if ((uniforms[0].x == 1u)) {
srcTexCoord.y = ((size.y - dstTexCoord.y) - 1);
srcTexCoord.y = ((size.y - dstTexCoord.y) - 1u);
}
float4 srcColor = src.Load(int3(srcTexCoord, 0));
float4 dstColor = tint_symbol.Load(int3(dstTexCoord, 0));
float4 srcColor = src.Load(uint3(srcTexCoord, 0u));
float4 dstColor = tint_symbol.Load(uint3(dstTexCoord, 0u));
bool success = true;
uint4 srcColorBits = uint4(0u, 0u, 0u, 0u);
uint4 dstColorBits = uint4(dstColor);

View File

@ -18,14 +18,14 @@ uint ConvertToFp16FloatValue(float fp32) {
uniform highp sampler2D src_1;
uniform highp sampler2D dst_1;
void tint_symbol_1(uvec3 GlobalInvocationID) {
ivec2 size = textureSize(src_1, 0);
ivec2 dstTexCoord = ivec2(GlobalInvocationID.xy);
ivec2 srcTexCoord = dstTexCoord;
uvec2 size = uvec2(textureSize(src_1, 0));
uvec2 dstTexCoord = GlobalInvocationID.xy;
uvec2 srcTexCoord = dstTexCoord;
if ((uniforms.dstTextureFlipY == 1u)) {
srcTexCoord.y = ((size.y - dstTexCoord.y) - 1);
srcTexCoord.y = ((size.y - dstTexCoord.y) - 1u);
}
vec4 srcColor = texelFetch(src_1, srcTexCoord, 0);
vec4 dstColor = texelFetch(dst_1, dstTexCoord, 0);
vec4 srcColor = texelFetch(src_1, ivec2(srcTexCoord), int(0u));
vec4 dstColor = texelFetch(dst_1, ivec2(dstTexCoord), int(0u));
bool success = true;
uvec4 srcColorBits = uvec4(0u, 0u, 0u, 0u);
uvec4 dstColorBits = uvec4(dstColor);

View File

@ -30,14 +30,14 @@ uint ConvertToFp16FloatValue(float fp32) {
}
void tint_symbol_inner(uint3 GlobalInvocationID, texture2d<float, access::sample> tint_symbol_2, const constant Uniforms* const tint_symbol_3, texture2d<float, access::sample> tint_symbol_4, device OutputBuf* const tint_symbol_5) {
int2 size = int2(tint_symbol_2.get_width(), tint_symbol_2.get_height());
int2 dstTexCoord = int2(uint3(GlobalInvocationID).xy);
int2 srcTexCoord = dstTexCoord;
uint2 size = uint2(tint_symbol_2.get_width(), tint_symbol_2.get_height());
uint2 dstTexCoord = uint3(GlobalInvocationID).xy;
uint2 srcTexCoord = dstTexCoord;
if (((*(tint_symbol_3)).dstTextureFlipY == 1u)) {
srcTexCoord[1] = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(size[1]) - as_type<uint>(dstTexCoord[1])))) - as_type<uint>(1)));
srcTexCoord[1] = ((size[1] - dstTexCoord[1]) - 1u);
}
float4 srcColor = tint_symbol_2.read(uint2(srcTexCoord), 0);
float4 dstColor = tint_symbol_4.read(uint2(dstTexCoord), 0);
float4 srcColor = tint_symbol_2.read(uint2(srcTexCoord), 0u);
float4 dstColor = tint_symbol_4.read(uint2(dstTexCoord), 0u);
bool success = true;
uint4 srcColorBits = 0u;
uint4 dstColorBits = uint4(dstColor);

View File

@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 138
; Bound: 133
; Schema: 0
OpCapability Shader
OpCapability ImageQuery
@ -73,33 +73,29 @@
%uint_1 = OpConstant %uint 1
%void = OpTypeVoid
%22 = OpTypeFunction %void %v3uint
%int = OpTypeInt 32 1
%v2int = OpTypeVector %int 2
%int_0 = OpConstant %int 0
%_ptr_Function_v2int = OpTypePointer Function %v2int
%34 = OpConstantNull %v2int
%v2uint = OpTypeVector %uint 2
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%_ptr_Function_v2uint = OpTypePointer Function %v2uint
%34 = OpConstantNull %v2uint
%uint_0 = OpConstant %uint 0
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
%bool = OpTypeBool
%_ptr_Function_int = OpTypePointer Function %int
%int_1 = OpConstant %int 1
%_ptr_Function_uint = OpTypePointer Function %uint
%v4float = OpTypeVector %float 4
%62 = OpConstantNull %int
%59 = OpConstantNull %uint
%_ptr_Function_v4float = OpTypePointer Function %v4float
%65 = OpConstantNull %v4float
%62 = OpConstantNull %v4float
%true = OpConstantTrue %bool
%_ptr_Function_bool = OpTypePointer Function %bool
%73 = OpConstantNull %bool
%70 = OpConstantNull %bool
%v4uint = OpTypeVector %uint 4
%_ptr_Function_v4uint = OpTypePointer Function %v4uint
%77 = OpConstantNull %v4uint
%81 = OpConstantNull %uint
%_ptr_Function_uint = OpTypePointer Function %uint
%74 = OpConstantNull %v4uint
%uint_3 = OpConstant %uint 3
%_ptr_Function_float = OpTypePointer Function %float
%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
%133 = OpTypeFunction %void
%128 = OpTypeFunction %void
%ConvertToFp16FloatValue = OpFunction %uint None %17
%fp32 = OpFunctionParameter %float
%20 = OpLabel
@ -108,124 +104,122 @@
%main_inner = OpFunction %void None %22
%GlobalInvocationID = OpFunctionParameter %v3uint
%26 = OpLabel
%size = OpVariable %_ptr_Function_v2int Function %34
%dstTexCoord = OpVariable %_ptr_Function_v2int Function %34
%srcTexCoord = OpVariable %_ptr_Function_v2int Function %34
%srcColor = OpVariable %_ptr_Function_v4float Function %65
%dstColor = OpVariable %_ptr_Function_v4float Function %65
%success = OpVariable %_ptr_Function_bool Function %73
%srcColorBits = OpVariable %_ptr_Function_v4uint Function %77
%dstColorBits = OpVariable %_ptr_Function_v4uint Function %77
%i = OpVariable %_ptr_Function_uint Function %81
%outputIndex = OpVariable %_ptr_Function_uint Function %81
%30 = OpLoad %7 %src
%27 = OpImageQuerySizeLod %v2int %30 %int_0
%size = OpVariable %_ptr_Function_v2uint Function %34
%dstTexCoord = OpVariable %_ptr_Function_v2uint Function %34
%srcTexCoord = OpVariable %_ptr_Function_v2uint Function %34
%srcColor = OpVariable %_ptr_Function_v4float Function %62
%dstColor = OpVariable %_ptr_Function_v4float Function %62
%success = OpVariable %_ptr_Function_bool Function %70
%srcColorBits = OpVariable %_ptr_Function_v4uint Function %74
%dstColorBits = OpVariable %_ptr_Function_v4uint Function %74
%i = OpVariable %_ptr_Function_uint Function %59
%outputIndex = OpVariable %_ptr_Function_uint Function %59
%29 = OpLoad %7 %src
%27 = OpImageQuerySizeLod %v2uint %29 %int_0
OpStore %size %27
%37 = OpVectorShuffle %v2uint %GlobalInvocationID %GlobalInvocationID 0 1
%35 = OpBitcast %v2int %37
%35 = OpVectorShuffle %v2uint %GlobalInvocationID %GlobalInvocationID 0 1
OpStore %dstTexCoord %35
%39 = OpLoad %v2int %dstTexCoord
OpStore %srcTexCoord %39
%43 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
%44 = OpLoad %uint %43
%45 = OpIEqual %bool %44 %uint_1
OpSelectionMerge %47 None
OpBranchConditional %45 %48 %47
%48 = OpLabel
%50 = OpAccessChain %_ptr_Function_int %srcTexCoord %uint_1
%51 = OpAccessChain %_ptr_Function_int %size %uint_1
%52 = OpLoad %int %51
%53 = OpAccessChain %_ptr_Function_int %dstTexCoord %uint_1
%54 = OpLoad %int %53
%55 = OpISub %int %52 %54
%57 = OpISub %int %55 %int_1
OpStore %50 %57
OpBranch %47
%47 = OpLabel
%60 = OpLoad %7 %src
%61 = OpLoad %v2int %srcTexCoord
%58 = OpImageFetch %v4float %60 %61 Lod %62
OpStore %srcColor %58
%67 = OpLoad %7 %dst
%68 = OpLoad %v2int %dstTexCoord
%66 = OpImageFetch %v4float %67 %68 Lod %62
OpStore %dstColor %66
%37 = OpLoad %v2uint %dstTexCoord
OpStore %srcTexCoord %37
%41 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
%42 = OpLoad %uint %41
%43 = OpIEqual %bool %42 %uint_1
OpSelectionMerge %45 None
OpBranchConditional %43 %46 %45
%46 = OpLabel
%48 = OpAccessChain %_ptr_Function_uint %srcTexCoord %uint_1
%49 = OpAccessChain %_ptr_Function_uint %size %uint_1
%50 = OpLoad %uint %49
%51 = OpAccessChain %_ptr_Function_uint %dstTexCoord %uint_1
%52 = OpLoad %uint %51
%53 = OpISub %uint %50 %52
%54 = OpISub %uint %53 %uint_1
OpStore %48 %54
OpBranch %45
%45 = OpLabel
%57 = OpLoad %7 %src
%58 = OpLoad %v2uint %srcTexCoord
%55 = OpImageFetch %v4float %57 %58 Lod %59
OpStore %srcColor %55
%64 = OpLoad %7 %dst
%65 = OpLoad %v2uint %dstTexCoord
%63 = OpImageFetch %v4float %64 %65 Lod %59
OpStore %dstColor %63
OpStore %success %true
%79 = OpLoad %v4float %dstColor
%78 = OpConvertFToU %v4uint %79
OpStore %dstColorBits %78
OpStore %i %81
OpBranch %84
%84 = OpLabel
OpLoopMerge %85 %86 None
OpBranch %87
%87 = OpLabel
%89 = OpLoad %uint %i
%91 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_3
%92 = OpLoad %uint %91
%93 = OpULessThan %bool %89 %92
%88 = OpLogicalNot %bool %93
OpSelectionMerge %94 None
OpBranchConditional %88 %95 %94
%95 = OpLabel
OpBranch %85
%94 = OpLabel
%97 = OpLoad %uint %i
%99 = OpAccessChain %_ptr_Function_float %srcColor %97
%100 = OpLoad %float %99
%96 = OpFunctionCall %uint %ConvertToFp16FloatValue %100
%76 = OpLoad %v4float %dstColor
%75 = OpConvertFToU %v4uint %76
OpStore %dstColorBits %75
OpStore %i %59
OpBranch %79
%79 = OpLabel
OpLoopMerge %80 %81 None
OpBranch %82
%82 = OpLabel
%84 = OpLoad %uint %i
%86 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_3
%87 = OpLoad %uint %86
%88 = OpULessThan %bool %84 %87
%83 = OpLogicalNot %bool %88
OpSelectionMerge %89 None
OpBranchConditional %83 %90 %89
%90 = OpLabel
OpBranch %80
%89 = OpLabel
%92 = OpLoad %uint %i
%94 = OpAccessChain %_ptr_Function_float %srcColor %92
%95 = OpLoad %float %94
%91 = OpFunctionCall %uint %ConvertToFp16FloatValue %95
%96 = OpLoad %uint %i
%97 = OpAccessChain %_ptr_Function_uint %srcColorBits %96
OpStore %97 %91
%98 = OpLoad %bool %success
OpSelectionMerge %99 None
OpBranchConditional %98 %100 %99
%100 = OpLabel
%101 = OpLoad %uint %i
%102 = OpAccessChain %_ptr_Function_uint %srcColorBits %101
OpStore %102 %96
%103 = OpLoad %bool %success
OpSelectionMerge %104 None
OpBranchConditional %103 %105 %104
%105 = OpLabel
%106 = OpLoad %uint %i
%107 = OpAccessChain %_ptr_Function_uint %srcColorBits %106
%108 = OpLoad %uint %107
%103 = OpLoad %uint %102
%104 = OpLoad %uint %i
%105 = OpAccessChain %_ptr_Function_uint %dstColorBits %104
%106 = OpLoad %uint %105
%107 = OpIEqual %bool %103 %106
OpBranch %99
%99 = OpLabel
%108 = OpPhi %bool %98 %89 %107 %100
OpStore %success %108
OpBranch %81
%81 = OpLabel
%109 = OpLoad %uint %i
%110 = OpAccessChain %_ptr_Function_uint %dstColorBits %109
%111 = OpLoad %uint %110
%112 = OpIEqual %bool %108 %111
OpBranch %104
%104 = OpLabel
%113 = OpPhi %bool %103 %94 %112 %105
OpStore %success %113
OpBranch %86
%86 = OpLabel
%114 = OpLoad %uint %i
%115 = OpIAdd %uint %114 %uint_1
OpStore %i %115
OpBranch %84
%85 = OpLabel
%116 = OpCompositeExtract %uint %GlobalInvocationID 1
%118 = OpAccessChain %_ptr_Function_int %size %uint_0
%119 = OpLoad %int %118
%117 = OpBitcast %uint %119
%120 = OpIMul %uint %116 %117
%121 = OpCompositeExtract %uint %GlobalInvocationID 0
%122 = OpIAdd %uint %120 %121
OpStore %outputIndex %122
%124 = OpLoad %bool %success
OpSelectionMerge %125 None
OpBranchConditional %124 %126 %127
%126 = OpLabel
%128 = OpLoad %uint %outputIndex
%130 = OpAccessChain %_ptr_StorageBuffer_uint %output %uint_0 %128
OpStore %130 %uint_1
OpBranch %125
%127 = OpLabel
%131 = OpLoad %uint %outputIndex
%132 = OpAccessChain %_ptr_StorageBuffer_uint %output %uint_0 %131
OpStore %132 %81
OpBranch %125
%125 = OpLabel
%110 = OpIAdd %uint %109 %uint_1
OpStore %i %110
OpBranch %79
%80 = OpLabel
%111 = OpCompositeExtract %uint %GlobalInvocationID 1
%113 = OpAccessChain %_ptr_Function_uint %size %uint_0
%114 = OpLoad %uint %113
%115 = OpIMul %uint %111 %114
%116 = OpCompositeExtract %uint %GlobalInvocationID 0
%117 = OpIAdd %uint %115 %116
OpStore %outputIndex %117
%119 = OpLoad %bool %success
OpSelectionMerge %120 None
OpBranchConditional %119 %121 %122
%121 = OpLabel
%123 = OpLoad %uint %outputIndex
%125 = OpAccessChain %_ptr_StorageBuffer_uint %output %uint_0 %123
OpStore %125 %uint_1
OpBranch %120
%122 = OpLabel
%126 = OpLoad %uint %outputIndex
%127 = OpAccessChain %_ptr_StorageBuffer_uint %output %uint_0 %126
OpStore %127 %59
OpBranch %120
%120 = OpLabel
OpReturn
OpFunctionEnd
%main = OpFunction %void None %133
%135 = OpLabel
%137 = OpLoad %v3uint %GlobalInvocationID_1
%136 = OpFunctionCall %void %main_inner %137
%main = OpFunction %void None %128
%130 = OpLabel
%132 = OpLoad %v3uint %GlobalInvocationID_1
%131 = OpFunctionCall %void %main_inner %132
OpReturn
OpFunctionEnd

View File

@ -23,9 +23,9 @@ fn ConvertToFp16FloatValue(fp32 : f32) -> u32 {
@compute @workgroup_size(1, 1, 1)
fn main(@builtin(global_invocation_id) GlobalInvocationID : vec3<u32>) {
var size : vec2<i32> = textureDimensions(src);
var dstTexCoord : vec2<i32> = vec2<i32>(GlobalInvocationID.xy);
var srcTexCoord : vec2<i32> = dstTexCoord;
var size = textureDimensions(src);
var dstTexCoord = GlobalInvocationID.xy;
var srcTexCoord = dstTexCoord;
if ((uniforms.dstTextureFlipY == 1u)) {
srcTexCoord.y = ((size.y - dstTexCoord.y) - 1);
}

File diff suppressed because it is too large Load Diff

View File

@ -18,8 +18,8 @@ fn aboutEqual(value : f32, expect : f32) -> bool {
}
@compute @workgroup_size(1, 1, 1)
fn main(@builtin(global_invocation_id) GlobalInvocationID : vec3<u32>) {
let srcSize : vec2<i32> = textureDimensions(src);
let dstSize : vec2<i32> = textureDimensions(dst);
let srcSize = textureDimensions(src);
let dstSize = textureDimensions(dst);
let dstTexCoord : vec2<u32> = vec2<u32>(GlobalInvocationID.xy);
let nonCoveredColor : vec4<f32> =
vec4<f32>(0.0, 1.0, 0.0, 1.0); // should be green
@ -40,7 +40,7 @@ fn main(@builtin(global_invocation_id) GlobalInvocationID : vec3<u32>) {
// which is flipped and need to unpack flip during the copy.
// We need to calculate the expect y coord based on this rule.
if (uniforms.dstTextureFlipY == 1u) {
srcTexCoord.y = u32(srcSize.y) - srcTexCoord.y - 1u;
srcTexCoord.y = srcSize.y - srcTexCoord.y - 1u;
}
let srcColor : vec4<f32> = textureLoad(src, vec2<i32>(srcTexCoord), 0);
@ -60,7 +60,7 @@ fn main(@builtin(global_invocation_id) GlobalInvocationID : vec3<u32>) {
aboutEqual(dstColor.a, srcColor.a);
}
}
let outputIndex : u32 = GlobalInvocationID.y * u32(dstSize.x) +
let outputIndex : u32 = GlobalInvocationID.y * dstSize.x +
GlobalInvocationID.x;
if (success) {
output.result[outputIndex] = 1u;

View File

@ -16,10 +16,10 @@ struct tint_symbol_2 {
void main_inner(uint3 GlobalInvocationID) {
int2 tint_tmp;
src.GetDimensions(tint_tmp.x, tint_tmp.y);
const int2 srcSize = tint_tmp;
const uint2 srcSize = tint_tmp;
int2 tint_tmp_1;
tint_symbol.GetDimensions(tint_tmp_1.x, tint_tmp_1.y);
const int2 dstSize = tint_tmp_1;
const uint2 dstSize = tint_tmp_1;
const uint2 dstTexCoord = uint2(GlobalInvocationID.xy);
const float4 nonCoveredColor = float4(0.0f, 1.0f, 0.0f, 1.0f);
bool success = true;
@ -44,7 +44,7 @@ void main_inner(uint3 GlobalInvocationID) {
} else {
uint2 srcTexCoord = ((dstTexCoord - uniforms[1].xy) + uniforms[0].zw);
if ((uniforms[0].x == 1u)) {
srcTexCoord.y = ((uint(srcSize.y) - srcTexCoord.y) - 1u);
srcTexCoord.y = ((srcSize.y - srcTexCoord.y) - 1u);
}
const float4 srcColor = src.Load(int3(int2(srcTexCoord), 0));
const float4 dstColor = tint_symbol.Load(int3(int2(dstTexCoord), 0));
@ -78,7 +78,7 @@ void main_inner(uint3 GlobalInvocationID) {
success = tint_symbol_5;
}
}
const uint outputIndex = ((GlobalInvocationID.y * uint(dstSize.x)) + GlobalInvocationID.x);
const uint outputIndex = ((GlobalInvocationID.y * dstSize.x) + GlobalInvocationID.x);
if (success) {
output.Store((4u * outputIndex), asuint(1u));
} else {

View File

@ -16,10 +16,10 @@ struct tint_symbol_2 {
void main_inner(uint3 GlobalInvocationID) {
int2 tint_tmp;
src.GetDimensions(tint_tmp.x, tint_tmp.y);
const int2 srcSize = tint_tmp;
const uint2 srcSize = tint_tmp;
int2 tint_tmp_1;
tint_symbol.GetDimensions(tint_tmp_1.x, tint_tmp_1.y);
const int2 dstSize = tint_tmp_1;
const uint2 dstSize = tint_tmp_1;
const uint2 dstTexCoord = uint2(GlobalInvocationID.xy);
const float4 nonCoveredColor = float4(0.0f, 1.0f, 0.0f, 1.0f);
bool success = true;
@ -44,7 +44,7 @@ void main_inner(uint3 GlobalInvocationID) {
} else {
uint2 srcTexCoord = ((dstTexCoord - uniforms[1].xy) + uniforms[0].zw);
if ((uniforms[0].x == 1u)) {
srcTexCoord.y = ((uint(srcSize.y) - srcTexCoord.y) - 1u);
srcTexCoord.y = ((srcSize.y - srcTexCoord.y) - 1u);
}
const float4 srcColor = src.Load(int3(int2(srcTexCoord), 0));
const float4 dstColor = tint_symbol.Load(int3(int2(dstTexCoord), 0));
@ -78,7 +78,7 @@ void main_inner(uint3 GlobalInvocationID) {
success = tint_symbol_5;
}
}
const uint outputIndex = ((GlobalInvocationID.y * uint(dstSize.x)) + GlobalInvocationID.x);
const uint outputIndex = ((GlobalInvocationID.y * dstSize.x) + GlobalInvocationID.x);
if (success) {
output.Store((4u * outputIndex), asuint(1u));
} else {

View File

@ -19,8 +19,8 @@ bool aboutEqual(float value, float expect) {
uniform highp sampler2D src_1;
uniform highp sampler2D dst_1;
void tint_symbol_1(uvec3 GlobalInvocationID) {
ivec2 srcSize = textureSize(src_1, 0);
ivec2 dstSize = textureSize(dst_1, 0);
uvec2 srcSize = uvec2(textureSize(src_1, 0));
uvec2 dstSize = uvec2(textureSize(dst_1, 0));
uvec2 dstTexCoord = uvec2(GlobalInvocationID.xy);
vec4 nonCoveredColor = vec4(0.0f, 1.0f, 0.0f, 1.0f);
bool success = true;
@ -45,7 +45,7 @@ void tint_symbol_1(uvec3 GlobalInvocationID) {
} else {
uvec2 srcTexCoord = ((dstTexCoord - uniforms.dstCopyOrigin) + uniforms.srcCopyOrigin);
if ((uniforms.dstTextureFlipY == 1u)) {
srcTexCoord.y = ((uint(srcSize.y) - srcTexCoord.y) - 1u);
srcTexCoord.y = ((srcSize.y - srcTexCoord.y) - 1u);
}
vec4 srcColor = texelFetch(src_1, ivec2(srcTexCoord), 0);
vec4 dstColor = texelFetch(dst_1, ivec2(dstTexCoord), 0);
@ -79,7 +79,7 @@ void tint_symbol_1(uvec3 GlobalInvocationID) {
success = tint_symbol_4;
}
}
uint outputIndex = ((GlobalInvocationID.y * uint(dstSize.x)) + GlobalInvocationID.x);
uint outputIndex = ((GlobalInvocationID.y * dstSize.x) + GlobalInvocationID.x);
if (success) {
tint_symbol.result[outputIndex] = 1u;
} else {

View File

@ -31,8 +31,8 @@ bool aboutEqual(float value, float expect) {
}
void tint_symbol_inner(uint3 GlobalInvocationID, texture2d<float, access::sample> tint_symbol_7, texture2d<float, access::sample> tint_symbol_8, const constant Uniforms* const tint_symbol_9, device OutputBuf* const tint_symbol_10) {
int2 const srcSize = int2(tint_symbol_7.get_width(), tint_symbol_7.get_height());
int2 const dstSize = int2(tint_symbol_8.get_width(), tint_symbol_8.get_height());
uint2 const srcSize = uint2(tint_symbol_7.get_width(), tint_symbol_7.get_height());
uint2 const dstSize = uint2(tint_symbol_8.get_width(), tint_symbol_8.get_height());
uint2 const dstTexCoord = uint2(uint3(GlobalInvocationID).xy);
float4 const nonCoveredColor = float4(0.0f, 1.0f, 0.0f, 1.0f);
bool success = true;
@ -41,7 +41,7 @@ void tint_symbol_inner(uint3 GlobalInvocationID, texture2d<float, access::sample
} else {
uint2 srcTexCoord = ((dstTexCoord - (*(tint_symbol_9)).dstCopyOrigin) + (*(tint_symbol_9)).srcCopyOrigin);
if (((*(tint_symbol_9)).dstTextureFlipY == 1u)) {
srcTexCoord[1] = ((uint(srcSize[1]) - srcTexCoord[1]) - 1u);
srcTexCoord[1] = ((srcSize[1] - srcTexCoord[1]) - 1u);
}
float4 const srcColor = tint_symbol_7.read(uint2(int2(srcTexCoord)), 0);
float4 const dstColor = tint_symbol_8.read(uint2(int2(dstTexCoord)), 0);
@ -75,7 +75,7 @@ void tint_symbol_inner(uint3 GlobalInvocationID, texture2d<float, access::sample
success = tint_symbol_3;
}
}
uint const outputIndex = ((GlobalInvocationID[1] * uint(dstSize[0])) + GlobalInvocationID[0]);
uint const outputIndex = ((GlobalInvocationID[1] * dstSize[0]) + GlobalInvocationID[0]);
if (success) {
(*(tint_symbol_10)).result[outputIndex] = 1u;
} else {

View File

@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 207
; Bound: 205
; Schema: 0
OpCapability Shader
OpCapability ImageQuery
@ -78,20 +78,20 @@
%void = OpTypeVoid
%29 = OpTypeFunction %void %v3uint
%int = OpTypeInt 32 1
%v2int = OpTypeVector %int 2
%int_0 = OpConstant %int 0
%v4float = OpTypeVector %float 4
%44 = OpConstantNull %float
%43 = OpConstantNull %float
%float_1 = OpConstant %float 1
%46 = OpConstantComposite %v4float %44 %float_1 %44 %float_1
%45 = OpConstantComposite %v4float %43 %float_1 %43 %float_1
%true = OpConstantTrue %bool
%_ptr_Function_bool = OpTypePointer Function %bool
%50 = OpConstantNull %bool
%49 = OpConstantNull %bool
%uint_3 = OpConstant %uint 3
%uint_0 = OpConstant %uint 0
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
%uint_1 = OpConstant %uint 1
%uint_4 = OpConstant %uint 4
%v2int = OpTypeVector %int 2
%97 = OpConstantNull %int
%v4bool = OpTypeVector %bool 4
%_ptr_Uniform_v2uint = OpTypePointer Uniform %v2uint
@ -100,8 +100,8 @@
%111 = OpConstantNull %v2uint
%_ptr_Function_uint = OpTypePointer Function %uint
%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
%201 = OpConstantNull %uint
%202 = OpTypeFunction %void
%199 = OpConstantNull %uint
%200 = OpTypeFunction %void
%aboutEqual = OpFunction %bool None %18
%value = OpFunctionParameter %float
%expect = OpFunctionParameter %float
@ -114,81 +114,81 @@
%main_inner = OpFunction %void None %29
%GlobalInvocationID = OpFunctionParameter %v3uint
%33 = OpLabel
%success = OpVariable %_ptr_Function_bool Function %50
%success = OpVariable %_ptr_Function_bool Function %49
%srcTexCoord = OpVariable %_ptr_Function_v2uint Function %111
%tint_symbol_1 = OpVariable %_ptr_Function_bool Function %50
%tint_symbol = OpVariable %_ptr_Function_bool Function %50
%tint_symbol_5 = OpVariable %_ptr_Function_bool Function %50
%tint_symbol_4 = OpVariable %_ptr_Function_bool Function %50
%tint_symbol_3 = OpVariable %_ptr_Function_bool Function %50
%tint_symbol_2 = OpVariable %_ptr_Function_bool Function %50
%37 = OpLoad %7 %src
%34 = OpImageQuerySizeLod %v2int %37 %int_0
%40 = OpLoad %7 %dst
%39 = OpImageQuerySizeLod %v2int %40 %int_0
%42 = OpVectorShuffle %v2uint %GlobalInvocationID %GlobalInvocationID 0 1
%tint_symbol_1 = OpVariable %_ptr_Function_bool Function %49
%tint_symbol = OpVariable %_ptr_Function_bool Function %49
%tint_symbol_5 = OpVariable %_ptr_Function_bool Function %49
%tint_symbol_4 = OpVariable %_ptr_Function_bool Function %49
%tint_symbol_3 = OpVariable %_ptr_Function_bool Function %49
%tint_symbol_2 = OpVariable %_ptr_Function_bool Function %49
%35 = OpLoad %7 %src
%34 = OpImageQuerySizeLod %v2uint %35 %int_0
%39 = OpLoad %7 %dst
%38 = OpImageQuerySizeLod %v2uint %39 %int_0
%41 = OpVectorShuffle %v2uint %GlobalInvocationID %GlobalInvocationID 0 1
OpStore %success %true
%51 = OpCompositeExtract %uint %42 0
%55 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_3 %uint_0
%56 = OpLoad %uint %55
%57 = OpULessThan %bool %51 %56
OpSelectionMerge %58 None
OpBranchConditional %57 %58 %59
%59 = OpLabel
%60 = OpCompositeExtract %uint %42 1
%62 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_3 %uint_1
%63 = OpLoad %uint %62
%64 = OpULessThan %bool %60 %63
OpBranch %58
%50 = OpCompositeExtract %uint %41 0
%54 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_3 %uint_0
%55 = OpLoad %uint %54
%56 = OpULessThan %bool %50 %55
OpSelectionMerge %57 None
OpBranchConditional %56 %57 %58
%58 = OpLabel
%65 = OpPhi %bool %57 %33 %64 %59
OpSelectionMerge %66 None
OpBranchConditional %65 %66 %67
%67 = OpLabel
%68 = OpCompositeExtract %uint %42 0
%69 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_3 %uint_0
%70 = OpLoad %uint %69
%72 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_4 %uint_0
%73 = OpLoad %uint %72
%74 = OpIAdd %uint %70 %73
%75 = OpUGreaterThanEqual %bool %68 %74
OpBranch %66
%59 = OpCompositeExtract %uint %41 1
%61 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_3 %uint_1
%62 = OpLoad %uint %61
%63 = OpULessThan %bool %59 %62
OpBranch %57
%57 = OpLabel
%64 = OpPhi %bool %56 %33 %63 %58
OpSelectionMerge %65 None
OpBranchConditional %64 %65 %66
%66 = OpLabel
%76 = OpPhi %bool %65 %58 %75 %67
OpSelectionMerge %77 None
OpBranchConditional %76 %77 %78
%78 = OpLabel
%79 = OpCompositeExtract %uint %42 1
%80 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_3 %uint_1
%81 = OpLoad %uint %80
%82 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_4 %uint_1
%83 = OpLoad %uint %82
%84 = OpIAdd %uint %81 %83
%85 = OpUGreaterThanEqual %bool %79 %84
OpBranch %77
%67 = OpCompositeExtract %uint %41 0
%68 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_3 %uint_0
%69 = OpLoad %uint %68
%71 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_4 %uint_0
%72 = OpLoad %uint %71
%73 = OpIAdd %uint %69 %72
%74 = OpUGreaterThanEqual %bool %67 %73
OpBranch %65
%65 = OpLabel
%75 = OpPhi %bool %64 %57 %74 %66
OpSelectionMerge %76 None
OpBranchConditional %75 %76 %77
%77 = OpLabel
%86 = OpPhi %bool %76 %66 %85 %78
OpSelectionMerge %87 None
OpBranchConditional %86 %88 %89
%88 = OpLabel
%90 = OpLoad %bool %success
OpSelectionMerge %91 None
OpBranchConditional %90 %92 %91
%92 = OpLabel
%95 = OpLoad %7 %dst
%96 = OpBitcast %v2int %42
%94 = OpImageFetch %v4float %95 %96 Lod %97
%98 = OpFOrdEqual %v4bool %94 %46
%93 = OpAll %bool %98
OpBranch %91
%78 = OpCompositeExtract %uint %41 1
%79 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_3 %uint_1
%80 = OpLoad %uint %79
%81 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_4 %uint_1
%82 = OpLoad %uint %81
%83 = OpIAdd %uint %80 %82
%84 = OpUGreaterThanEqual %bool %78 %83
OpBranch %76
%76 = OpLabel
%85 = OpPhi %bool %75 %65 %84 %77
OpSelectionMerge %86 None
OpBranchConditional %85 %87 %88
%87 = OpLabel
%89 = OpLoad %bool %success
OpSelectionMerge %90 None
OpBranchConditional %89 %91 %90
%91 = OpLabel
%100 = OpPhi %bool %90 %88 %93 %92
%94 = OpLoad %7 %dst
%95 = OpBitcast %v2int %41
%93 = OpImageFetch %v4float %94 %95 Lod %97
%98 = OpFOrdEqual %v4bool %93 %45
%92 = OpAll %bool %98
OpBranch %90
%90 = OpLabel
%100 = OpPhi %bool %89 %87 %92 %91
OpStore %success %100
OpBranch %87
%89 = OpLabel
OpBranch %86
%88 = OpLabel
%102 = OpAccessChain %_ptr_Uniform_v2uint %uniforms %uint_3
%103 = OpLoad %v2uint %102
%104 = OpISub %v2uint %42 %103
%104 = OpISub %v2uint %41 %103
%106 = OpAccessChain %_ptr_Uniform_v2uint %uniforms %uint_2
%107 = OpLoad %v2uint %106
%108 = OpIAdd %v2uint %104 %107
@ -200,133 +200,131 @@
OpBranchConditional %114 %116 %115
%116 = OpLabel
%118 = OpAccessChain %_ptr_Function_uint %srcTexCoord %uint_1
%120 = OpCompositeExtract %int %34 1
%119 = OpBitcast %uint %120
%121 = OpAccessChain %_ptr_Function_uint %srcTexCoord %uint_1
%122 = OpLoad %uint %121
%123 = OpISub %uint %119 %122
%124 = OpISub %uint %123 %uint_1
OpStore %118 %124
%119 = OpCompositeExtract %uint %34 1
%120 = OpAccessChain %_ptr_Function_uint %srcTexCoord %uint_1
%121 = OpLoad %uint %120
%122 = OpISub %uint %119 %121
%123 = OpISub %uint %122 %uint_1
OpStore %118 %123
OpBranch %115
%115 = OpLabel
%126 = OpLoad %7 %src
%128 = OpLoad %v2uint %srcTexCoord
%127 = OpBitcast %v2int %128
%125 = OpImageFetch %v4float %126 %127 Lod %97
%130 = OpLoad %7 %dst
%131 = OpBitcast %v2int %42
%129 = OpImageFetch %v4float %130 %131 Lod %97
%132 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
%133 = OpLoad %uint %132
%134 = OpIEqual %bool %133 %uint_2
OpSelectionMerge %135 None
OpBranchConditional %134 %136 %137
%136 = OpLabel
%138 = OpLoad %bool %success
OpStore %tint_symbol_1 %138
%140 = OpLoad %bool %tint_symbol_1
OpSelectionMerge %141 None
OpBranchConditional %140 %142 %141
%142 = OpLabel
%144 = OpCompositeExtract %float %129 0
%145 = OpCompositeExtract %float %125 0
%143 = OpFunctionCall %bool %aboutEqual %144 %145
OpStore %tint_symbol_1 %143
OpBranch %141
%141 = OpLabel
%146 = OpLoad %bool %tint_symbol_1
OpStore %tint_symbol %146
%148 = OpLoad %bool %tint_symbol
OpSelectionMerge %149 None
OpBranchConditional %148 %150 %149
%150 = OpLabel
%152 = OpCompositeExtract %float %129 1
%153 = OpCompositeExtract %float %125 1
%151 = OpFunctionCall %bool %aboutEqual %152 %153
OpStore %tint_symbol %151
OpBranch %149
%149 = OpLabel
%154 = OpLoad %bool %tint_symbol
OpStore %success %154
OpBranch %135
%137 = OpLabel
%155 = OpLoad %bool %success
OpStore %tint_symbol_5 %155
%157 = OpLoad %bool %tint_symbol_5
OpSelectionMerge %158 None
OpBranchConditional %157 %159 %158
%159 = OpLabel
%161 = OpCompositeExtract %float %129 0
%162 = OpCompositeExtract %float %125 0
%160 = OpFunctionCall %bool %aboutEqual %161 %162
OpStore %tint_symbol_5 %160
OpBranch %158
%158 = OpLabel
%163 = OpLoad %bool %tint_symbol_5
OpStore %tint_symbol_4 %163
%165 = OpLoad %bool %tint_symbol_4
OpSelectionMerge %166 None
OpBranchConditional %165 %167 %166
%167 = OpLabel
%169 = OpCompositeExtract %float %129 1
%170 = OpCompositeExtract %float %125 1
%168 = OpFunctionCall %bool %aboutEqual %169 %170
OpStore %tint_symbol_4 %168
OpBranch %166
%166 = OpLabel
%171 = OpLoad %bool %tint_symbol_4
OpStore %tint_symbol_3 %171
%173 = OpLoad %bool %tint_symbol_3
OpSelectionMerge %174 None
OpBranchConditional %173 %175 %174
%175 = OpLabel
%177 = OpCompositeExtract %float %129 2
%178 = OpCompositeExtract %float %125 2
%176 = OpFunctionCall %bool %aboutEqual %177 %178
OpStore %tint_symbol_3 %176
OpBranch %174
%174 = OpLabel
%179 = OpLoad %bool %tint_symbol_3
OpStore %tint_symbol_2 %179
%181 = OpLoad %bool %tint_symbol_2
OpSelectionMerge %182 None
OpBranchConditional %181 %183 %182
%183 = OpLabel
%185 = OpCompositeExtract %float %129 3
%186 = OpCompositeExtract %float %125 3
%184 = OpFunctionCall %bool %aboutEqual %185 %186
OpStore %tint_symbol_2 %184
OpBranch %182
%182 = OpLabel
%187 = OpLoad %bool %tint_symbol_2
OpStore %success %187
OpBranch %135
%125 = OpLoad %7 %src
%127 = OpLoad %v2uint %srcTexCoord
%126 = OpBitcast %v2int %127
%124 = OpImageFetch %v4float %125 %126 Lod %97
%129 = OpLoad %7 %dst
%130 = OpBitcast %v2int %41
%128 = OpImageFetch %v4float %129 %130 Lod %97
%131 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
%132 = OpLoad %uint %131
%133 = OpIEqual %bool %132 %uint_2
OpSelectionMerge %134 None
OpBranchConditional %133 %135 %136
%135 = OpLabel
OpBranch %87
%87 = OpLabel
%188 = OpCompositeExtract %uint %GlobalInvocationID 1
%190 = OpCompositeExtract %int %39 0
%189 = OpBitcast %uint %190
%191 = OpIMul %uint %188 %189
%192 = OpCompositeExtract %uint %GlobalInvocationID 0
%193 = OpIAdd %uint %191 %192
%194 = OpLoad %bool %success
OpSelectionMerge %195 None
OpBranchConditional %194 %196 %197
%196 = OpLabel
%199 = OpAccessChain %_ptr_StorageBuffer_uint %output %uint_0 %193
OpStore %199 %uint_1
OpBranch %195
%197 = OpLabel
%200 = OpAccessChain %_ptr_StorageBuffer_uint %output %uint_0 %193
OpStore %200 %201
OpBranch %195
%137 = OpLoad %bool %success
OpStore %tint_symbol_1 %137
%139 = OpLoad %bool %tint_symbol_1
OpSelectionMerge %140 None
OpBranchConditional %139 %141 %140
%141 = OpLabel
%143 = OpCompositeExtract %float %128 0
%144 = OpCompositeExtract %float %124 0
%142 = OpFunctionCall %bool %aboutEqual %143 %144
OpStore %tint_symbol_1 %142
OpBranch %140
%140 = OpLabel
%145 = OpLoad %bool %tint_symbol_1
OpStore %tint_symbol %145
%147 = OpLoad %bool %tint_symbol
OpSelectionMerge %148 None
OpBranchConditional %147 %149 %148
%149 = OpLabel
%151 = OpCompositeExtract %float %128 1
%152 = OpCompositeExtract %float %124 1
%150 = OpFunctionCall %bool %aboutEqual %151 %152
OpStore %tint_symbol %150
OpBranch %148
%148 = OpLabel
%153 = OpLoad %bool %tint_symbol
OpStore %success %153
OpBranch %134
%136 = OpLabel
%154 = OpLoad %bool %success
OpStore %tint_symbol_5 %154
%156 = OpLoad %bool %tint_symbol_5
OpSelectionMerge %157 None
OpBranchConditional %156 %158 %157
%158 = OpLabel
%160 = OpCompositeExtract %float %128 0
%161 = OpCompositeExtract %float %124 0
%159 = OpFunctionCall %bool %aboutEqual %160 %161
OpStore %tint_symbol_5 %159
OpBranch %157
%157 = OpLabel
%162 = OpLoad %bool %tint_symbol_5
OpStore %tint_symbol_4 %162
%164 = OpLoad %bool %tint_symbol_4
OpSelectionMerge %165 None
OpBranchConditional %164 %166 %165
%166 = OpLabel
%168 = OpCompositeExtract %float %128 1
%169 = OpCompositeExtract %float %124 1
%167 = OpFunctionCall %bool %aboutEqual %168 %169
OpStore %tint_symbol_4 %167
OpBranch %165
%165 = OpLabel
%170 = OpLoad %bool %tint_symbol_4
OpStore %tint_symbol_3 %170
%172 = OpLoad %bool %tint_symbol_3
OpSelectionMerge %173 None
OpBranchConditional %172 %174 %173
%174 = OpLabel
%176 = OpCompositeExtract %float %128 2
%177 = OpCompositeExtract %float %124 2
%175 = OpFunctionCall %bool %aboutEqual %176 %177
OpStore %tint_symbol_3 %175
OpBranch %173
%173 = OpLabel
%178 = OpLoad %bool %tint_symbol_3
OpStore %tint_symbol_2 %178
%180 = OpLoad %bool %tint_symbol_2
OpSelectionMerge %181 None
OpBranchConditional %180 %182 %181
%182 = OpLabel
%184 = OpCompositeExtract %float %128 3
%185 = OpCompositeExtract %float %124 3
%183 = OpFunctionCall %bool %aboutEqual %184 %185
OpStore %tint_symbol_2 %183
OpBranch %181
%181 = OpLabel
%186 = OpLoad %bool %tint_symbol_2
OpStore %success %186
OpBranch %134
%134 = OpLabel
OpBranch %86
%86 = OpLabel
%187 = OpCompositeExtract %uint %GlobalInvocationID 1
%188 = OpCompositeExtract %uint %38 0
%189 = OpIMul %uint %187 %188
%190 = OpCompositeExtract %uint %GlobalInvocationID 0
%191 = OpIAdd %uint %189 %190
%192 = OpLoad %bool %success
OpSelectionMerge %193 None
OpBranchConditional %192 %194 %195
%194 = OpLabel
%197 = OpAccessChain %_ptr_StorageBuffer_uint %output %uint_0 %191
OpStore %197 %uint_1
OpBranch %193
%195 = OpLabel
%198 = OpAccessChain %_ptr_StorageBuffer_uint %output %uint_0 %191
OpStore %198 %199
OpBranch %193
%193 = OpLabel
OpReturn
OpFunctionEnd
%main = OpFunction %void None %202
%204 = OpLabel
%206 = OpLoad %v3uint %GlobalInvocationID_1
%205 = OpFunctionCall %void %main_inner %206
%main = OpFunction %void None %200
%202 = OpLabel
%204 = OpLoad %v3uint %GlobalInvocationID_1
%203 = OpFunctionCall %void %main_inner %204
OpReturn
OpFunctionEnd

View File

@ -24,8 +24,8 @@ fn aboutEqual(value : f32, expect : f32) -> bool {
@compute @workgroup_size(1, 1, 1)
fn main(@builtin(global_invocation_id) GlobalInvocationID : vec3<u32>) {
let srcSize : vec2<i32> = textureDimensions(src);
let dstSize : vec2<i32> = textureDimensions(dst);
let srcSize = textureDimensions(src);
let dstSize = textureDimensions(dst);
let dstTexCoord : vec2<u32> = vec2<u32>(GlobalInvocationID.xy);
let nonCoveredColor : vec4<f32> = vec4<f32>(0.0, 1.0, 0.0, 1.0);
var success : bool = true;
@ -34,7 +34,7 @@ fn main(@builtin(global_invocation_id) GlobalInvocationID : vec3<u32>) {
} else {
var srcTexCoord : vec2<u32> = ((dstTexCoord - uniforms.dstCopyOrigin) + uniforms.srcCopyOrigin);
if ((uniforms.dstTextureFlipY == 1u)) {
srcTexCoord.y = ((u32(srcSize.y) - srcTexCoord.y) - 1u);
srcTexCoord.y = ((srcSize.y - srcTexCoord.y) - 1u);
}
let srcColor : vec4<f32> = textureLoad(src, vec2<i32>(srcTexCoord), 0);
let dstColor : vec4<f32> = textureLoad(dst, vec2<i32>(dstTexCoord), 0);
@ -44,7 +44,7 @@ fn main(@builtin(global_invocation_id) GlobalInvocationID : vec3<u32>) {
success = ((((success && aboutEqual(dstColor.r, srcColor.r)) && aboutEqual(dstColor.g, srcColor.g)) && aboutEqual(dstColor.b, srcColor.b)) && aboutEqual(dstColor.a, srcColor.a));
}
}
let outputIndex : u32 = ((GlobalInvocationID.y * u32(dstSize.x)) + GlobalInvocationID.x);
let outputIndex : u32 = ((GlobalInvocationID.y * dstSize.x) + GlobalInvocationID.x);
if (success) {
output.result[outputIndex] = 1u;
} else {

View File

@ -1,4 +1,4 @@
struct Params {
struct Params {
filterDim : u32,
blockDim : u32,
};
@ -35,16 +35,13 @@ fn main(
@builtin(local_invocation_id) LocalInvocationID : vec3<u32>
) {
let filterOffset : u32 = (params.filterDim - 1u) / 2u;
let dims : vec2<i32> = textureDimensions(inputTex, 0);
let dims = textureDimensions(inputTex, 0);
let baseIndex = vec2<i32>(
WorkGroupID.xy * vec2<u32>(params.blockDim, 4u) +
LocalInvocationID.xy * vec2<u32>(4u, 1u)
) - vec2<i32>(i32(filterOffset), 0);
let baseIndex = (WorkGroupID.xy * vec2(params.blockDim, 4) + LocalInvocationID.xy * vec2(4u, 1u)) - vec2(filterOffset, 0);
for (var r : u32 = 0u; r < 4u; r = r + 1u) {
for (var c : u32 = 0u; c < 4u; c = c + 1u) {
var loadIndex = baseIndex + vec2<i32>(i32(c), i32(r));
var loadIndex = baseIndex + vec2(c, r);
if (flip.value != 0u) {
loadIndex = loadIndex.yx;
}
@ -59,7 +56,7 @@ fn main(
for (var r : u32 = 0u; r < 4u; r = r + 1u) {
for (var c : u32 = 0u; c < 4u; c = c + 1u) {
var writeIndex = baseIndex + vec2<i32>(i32(c), i32(r));
var writeIndex = baseIndex + vec2(c, r);
if (flip.value != 0u) {
writeIndex = writeIndex.yx;
}

View File

@ -28,13 +28,13 @@ void main_inner(uint3 WorkGroupID, uint3 LocalInvocationID, uint local_invocatio
const uint filterOffset = ((params[0].x - 1u) / 2u);
int3 tint_tmp;
inputTex.GetDimensions(0, tint_tmp.x, tint_tmp.y, tint_tmp.z);
const int2 dims = tint_tmp.xy;
const int2 baseIndex = (int2(((WorkGroupID.xy * uint2(params[0].y, 4u)) + (LocalInvocationID.xy * uint2(4u, 1u)))) - int2(int(filterOffset), 0));
const uint2 dims = tint_tmp.xy;
const uint2 baseIndex = (((WorkGroupID.xy * uint2(params[0].y, 4u)) + (LocalInvocationID.xy * uint2(4u, 1u))) - uint2(filterOffset, 0u));
{
for(uint r = 0u; (r < 4u); r = (r + 1u)) {
{
for(uint c = 0u; (c < 4u); c = (c + 1u)) {
int2 loadIndex = (baseIndex + int2(int(c), int(r)));
uint2 loadIndex = (baseIndex + uint2(c, r));
if ((flip[0].x != 0u)) {
loadIndex = loadIndex.yx;
}
@ -48,7 +48,7 @@ void main_inner(uint3 WorkGroupID, uint3 LocalInvocationID, uint local_invocatio
for(uint r = 0u; (r < 4u); r = (r + 1u)) {
{
for(uint c = 0u; (c < 4u); c = (c + 1u)) {
int2 writeIndex = (baseIndex + int2(int(c), int(r)));
uint2 writeIndex = (baseIndex + uint2(c, r));
if ((flip[0].x != 0u)) {
writeIndex = writeIndex.yx;
}

View File

@ -28,13 +28,13 @@ void main_inner(uint3 WorkGroupID, uint3 LocalInvocationID, uint local_invocatio
const uint filterOffset = ((params[0].x - 1u) / 2u);
int3 tint_tmp;
inputTex.GetDimensions(0, tint_tmp.x, tint_tmp.y, tint_tmp.z);
const int2 dims = tint_tmp.xy;
const int2 baseIndex = (int2(((WorkGroupID.xy * uint2(params[0].y, 4u)) + (LocalInvocationID.xy * uint2(4u, 1u)))) - int2(int(filterOffset), 0));
const uint2 dims = tint_tmp.xy;
const uint2 baseIndex = (((WorkGroupID.xy * uint2(params[0].y, 4u)) + (LocalInvocationID.xy * uint2(4u, 1u))) - uint2(filterOffset, 0u));
{
for(uint r = 0u; (r < 4u); r = (r + 1u)) {
{
for(uint c = 0u; (c < 4u); c = (c + 1u)) {
int2 loadIndex = (baseIndex + int2(int(c), int(r)));
uint2 loadIndex = (baseIndex + uint2(c, r));
if ((flip[0].x != 0u)) {
loadIndex = loadIndex.yx;
}
@ -48,7 +48,7 @@ void main_inner(uint3 WorkGroupID, uint3 LocalInvocationID, uint local_invocatio
for(uint r = 0u; (r < 4u); r = (r + 1u)) {
{
for(uint c = 0u; (c < 4u); c = (c + 1u)) {
int2 writeIndex = (baseIndex + int2(int(c), int(r)));
uint2 writeIndex = (baseIndex + uint2(c, r));
if ((flip[0].x != 0u)) {
writeIndex = writeIndex.yx;
}

View File

@ -29,13 +29,13 @@ void tint_symbol(uvec3 WorkGroupID, uvec3 LocalInvocationID, uint local_invocati
}
barrier();
uint filterOffset = ((params.filterDim - 1u) / 2u);
ivec2 dims = textureSize(inputTex_1, 0);
ivec2 baseIndex = (ivec2(((WorkGroupID.xy * uvec2(params.blockDim, 4u)) + (LocalInvocationID.xy * uvec2(4u, 1u)))) - ivec2(int(filterOffset), 0));
uvec2 dims = uvec2(textureSize(inputTex_1, 0));
uvec2 baseIndex = (((WorkGroupID.xy * uvec2(params.blockDim, 4u)) + (LocalInvocationID.xy * uvec2(4u, 1u))) - uvec2(filterOffset, 0u));
{
for(uint r = 0u; (r < 4u); r = (r + 1u)) {
{
for(uint c = 0u; (c < 4u); c = (c + 1u)) {
ivec2 loadIndex = (baseIndex + ivec2(int(c), int(r)));
uvec2 loadIndex = (baseIndex + uvec2(c, r));
if ((flip.value != 0u)) {
loadIndex = loadIndex.yx;
}
@ -49,7 +49,7 @@ void tint_symbol(uvec3 WorkGroupID, uvec3 LocalInvocationID, uint local_invocati
for(uint r = 0u; (r < 4u); r = (r + 1u)) {
{
for(uint c = 0u; (c < 4u); c = (c + 1u)) {
ivec2 writeIndex = (baseIndex + ivec2(int(c), int(r)));
uvec2 writeIndex = (baseIndex + uvec2(c, r));
if ((flip.value != 0u)) {
writeIndex = writeIndex.yx;
}
@ -70,7 +70,7 @@ void tint_symbol(uvec3 WorkGroupID, uvec3 LocalInvocationID, uint local_invocati
acc = (acc + ((1.0f / float(params.filterDim)) * tile[r][i]));
}
}
imageStore(outputTex, writeIndex, vec4(acc, 1.0f));
imageStore(outputTex, ivec2(writeIndex), vec4(acc, 1.0f));
}
}
}

View File

@ -31,13 +31,13 @@ void tint_symbol_inner(uint3 WorkGroupID, uint3 LocalInvocationID, uint local_in
}
threadgroup_barrier(mem_flags::mem_threadgroup);
uint const filterOffset = (((*(tint_symbol_2)).filterDim - 1u) / 2u);
int2 const dims = int2(tint_symbol_3.get_width(0), tint_symbol_3.get_height(0));
int2 const baseIndex = as_type<int2>((as_type<uint2>(int2(int2(((uint3(WorkGroupID).xy * uint2((*(tint_symbol_2)).blockDim, 4u)) + (uint3(LocalInvocationID).xy * uint2(4u, 1u)))))) - as_type<uint2>(int2(int2(int(filterOffset), 0)))));
uint2 const dims = uint2(tint_symbol_3.get_width(0), tint_symbol_3.get_height(0));
uint2 const baseIndex = (((uint3(WorkGroupID).xy * uint2((*(tint_symbol_2)).blockDim, 4u)) + (uint3(LocalInvocationID).xy * uint2(4u, 1u))) - uint2(filterOffset, 0u));
for(uint r = 0u; (r < 4u); r = (r + 1u)) {
for(uint c = 0u; (c < 4u); c = (c + 1u)) {
int2 loadIndex = as_type<int2>((as_type<uint2>(int2(baseIndex)) + as_type<uint2>(int2(int2(int(c), int(r))))));
uint2 loadIndex = (baseIndex + uint2(c, r));
if (((*(tint_symbol_4)).value != 0u)) {
loadIndex = int2(loadIndex).yx;
loadIndex = uint2(loadIndex).yx;
}
(*(tint_symbol_1))[r][((4u * LocalInvocationID[0]) + c)] = float4(tint_symbol_3.sample(tint_symbol_5, ((float2(loadIndex) + float2(0.25f)) / float2(dims)), level(0.0f))).rgb;
}
@ -45,9 +45,9 @@ void tint_symbol_inner(uint3 WorkGroupID, uint3 LocalInvocationID, uint local_in
threadgroup_barrier(mem_flags::mem_threadgroup);
for(uint r = 0u; (r < 4u); r = (r + 1u)) {
for(uint c = 0u; (c < 4u); c = (c + 1u)) {
int2 writeIndex = as_type<int2>((as_type<uint2>(int2(baseIndex)) + as_type<uint2>(int2(int2(int(c), int(r))))));
uint2 writeIndex = (baseIndex + uint2(c, r));
if (((*(tint_symbol_4)).value != 0u)) {
writeIndex = int2(writeIndex).yx;
writeIndex = uint2(writeIndex).yx;
}
uint const center = ((4u * LocalInvocationID[0]) + c);
if ((((center >= filterOffset) && (center < (256u - filterOffset))) && all((writeIndex < dims)))) {

View File

@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 258
; Bound: 251
; Schema: 0
OpCapability Shader
OpCapability ImageQuery
@ -104,23 +104,22 @@
%uint_0 = OpConstant %uint 0
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
%uint_1 = OpConstant %uint 1
%int = OpTypeInt 32 1
%v2int = OpTypeVector %int 2
%76 = OpConstantNull %int
%v2uint = OpTypeVector %uint 2
%85 = OpConstantComposite %v2uint %uint_4 %uint_1
%_ptr_Function_v2int = OpTypePointer Function %v2int
%119 = OpConstantNull %v2int
%int = OpTypeInt 32 1
%76 = OpConstantNull %int
%83 = OpConstantComposite %v2uint %uint_4 %uint_1
%_ptr_Function_v2uint = OpTypePointer Function %v2uint
%114 = OpConstantNull %v2uint
%v4float = OpTypeVector %float 4
%137 = OpTypeSampledImage %16
%132 = OpTypeSampledImage %16
%v2float = OpTypeVector %float 2
%float_0_25 = OpConstant %float 0.25
%143 = OpConstantComposite %v2float %float_0_25 %float_0_25
%147 = OpConstantNull %float
%138 = OpConstantComposite %v2float %float_0_25 %float_0_25
%142 = OpConstantNull %float
%v2bool = OpTypeVector %bool 2
%_ptr_Function_v3float = OpTypePointer Function %v3float
%float_1 = OpConstant %float 1
%251 = OpTypeFunction %void
%244 = OpTypeFunction %void
%main_inner = OpFunction %void None %31
%WorkGroupID = OpFunctionParameter %v3uint
%LocalInvocationID = OpFunctionParameter %v3uint
@ -129,10 +128,10 @@
%idx = OpVariable %_ptr_Function_uint Function %40
%r = OpVariable %_ptr_Function_uint Function %40
%c = OpVariable %_ptr_Function_uint Function %40
%loadIndex = OpVariable %_ptr_Function_v2int Function %119
%loadIndex = OpVariable %_ptr_Function_v2uint Function %114
%r_0 = OpVariable %_ptr_Function_uint Function %40
%c_0 = OpVariable %_ptr_Function_uint Function %40
%writeIndex = OpVariable %_ptr_Function_v2int Function %119
%writeIndex = OpVariable %_ptr_Function_v2uint Function %114
%acc = OpVariable %_ptr_Function_v3float Function %58
%f = OpVariable %_ptr_Function_uint Function %40
%i = OpVariable %_ptr_Function_uint Function %40
@ -168,240 +167,234 @@
%68 = OpLoad %uint %67
%70 = OpISub %uint %68 %uint_1
%71 = OpUDiv %uint %70 %uint_2
%75 = OpLoad %16 %inputTex
%72 = OpImageQuerySizeLod %v2int %75 %76
%79 = OpVectorShuffle %v2uint %WorkGroupID %WorkGroupID 0 1
%80 = OpAccessChain %_ptr_Uniform_uint %params %uint_1
%81 = OpLoad %uint %80
%82 = OpCompositeConstruct %v2uint %81 %uint_4
%83 = OpIMul %v2uint %79 %82
%84 = OpVectorShuffle %v2uint %LocalInvocationID %LocalInvocationID 0 1
%86 = OpIMul %v2uint %84 %85
%87 = OpIAdd %v2uint %83 %86
%77 = OpBitcast %v2int %87
%88 = OpBitcast %int %71
%89 = OpCompositeConstruct %v2int %88 %76
%90 = OpISub %v2int %77 %89
%74 = OpLoad %16 %inputTex
%72 = OpImageQuerySizeLod %v2uint %74 %76
%77 = OpVectorShuffle %v2uint %WorkGroupID %WorkGroupID 0 1
%78 = OpAccessChain %_ptr_Uniform_uint %params %uint_1
%79 = OpLoad %uint %78
%80 = OpCompositeConstruct %v2uint %79 %uint_4
%81 = OpIMul %v2uint %77 %80
%82 = OpVectorShuffle %v2uint %LocalInvocationID %LocalInvocationID 0 1
%84 = OpIMul %v2uint %82 %83
%85 = OpIAdd %v2uint %81 %84
%86 = OpCompositeConstruct %v2uint %71 %40
%87 = OpISub %v2uint %85 %86
OpStore %r %40
OpBranch %89
%89 = OpLabel
OpLoopMerge %90 %91 None
OpBranch %92
%92 = OpLabel
OpLoopMerge %93 %94 None
OpBranch %95
%95 = OpLabel
%97 = OpLoad %uint %r
%98 = OpULessThan %bool %97 %uint_4
%96 = OpLogicalNot %bool %98
OpSelectionMerge %99 None
OpBranchConditional %96 %100 %99
%100 = OpLabel
OpBranch %93
%99 = OpLabel
%94 = OpLoad %uint %r
%95 = OpULessThan %bool %94 %uint_4
%93 = OpLogicalNot %bool %95
OpSelectionMerge %96 None
OpBranchConditional %93 %97 %96
%97 = OpLabel
OpBranch %90
%96 = OpLabel
OpStore %c %40
OpBranch %99
%99 = OpLabel
OpLoopMerge %100 %101 None
OpBranch %102
%102 = OpLabel
OpLoopMerge %103 %104 None
OpBranch %105
%105 = OpLabel
%107 = OpLoad %uint %c
%108 = OpULessThan %bool %107 %uint_4
%106 = OpLogicalNot %bool %108
OpSelectionMerge %109 None
OpBranchConditional %106 %110 %109
%110 = OpLabel
OpBranch %103
%109 = OpLabel
%112 = OpLoad %uint %c
%111 = OpBitcast %int %112
%114 = OpLoad %uint %r
%113 = OpBitcast %int %114
%115 = OpCompositeConstruct %v2int %111 %113
%116 = OpIAdd %v2int %90 %115
OpStore %loadIndex %116
%120 = OpAccessChain %_ptr_Uniform_uint %flip %uint_0
%121 = OpLoad %uint %120
%122 = OpINotEqual %bool %121 %40
OpSelectionMerge %123 None
OpBranchConditional %122 %124 %123
%124 = OpLabel
%125 = OpLoad %v2int %loadIndex
%126 = OpVectorShuffle %v2int %125 %125 1 0
OpStore %loadIndex %126
OpBranch %123
%123 = OpLabel
%127 = OpLoad %uint %r
%128 = OpCompositeExtract %uint %LocalInvocationID 0
%129 = OpIMul %uint %uint_4 %128
%130 = OpLoad %uint %c
%131 = OpIAdd %uint %129 %130
%132 = OpAccessChain %_ptr_Workgroup_v3float %tile %127 %131
%135 = OpLoad %10 %samp
%136 = OpLoad %16 %inputTex
%138 = OpSampledImage %137 %136 %135
%141 = OpLoad %v2int %loadIndex
%139 = OpConvertSToF %v2float %141
%144 = OpFAdd %v2float %139 %143
%145 = OpConvertSToF %v2float %72
%146 = OpFDiv %v2float %144 %145
%133 = OpImageSampleExplicitLod %v4float %138 %146 Lod %147
%148 = OpVectorShuffle %v3float %133 %133 0 1 2
OpStore %132 %148
OpBranch %104
%104 = OpLabel
%149 = OpLoad %uint %c
%150 = OpIAdd %uint %149 %uint_1
OpStore %c %150
OpBranch %102
%103 = OpLabel
OpBranch %94
%94 = OpLabel
%151 = OpLoad %uint %r
%152 = OpIAdd %uint %151 %uint_1
OpStore %r %152
OpBranch %92
%93 = OpLabel
%104 = OpLoad %uint %c
%105 = OpULessThan %bool %104 %uint_4
%103 = OpLogicalNot %bool %105
OpSelectionMerge %106 None
OpBranchConditional %103 %107 %106
%107 = OpLabel
OpBranch %100
%106 = OpLabel
%108 = OpLoad %uint %c
%109 = OpLoad %uint %r
%110 = OpCompositeConstruct %v2uint %108 %109
%111 = OpIAdd %v2uint %87 %110
OpStore %loadIndex %111
%115 = OpAccessChain %_ptr_Uniform_uint %flip %uint_0
%116 = OpLoad %uint %115
%117 = OpINotEqual %bool %116 %40
OpSelectionMerge %118 None
OpBranchConditional %117 %119 %118
%119 = OpLabel
%120 = OpLoad %v2uint %loadIndex
%121 = OpVectorShuffle %v2uint %120 %120 1 0
OpStore %loadIndex %121
OpBranch %118
%118 = OpLabel
%122 = OpLoad %uint %r
%123 = OpCompositeExtract %uint %LocalInvocationID 0
%124 = OpIMul %uint %uint_4 %123
%125 = OpLoad %uint %c
%126 = OpIAdd %uint %124 %125
%127 = OpAccessChain %_ptr_Workgroup_v3float %tile %122 %126
%130 = OpLoad %10 %samp
%131 = OpLoad %16 %inputTex
%133 = OpSampledImage %132 %131 %130
%136 = OpLoad %v2uint %loadIndex
%134 = OpConvertUToF %v2float %136
%139 = OpFAdd %v2float %134 %138
%140 = OpConvertUToF %v2float %72
%141 = OpFDiv %v2float %139 %140
%128 = OpImageSampleExplicitLod %v4float %133 %141 Lod %142
%143 = OpVectorShuffle %v3float %128 %128 0 1 2
OpStore %127 %143
OpBranch %101
%101 = OpLabel
%144 = OpLoad %uint %c
%145 = OpIAdd %uint %144 %uint_1
OpStore %c %145
OpBranch %99
%100 = OpLabel
OpBranch %91
%91 = OpLabel
%146 = OpLoad %uint %r
%147 = OpIAdd %uint %146 %uint_1
OpStore %r %147
OpBranch %89
%90 = OpLabel
OpControlBarrier %uint_2 %uint_2 %uint_264
OpStore %r_0 %40
OpBranch %155
%155 = OpLabel
OpLoopMerge %156 %157 None
OpBranch %158
OpBranch %150
%150 = OpLabel
OpLoopMerge %151 %152 None
OpBranch %153
%153 = OpLabel
%155 = OpLoad %uint %r_0
%156 = OpULessThan %bool %155 %uint_4
%154 = OpLogicalNot %bool %156
OpSelectionMerge %157 None
OpBranchConditional %154 %158 %157
%158 = OpLabel
%160 = OpLoad %uint %r_0
%161 = OpULessThan %bool %160 %uint_4
%159 = OpLogicalNot %bool %161
OpSelectionMerge %162 None
OpBranchConditional %159 %163 %162
%163 = OpLabel
OpBranch %156
%162 = OpLabel
OpBranch %151
%157 = OpLabel
OpStore %c_0 %40
OpBranch %165
%165 = OpLabel
OpLoopMerge %166 %167 None
OpBranch %168
OpBranch %160
%160 = OpLabel
OpLoopMerge %161 %162 None
OpBranch %163
%163 = OpLabel
%165 = OpLoad %uint %c_0
%166 = OpULessThan %bool %165 %uint_4
%164 = OpLogicalNot %bool %166
OpSelectionMerge %167 None
OpBranchConditional %164 %168 %167
%168 = OpLabel
%170 = OpLoad %uint %c_0
%171 = OpULessThan %bool %170 %uint_4
%169 = OpLogicalNot %bool %171
OpSelectionMerge %172 None
OpBranchConditional %169 %173 %172
%173 = OpLabel
OpBranch %166
%172 = OpLabel
%175 = OpLoad %uint %c_0
%174 = OpBitcast %int %175
%177 = OpLoad %uint %r_0
%176 = OpBitcast %int %177
%178 = OpCompositeConstruct %v2int %174 %176
%179 = OpIAdd %v2int %90 %178
OpStore %writeIndex %179
%181 = OpAccessChain %_ptr_Uniform_uint %flip %uint_0
%182 = OpLoad %uint %181
%183 = OpINotEqual %bool %182 %40
OpSelectionMerge %184 None
OpBranchConditional %183 %185 %184
%185 = OpLabel
%186 = OpLoad %v2int %writeIndex
%187 = OpVectorShuffle %v2int %186 %186 1 0
OpStore %writeIndex %187
OpBranch %184
%184 = OpLabel
%188 = OpCompositeExtract %uint %LocalInvocationID 0
%189 = OpIMul %uint %uint_4 %188
%190 = OpLoad %uint %c_0
%191 = OpIAdd %uint %189 %190
%192 = OpUGreaterThanEqual %bool %191 %71
OpSelectionMerge %193 None
OpBranchConditional %192 %194 %193
%194 = OpLabel
%195 = OpISub %uint %uint_256 %71
%196 = OpULessThan %bool %191 %195
OpBranch %193
%193 = OpLabel
%197 = OpPhi %bool %192 %184 %196 %194
OpBranch %161
%167 = OpLabel
%169 = OpLoad %uint %c_0
%170 = OpLoad %uint %r_0
%171 = OpCompositeConstruct %v2uint %169 %170
%172 = OpIAdd %v2uint %87 %171
OpStore %writeIndex %172
%174 = OpAccessChain %_ptr_Uniform_uint %flip %uint_0
%175 = OpLoad %uint %174
%176 = OpINotEqual %bool %175 %40
OpSelectionMerge %177 None
OpBranchConditional %176 %178 %177
%178 = OpLabel
%179 = OpLoad %v2uint %writeIndex
%180 = OpVectorShuffle %v2uint %179 %179 1 0
OpStore %writeIndex %180
OpBranch %177
%177 = OpLabel
%181 = OpCompositeExtract %uint %LocalInvocationID 0
%182 = OpIMul %uint %uint_4 %181
%183 = OpLoad %uint %c_0
%184 = OpIAdd %uint %182 %183
%185 = OpUGreaterThanEqual %bool %184 %71
OpSelectionMerge %186 None
OpBranchConditional %185 %187 %186
%187 = OpLabel
%188 = OpISub %uint %uint_256 %71
%189 = OpULessThan %bool %184 %188
OpBranch %186
%186 = OpLabel
%190 = OpPhi %bool %185 %177 %189 %187
OpSelectionMerge %191 None
OpBranchConditional %190 %192 %191
%192 = OpLabel
%194 = OpLoad %v2uint %writeIndex
%195 = OpULessThan %v2bool %194 %72
%193 = OpAll %bool %195
OpBranch %191
%191 = OpLabel
%197 = OpPhi %bool %190 %186 %193 %192
OpSelectionMerge %198 None
OpBranchConditional %197 %199 %198
%199 = OpLabel
%201 = OpLoad %v2int %writeIndex
%202 = OpSLessThan %v2bool %201 %72
%200 = OpAll %bool %202
OpBranch %198
%198 = OpLabel
%204 = OpPhi %bool %197 %193 %200 %199
OpSelectionMerge %205 None
OpBranchConditional %204 %206 %205
%206 = OpLabel
OpStore %acc %58
OpStore %f %40
OpBranch %210
%210 = OpLabel
OpLoopMerge %211 %212 None
OpBranch %213
OpBranch %203
%203 = OpLabel
OpLoopMerge %204 %205 None
OpBranch %206
%206 = OpLabel
%208 = OpLoad %uint %f
%209 = OpAccessChain %_ptr_Uniform_uint %params %uint_0
%210 = OpLoad %uint %209
%211 = OpULessThan %bool %208 %210
%207 = OpLogicalNot %bool %211
OpSelectionMerge %212 None
OpBranchConditional %207 %213 %212
%213 = OpLabel
%215 = OpLoad %uint %f
%216 = OpAccessChain %_ptr_Uniform_uint %params %uint_0
%217 = OpLoad %uint %216
%218 = OpULessThan %bool %215 %217
%214 = OpLogicalNot %bool %218
OpSelectionMerge %219 None
OpBranchConditional %214 %220 %219
%220 = OpLabel
OpBranch %211
%219 = OpLabel
%221 = OpLoad %uint %f
%222 = OpIAdd %uint %191 %221
%223 = OpISub %uint %222 %71
OpStore %i %223
%225 = OpLoad %v3float %acc
%228 = OpAccessChain %_ptr_Uniform_uint %params %uint_0
%229 = OpLoad %uint %228
%227 = OpConvertUToF %float %229
%230 = OpFDiv %float %float_1 %227
%231 = OpLoad %uint %r_0
%232 = OpLoad %uint %i
%233 = OpAccessChain %_ptr_Workgroup_v3float %tile %231 %232
%234 = OpLoad %v3float %233
%235 = OpVectorTimesScalar %v3float %234 %230
%236 = OpFAdd %v3float %225 %235
OpStore %acc %236
OpBranch %212
OpBranch %204
%212 = OpLabel
%237 = OpLoad %uint %f
%238 = OpIAdd %uint %237 %uint_1
OpStore %f %238
OpBranch %210
%211 = OpLabel
%240 = OpLoad %20 %outputTex
%241 = OpLoad %v2int %writeIndex
%242 = OpLoad %v3float %acc
%243 = OpCompositeExtract %float %242 0
%244 = OpCompositeExtract %float %242 1
%245 = OpCompositeExtract %float %242 2
%246 = OpCompositeConstruct %v4float %243 %244 %245 %float_1
OpImageWrite %240 %241 %246
%214 = OpLoad %uint %f
%215 = OpIAdd %uint %184 %214
%216 = OpISub %uint %215 %71
OpStore %i %216
%218 = OpLoad %v3float %acc
%221 = OpAccessChain %_ptr_Uniform_uint %params %uint_0
%222 = OpLoad %uint %221
%220 = OpConvertUToF %float %222
%223 = OpFDiv %float %float_1 %220
%224 = OpLoad %uint %r_0
%225 = OpLoad %uint %i
%226 = OpAccessChain %_ptr_Workgroup_v3float %tile %224 %225
%227 = OpLoad %v3float %226
%228 = OpVectorTimesScalar %v3float %227 %223
%229 = OpFAdd %v3float %218 %228
OpStore %acc %229
OpBranch %205
%205 = OpLabel
OpBranch %167
%167 = OpLabel
%247 = OpLoad %uint %c_0
%248 = OpIAdd %uint %247 %uint_1
OpStore %c_0 %248
OpBranch %165
%166 = OpLabel
OpBranch %157
%157 = OpLabel
%249 = OpLoad %uint %r_0
%250 = OpIAdd %uint %249 %uint_1
OpStore %r_0 %250
OpBranch %155
%156 = OpLabel
%230 = OpLoad %uint %f
%231 = OpIAdd %uint %230 %uint_1
OpStore %f %231
OpBranch %203
%204 = OpLabel
%233 = OpLoad %20 %outputTex
%234 = OpLoad %v2uint %writeIndex
%235 = OpLoad %v3float %acc
%236 = OpCompositeExtract %float %235 0
%237 = OpCompositeExtract %float %235 1
%238 = OpCompositeExtract %float %235 2
%239 = OpCompositeConstruct %v4float %236 %237 %238 %float_1
OpImageWrite %233 %234 %239
OpBranch %198
%198 = OpLabel
OpBranch %162
%162 = OpLabel
%240 = OpLoad %uint %c_0
%241 = OpIAdd %uint %240 %uint_1
OpStore %c_0 %241
OpBranch %160
%161 = OpLabel
OpBranch %152
%152 = OpLabel
%242 = OpLoad %uint %r_0
%243 = OpIAdd %uint %242 %uint_1
OpStore %r_0 %243
OpBranch %150
%151 = OpLabel
OpReturn
OpFunctionEnd
%main = OpFunction %void None %251
%253 = OpLabel
%255 = OpLoad %v3uint %WorkGroupID_1
%256 = OpLoad %v3uint %LocalInvocationID_1
%257 = OpLoad %uint %local_invocation_index_1
%254 = OpFunctionCall %void %main_inner %255 %256 %257
%main = OpFunction %void None %244
%246 = OpLabel
%248 = OpLoad %v3uint %WorkGroupID_1
%249 = OpLoad %v3uint %LocalInvocationID_1
%250 = OpLoad %uint %local_invocation_index_1
%247 = OpFunctionCall %void %main_inner %248 %249 %250
OpReturn
OpFunctionEnd

View File

@ -22,11 +22,11 @@ var<workgroup> tile : array<array<vec3<f32>, 256>, 4>;
@compute @workgroup_size(64, 1, 1)
fn main(@builtin(workgroup_id) WorkGroupID : vec3<u32>, @builtin(local_invocation_id) LocalInvocationID : vec3<u32>) {
let filterOffset : u32 = ((params.filterDim - 1u) / 2u);
let dims : vec2<i32> = textureDimensions(inputTex, 0);
let baseIndex = (vec2<i32>(((WorkGroupID.xy * vec2<u32>(params.blockDim, 4u)) + (LocalInvocationID.xy * vec2<u32>(4u, 1u)))) - vec2<i32>(i32(filterOffset), 0));
let dims = textureDimensions(inputTex, 0);
let baseIndex = (((WorkGroupID.xy * vec2(params.blockDim, 4)) + (LocalInvocationID.xy * vec2(4u, 1u))) - vec2(filterOffset, 0));
for(var r : u32 = 0u; (r < 4u); r = (r + 1u)) {
for(var c : u32 = 0u; (c < 4u); c = (c + 1u)) {
var loadIndex = (baseIndex + vec2<i32>(i32(c), i32(r)));
var loadIndex = (baseIndex + vec2(c, r));
if ((flip.value != 0u)) {
loadIndex = loadIndex.yx;
}
@ -36,7 +36,7 @@ fn main(@builtin(workgroup_id) WorkGroupID : vec3<u32>, @builtin(local_invocatio
workgroupBarrier();
for(var r : u32 = 0u; (r < 4u); r = (r + 1u)) {
for(var c : u32 = 0u; (c < 4u); c = (c + 1u)) {
var writeIndex = (baseIndex + vec2<i32>(i32(c), i32(r)));
var writeIndex = (baseIndex + vec2(c, r));
if ((flip.value != 0u)) {
writeIndex = writeIndex.yx;
}

View File

@ -1,44 +0,0 @@
// Copyright 2021 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
////////////////////////////////////////////////////////////////////////////////
// File generated by tools/src/cmd/gen
// using the template:
// test/tint/builtins/gen/gen.wgsl.tmpl
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
@group(1) @binding(0) var arg_0: texture_1d<f32>;
// fn textureDimensions(texture: texture_1d<f32>) -> i32
fn textureDimensions_002b2a() {
var res: i32 = textureDimensions(arg_0);
}
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
textureDimensions_002b2a();
return vec4<f32>();
}
@fragment
fn fragment_main() {
textureDimensions_002b2a();
}
@compute @workgroup_size(1)
fn compute_main() {
textureDimensions_002b2a();
}

View File

@ -1,34 +0,0 @@
Texture1D<float4> arg_0 : register(t0, space1);
void textureDimensions_002b2a() {
int tint_tmp;
arg_0.GetDimensions(tint_tmp);
int res = tint_tmp;
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
textureDimensions_002b2a();
return (0.0f).xxxx;
}
tint_symbol vertex_main() {
const float4 inner_result = vertex_main_inner();
tint_symbol wrapper_result = (tint_symbol)0;
wrapper_result.value = inner_result;
return wrapper_result;
}
void fragment_main() {
textureDimensions_002b2a();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
textureDimensions_002b2a();
return;
}

View File

@ -1,34 +0,0 @@
Texture1D<float4> arg_0 : register(t0, space1);
void textureDimensions_002b2a() {
int tint_tmp;
arg_0.GetDimensions(tint_tmp);
int res = tint_tmp;
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
textureDimensions_002b2a();
return (0.0f).xxxx;
}
tint_symbol vertex_main() {
const float4 inner_result = vertex_main_inner();
tint_symbol wrapper_result = (tint_symbol)0;
wrapper_result.value = inner_result;
return wrapper_result;
}
void fragment_main() {
textureDimensions_002b2a();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
textureDimensions_002b2a();
return;
}

View File

@ -1,75 +0,0 @@
SKIP: FAILED
#version 310 es
uniform highp sampler1D arg_0_1;
void textureDimensions_002b2a() {
int res = textureSize(arg_0_1, 0);
}
vec4 vertex_main() {
textureDimensions_002b2a();
return vec4(0.0f);
}
void main() {
gl_PointSize = 1.0;
vec4 inner_result = vertex_main();
gl_Position = inner_result;
gl_Position.y = -(gl_Position.y);
gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
return;
}
Error parsing GLSL shader:
ERROR: 0:3: 'sampler1D' : Reserved word.
ERROR: 0:3: '' : compilation terminated
ERROR: 2 compilation errors. No code generated.
#version 310 es
precision mediump float;
uniform highp sampler1D arg_0_1;
void textureDimensions_002b2a() {
int res = textureSize(arg_0_1, 0);
}
void fragment_main() {
textureDimensions_002b2a();
}
void main() {
fragment_main();
return;
}
Error parsing GLSL shader:
ERROR: 0:4: 'sampler1D' : Reserved word.
ERROR: 0:4: '' : compilation terminated
ERROR: 2 compilation errors. No code generated.
#version 310 es
uniform highp sampler1D arg_0_1;
void textureDimensions_002b2a() {
int res = textureSize(arg_0_1, 0);
}
void compute_main() {
textureDimensions_002b2a();
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
compute_main();
return;
}
Error parsing GLSL shader:
ERROR: 0:3: 'sampler1D' : Reserved word.
ERROR: 0:3: '' : compilation terminated
ERROR: 2 compilation errors. No code generated.

View File

@ -1,33 +0,0 @@
#include <metal_stdlib>
using namespace metal;
void textureDimensions_002b2a(texture1d<float, access::sample> tint_symbol_1) {
int res = int(tint_symbol_1.get_width(0));
}
struct tint_symbol {
float4 value [[position]];
};
float4 vertex_main_inner(texture1d<float, access::sample> tint_symbol_2) {
textureDimensions_002b2a(tint_symbol_2);
return float4(0.0f);
}
vertex tint_symbol vertex_main(texture1d<float, access::sample> tint_symbol_3 [[texture(0)]]) {
float4 const inner_result = vertex_main_inner(tint_symbol_3);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
fragment void fragment_main(texture1d<float, access::sample> tint_symbol_4 [[texture(0)]]) {
textureDimensions_002b2a(tint_symbol_4);
return;
}
kernel void compute_main(texture1d<float, access::sample> tint_symbol_5 [[texture(0)]]) {
textureDimensions_002b2a(tint_symbol_5);
return;
}

View File

@ -1,76 +0,0 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 37
; Schema: 0
OpCapability Shader
OpCapability Sampled1D
OpCapability ImageQuery
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
OpEntryPoint GLCompute %compute_main "compute_main"
OpExecutionMode %fragment_main OriginUpperLeft
OpExecutionMode %compute_main LocalSize 1 1 1
OpName %value "value"
OpName %vertex_point_size "vertex_point_size"
OpName %arg_0 "arg_0"
OpName %textureDimensions_002b2a "textureDimensions_002b2a"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
OpName %vertex_main "vertex_main"
OpName %fragment_main "fragment_main"
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
OpDecorate %arg_0 DescriptorSet 1
OpDecorate %arg_0 Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
%5 = OpConstantNull %v4float
%value = OpVariable %_ptr_Output_v4float Output %5
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
%11 = OpTypeImage %float 1D 0 0 0 1 Unknown
%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
%arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant
%void = OpTypeVoid
%12 = OpTypeFunction %void
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%_ptr_Function_int = OpTypePointer Function %int
%22 = OpConstantNull %int
%23 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%textureDimensions_002b2a = OpFunction %void None %12
%15 = OpLabel
%res = OpVariable %_ptr_Function_int Function %22
%18 = OpLoad %11 %arg_0
%16 = OpImageQuerySizeLod %int %18 %int_0
OpStore %res %16
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %23
%25 = OpLabel
%26 = OpFunctionCall %void %textureDimensions_002b2a
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %12
%28 = OpLabel
%29 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %29
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %12
%32 = OpLabel
%33 = OpFunctionCall %void %textureDimensions_002b2a
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %12
%35 = OpLabel
%36 = OpFunctionCall %void %textureDimensions_002b2a
OpReturn
OpFunctionEnd

View File

@ -1,21 +0,0 @@
@group(1) @binding(0) var arg_0 : texture_1d<f32>;
fn textureDimensions_002b2a() {
var res : i32 = textureDimensions(arg_0);
}
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
textureDimensions_002b2a();
return vec4<f32>();
}
@fragment
fn fragment_main() {
textureDimensions_002b2a();
}
@compute @workgroup_size(1)
fn compute_main() {
textureDimensions_002b2a();
}

View File

@ -0,0 +1,44 @@
// Copyright 2022 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
////////////////////////////////////////////////////////////////////////////////
// File generated by tools/src/cmd/gen
// using the template:
// test/tint/builtins/gen/gen.wgsl.tmpl
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
@group(1) @binding(0) var arg_0: texture_multisampled_2d<i32>;
// fn textureDimensions(texture: texture_multisampled_2d<i32>) -> vec2<u32>
fn textureDimensions_00348c() {
var res: vec2<u32> = textureDimensions(arg_0);
}
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
textureDimensions_00348c();
return vec4<f32>();
}
@fragment
fn fragment_main() {
textureDimensions_00348c();
}
@compute @workgroup_size(1)
fn compute_main() {
textureDimensions_00348c();
}

View File

@ -0,0 +1,34 @@
Texture2DMS<int4> arg_0 : register(t0, space1);
void textureDimensions_00348c() {
int3 tint_tmp;
arg_0.GetDimensions(tint_tmp.x, tint_tmp.y, tint_tmp.z);
uint2 res = tint_tmp.xy;
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
textureDimensions_00348c();
return (0.0f).xxxx;
}
tint_symbol vertex_main() {
const float4 inner_result = vertex_main_inner();
tint_symbol wrapper_result = (tint_symbol)0;
wrapper_result.value = inner_result;
return wrapper_result;
}
void fragment_main() {
textureDimensions_00348c();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
textureDimensions_00348c();
return;
}

View File

@ -0,0 +1,34 @@
Texture2DMS<int4> arg_0 : register(t0, space1);
void textureDimensions_00348c() {
int3 tint_tmp;
arg_0.GetDimensions(tint_tmp.x, tint_tmp.y, tint_tmp.z);
uint2 res = tint_tmp.xy;
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
textureDimensions_00348c();
return (0.0f).xxxx;
}
tint_symbol vertex_main() {
const float4 inner_result = vertex_main_inner();
tint_symbol wrapper_result = (tint_symbol)0;
wrapper_result.value = inner_result;
return wrapper_result;
}
void fragment_main() {
textureDimensions_00348c();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
textureDimensions_00348c();
return;
}

View File

@ -0,0 +1,52 @@
#version 310 es
uniform highp isampler2DMS arg_0_1;
void textureDimensions_00348c() {
uvec2 res = uvec2(textureSize(arg_0_1));
}
vec4 vertex_main() {
textureDimensions_00348c();
return vec4(0.0f);
}
void main() {
gl_PointSize = 1.0;
vec4 inner_result = vertex_main();
gl_Position = inner_result;
gl_Position.y = -(gl_Position.y);
gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
return;
}
#version 310 es
precision mediump float;
uniform highp isampler2DMS arg_0_1;
void textureDimensions_00348c() {
uvec2 res = uvec2(textureSize(arg_0_1));
}
void fragment_main() {
textureDimensions_00348c();
}
void main() {
fragment_main();
return;
}
#version 310 es
uniform highp isampler2DMS arg_0_1;
void textureDimensions_00348c() {
uvec2 res = uvec2(textureSize(arg_0_1));
}
void compute_main() {
textureDimensions_00348c();
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
compute_main();
return;
}

View File

@ -0,0 +1,33 @@
#include <metal_stdlib>
using namespace metal;
void textureDimensions_00348c(texture2d_ms<int, access::read> tint_symbol_1) {
uint2 res = uint2(tint_symbol_1.get_width(), tint_symbol_1.get_height());
}
struct tint_symbol {
float4 value [[position]];
};
float4 vertex_main_inner(texture2d_ms<int, access::read> tint_symbol_2) {
textureDimensions_00348c(tint_symbol_2);
return float4(0.0f);
}
vertex tint_symbol vertex_main(texture2d_ms<int, access::read> tint_symbol_3 [[texture(0)]]) {
float4 const inner_result = vertex_main_inner(tint_symbol_3);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
fragment void fragment_main(texture2d_ms<int, access::read> tint_symbol_4 [[texture(0)]]) {
textureDimensions_00348c(tint_symbol_4);
return;
}
kernel void compute_main(texture2d_ms<int, access::read> tint_symbol_5 [[texture(0)]]) {
textureDimensions_00348c(tint_symbol_5);
return;
}

View File

@ -0,0 +1,76 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 38
; Schema: 0
OpCapability Shader
OpCapability ImageQuery
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
OpEntryPoint GLCompute %compute_main "compute_main"
OpExecutionMode %fragment_main OriginUpperLeft
OpExecutionMode %compute_main LocalSize 1 1 1
OpName %value "value"
OpName %vertex_point_size "vertex_point_size"
OpName %arg_0 "arg_0"
OpName %textureDimensions_00348c "textureDimensions_00348c"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
OpName %vertex_main "vertex_main"
OpName %fragment_main "fragment_main"
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
OpDecorate %arg_0 DescriptorSet 1
OpDecorate %arg_0 Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
%5 = OpConstantNull %v4float
%value = OpVariable %_ptr_Output_v4float Output %5
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
%int = OpTypeInt 32 1
%11 = OpTypeImage %int 2D 0 0 1 1 Unknown
%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
%arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant
%void = OpTypeVoid
%13 = OpTypeFunction %void
%uint = OpTypeInt 32 0
%v2uint = OpTypeVector %uint 2
%_ptr_Function_v2uint = OpTypePointer Function %v2uint
%23 = OpConstantNull %v2uint
%24 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%textureDimensions_00348c = OpFunction %void None %13
%16 = OpLabel
%res = OpVariable %_ptr_Function_v2uint Function %23
%20 = OpLoad %11 %arg_0
%17 = OpImageQuerySize %v2uint %20
OpStore %res %17
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %24
%26 = OpLabel
%27 = OpFunctionCall %void %textureDimensions_00348c
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %13
%29 = OpLabel
%30 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %30
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %13
%33 = OpLabel
%34 = OpFunctionCall %void %textureDimensions_00348c
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %13
%36 = OpLabel
%37 = OpFunctionCall %void %textureDimensions_00348c
OpReturn
OpFunctionEnd

View File

@ -0,0 +1,21 @@
@group(1) @binding(0) var arg_0 : texture_multisampled_2d<i32>;
fn textureDimensions_00348c() {
var res : vec2<u32> = textureDimensions(arg_0);
}
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
textureDimensions_00348c();
return vec4<f32>();
}
@fragment
fn fragment_main() {
textureDimensions_00348c();
}
@compute @workgroup_size(1)
fn compute_main() {
textureDimensions_00348c();
}

View File

@ -1,44 +0,0 @@
// Copyright 2021 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
////////////////////////////////////////////////////////////////////////////////
// File generated by tools/src/cmd/gen
// using the template:
// test/tint/builtins/gen/gen.wgsl.tmpl
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
@group(1) @binding(0) var arg_0: texture_storage_2d_array<r32sint, write>;
// fn textureDimensions(texture: texture_storage_2d_array<r32sint, write>) -> vec2<i32>
fn textureDimensions_012b82() {
var res: vec2<i32> = textureDimensions(arg_0);
}
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
textureDimensions_012b82();
return vec4<f32>();
}
@fragment
fn fragment_main() {
textureDimensions_012b82();
}
@compute @workgroup_size(1)
fn compute_main() {
textureDimensions_012b82();
}

View File

@ -1,34 +0,0 @@
RWTexture2DArray<int4> arg_0 : register(u0, space1);
void textureDimensions_012b82() {
int3 tint_tmp;
arg_0.GetDimensions(tint_tmp.x, tint_tmp.y, tint_tmp.z);
int2 res = tint_tmp.xy;
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
textureDimensions_012b82();
return (0.0f).xxxx;
}
tint_symbol vertex_main() {
const float4 inner_result = vertex_main_inner();
tint_symbol wrapper_result = (tint_symbol)0;
wrapper_result.value = inner_result;
return wrapper_result;
}
void fragment_main() {
textureDimensions_012b82();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
textureDimensions_012b82();
return;
}

View File

@ -1,34 +0,0 @@
RWTexture2DArray<int4> arg_0 : register(u0, space1);
void textureDimensions_012b82() {
int3 tint_tmp;
arg_0.GetDimensions(tint_tmp.x, tint_tmp.y, tint_tmp.z);
int2 res = tint_tmp.xy;
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
textureDimensions_012b82();
return (0.0f).xxxx;
}
tint_symbol vertex_main() {
const float4 inner_result = vertex_main_inner();
tint_symbol wrapper_result = (tint_symbol)0;
wrapper_result.value = inner_result;
return wrapper_result;
}
void fragment_main() {
textureDimensions_012b82();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
textureDimensions_012b82();
return;
}

View File

@ -1,52 +0,0 @@
#version 310 es
layout(r32i) uniform highp writeonly iimage2DArray arg_0;
void textureDimensions_012b82() {
ivec2 res = imageSize(arg_0).xy;
}
vec4 vertex_main() {
textureDimensions_012b82();
return vec4(0.0f);
}
void main() {
gl_PointSize = 1.0;
vec4 inner_result = vertex_main();
gl_Position = inner_result;
gl_Position.y = -(gl_Position.y);
gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
return;
}
#version 310 es
precision mediump float;
layout(r32i) uniform highp writeonly iimage2DArray arg_0;
void textureDimensions_012b82() {
ivec2 res = imageSize(arg_0).xy;
}
void fragment_main() {
textureDimensions_012b82();
}
void main() {
fragment_main();
return;
}
#version 310 es
layout(r32i) uniform highp writeonly iimage2DArray arg_0;
void textureDimensions_012b82() {
ivec2 res = imageSize(arg_0).xy;
}
void compute_main() {
textureDimensions_012b82();
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
compute_main();
return;
}

View File

@ -1,33 +0,0 @@
#include <metal_stdlib>
using namespace metal;
void textureDimensions_012b82(texture2d_array<int, access::write> tint_symbol_1) {
int2 res = int2(tint_symbol_1.get_width(), tint_symbol_1.get_height());
}
struct tint_symbol {
float4 value [[position]];
};
float4 vertex_main_inner(texture2d_array<int, access::write> tint_symbol_2) {
textureDimensions_012b82(tint_symbol_2);
return float4(0.0f);
}
vertex tint_symbol vertex_main(texture2d_array<int, access::write> tint_symbol_3 [[texture(0)]]) {
float4 const inner_result = vertex_main_inner(tint_symbol_3);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
fragment void fragment_main(texture2d_array<int, access::write> tint_symbol_4 [[texture(0)]]) {
textureDimensions_012b82(tint_symbol_4);
return;
}
kernel void compute_main(texture2d_array<int, access::write> tint_symbol_5 [[texture(0)]]) {
textureDimensions_012b82(tint_symbol_5);
return;
}

View File

@ -1,78 +0,0 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 39
; Schema: 0
OpCapability Shader
OpCapability ImageQuery
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
OpEntryPoint GLCompute %compute_main "compute_main"
OpExecutionMode %fragment_main OriginUpperLeft
OpExecutionMode %compute_main LocalSize 1 1 1
OpName %value "value"
OpName %vertex_point_size "vertex_point_size"
OpName %arg_0 "arg_0"
OpName %textureDimensions_012b82 "textureDimensions_012b82"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
OpName %vertex_main "vertex_main"
OpName %fragment_main "fragment_main"
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
OpDecorate %arg_0 NonReadable
OpDecorate %arg_0 DescriptorSet 1
OpDecorate %arg_0 Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
%5 = OpConstantNull %v4float
%value = OpVariable %_ptr_Output_v4float Output %5
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
%int = OpTypeInt 32 1
%11 = OpTypeImage %int 2D 0 1 0 2 R32i
%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
%arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant
%void = OpTypeVoid
%13 = OpTypeFunction %void
%v2int = OpTypeVector %int 2
%v3int = OpTypeVector %int 3
%_ptr_Function_v2int = OpTypePointer Function %v2int
%24 = OpConstantNull %v2int
%25 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%textureDimensions_012b82 = OpFunction %void None %13
%16 = OpLabel
%res = OpVariable %_ptr_Function_v2int Function %24
%21 = OpLoad %11 %arg_0
%19 = OpImageQuerySize %v3int %21
%17 = OpVectorShuffle %v2int %19 %19 0 1
OpStore %res %17
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %25
%27 = OpLabel
%28 = OpFunctionCall %void %textureDimensions_012b82
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %13
%30 = OpLabel
%31 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %31
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %13
%34 = OpLabel
%35 = OpFunctionCall %void %textureDimensions_012b82
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %13
%37 = OpLabel
%38 = OpFunctionCall %void %textureDimensions_012b82
OpReturn
OpFunctionEnd

View File

@ -1,21 +0,0 @@
@group(1) @binding(0) var arg_0 : texture_storage_2d_array<r32sint, write>;
fn textureDimensions_012b82() {
var res : vec2<i32> = textureDimensions(arg_0);
}
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
textureDimensions_012b82();
return vec4<f32>();
}
@fragment
fn fragment_main() {
textureDimensions_012b82();
}
@compute @workgroup_size(1)
fn compute_main() {
textureDimensions_012b82();
}

View File

@ -0,0 +1,44 @@
// Copyright 2022 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
////////////////////////////////////////////////////////////////////////////////
// File generated by tools/src/cmd/gen
// using the template:
// test/tint/builtins/gen/gen.wgsl.tmpl
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
@group(1) @binding(0) var arg_0: texture_1d<i32>;
// fn textureDimensions(texture: texture_1d<i32>, level: u32) -> u32
fn textureDimensions_022903() {
var res: u32 = textureDimensions(arg_0, 1u);
}
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
textureDimensions_022903();
return vec4<f32>();
}
@fragment
fn fragment_main() {
textureDimensions_022903();
}
@compute @workgroup_size(1)
fn compute_main() {
textureDimensions_022903();
}

View File

@ -0,0 +1,34 @@
Texture1D<int4> arg_0 : register(t0, space1);
void textureDimensions_022903() {
int2 tint_tmp;
arg_0.GetDimensions(1u, tint_tmp.x, tint_tmp.y);
uint res = tint_tmp.x;
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
textureDimensions_022903();
return (0.0f).xxxx;
}
tint_symbol vertex_main() {
const float4 inner_result = vertex_main_inner();
tint_symbol wrapper_result = (tint_symbol)0;
wrapper_result.value = inner_result;
return wrapper_result;
}
void fragment_main() {
textureDimensions_022903();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
textureDimensions_022903();
return;
}

View File

@ -0,0 +1,34 @@
Texture1D<int4> arg_0 : register(t0, space1);
void textureDimensions_022903() {
int2 tint_tmp;
arg_0.GetDimensions(1u, tint_tmp.x, tint_tmp.y);
uint res = tint_tmp.x;
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
textureDimensions_022903();
return (0.0f).xxxx;
}
tint_symbol vertex_main() {
const float4 inner_result = vertex_main_inner();
tint_symbol wrapper_result = (tint_symbol)0;
wrapper_result.value = inner_result;
return wrapper_result;
}
void fragment_main() {
textureDimensions_022903();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
textureDimensions_022903();
return;
}

View File

@ -0,0 +1,75 @@
SKIP: FAILED
#version 310 es
uniform highp isampler1D arg_0_1;
void textureDimensions_022903() {
uint res = uint(textureSize(arg_0_1, int(1u)));
}
vec4 vertex_main() {
textureDimensions_022903();
return vec4(0.0f);
}
void main() {
gl_PointSize = 1.0;
vec4 inner_result = vertex_main();
gl_Position = inner_result;
gl_Position.y = -(gl_Position.y);
gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
return;
}
Error parsing GLSL shader:
ERROR: 0:3: 'isampler1D' : Reserved word.
ERROR: 0:3: '' : compilation terminated
ERROR: 2 compilation errors. No code generated.
#version 310 es
precision mediump float;
uniform highp isampler1D arg_0_1;
void textureDimensions_022903() {
uint res = uint(textureSize(arg_0_1, int(1u)));
}
void fragment_main() {
textureDimensions_022903();
}
void main() {
fragment_main();
return;
}
Error parsing GLSL shader:
ERROR: 0:4: 'isampler1D' : Reserved word.
ERROR: 0:4: '' : compilation terminated
ERROR: 2 compilation errors. No code generated.
#version 310 es
uniform highp isampler1D arg_0_1;
void textureDimensions_022903() {
uint res = uint(textureSize(arg_0_1, int(1u)));
}
void compute_main() {
textureDimensions_022903();
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
compute_main();
return;
}
Error parsing GLSL shader:
ERROR: 0:3: 'isampler1D' : Reserved word.
ERROR: 0:3: '' : compilation terminated
ERROR: 2 compilation errors. No code generated.

View File

@ -0,0 +1,33 @@
#include <metal_stdlib>
using namespace metal;
void textureDimensions_022903(texture1d<int, access::sample> tint_symbol_1) {
uint res = tint_symbol_1.get_width(0);
}
struct tint_symbol {
float4 value [[position]];
};
float4 vertex_main_inner(texture1d<int, access::sample> tint_symbol_2) {
textureDimensions_022903(tint_symbol_2);
return float4(0.0f);
}
vertex tint_symbol vertex_main(texture1d<int, access::sample> tint_symbol_3 [[texture(0)]]) {
float4 const inner_result = vertex_main_inner(tint_symbol_3);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
fragment void fragment_main(texture1d<int, access::sample> tint_symbol_4 [[texture(0)]]) {
textureDimensions_022903(tint_symbol_4);
return;
}
kernel void compute_main(texture1d<int, access::sample> tint_symbol_5 [[texture(0)]]) {
textureDimensions_022903(tint_symbol_5);
return;
}

View File

@ -0,0 +1,77 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 38
; Schema: 0
OpCapability Shader
OpCapability Sampled1D
OpCapability ImageQuery
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
OpEntryPoint GLCompute %compute_main "compute_main"
OpExecutionMode %fragment_main OriginUpperLeft
OpExecutionMode %compute_main LocalSize 1 1 1
OpName %value "value"
OpName %vertex_point_size "vertex_point_size"
OpName %arg_0 "arg_0"
OpName %textureDimensions_022903 "textureDimensions_022903"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
OpName %vertex_main "vertex_main"
OpName %fragment_main "fragment_main"
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
OpDecorate %arg_0 DescriptorSet 1
OpDecorate %arg_0 Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
%5 = OpConstantNull %v4float
%value = OpVariable %_ptr_Output_v4float Output %5
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
%int = OpTypeInt 32 1
%11 = OpTypeImage %int 1D 0 0 0 1 Unknown
%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
%arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant
%void = OpTypeVoid
%13 = OpTypeFunction %void
%uint = OpTypeInt 32 0
%uint_1 = OpConstant %uint 1
%_ptr_Function_uint = OpTypePointer Function %uint
%23 = OpConstantNull %uint
%24 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%textureDimensions_022903 = OpFunction %void None %13
%16 = OpLabel
%res = OpVariable %_ptr_Function_uint Function %23
%19 = OpLoad %11 %arg_0
%17 = OpImageQuerySizeLod %uint %19 %uint_1
OpStore %res %17
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %24
%26 = OpLabel
%27 = OpFunctionCall %void %textureDimensions_022903
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %13
%29 = OpLabel
%30 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %30
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %13
%33 = OpLabel
%34 = OpFunctionCall %void %textureDimensions_022903
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %13
%36 = OpLabel
%37 = OpFunctionCall %void %textureDimensions_022903
OpReturn
OpFunctionEnd

View File

@ -0,0 +1,21 @@
@group(1) @binding(0) var arg_0 : texture_1d<i32>;
fn textureDimensions_022903() {
var res : u32 = textureDimensions(arg_0, 1u);
}
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
textureDimensions_022903();
return vec4<f32>();
}
@fragment
fn fragment_main() {
textureDimensions_022903();
}
@compute @workgroup_size(1)
fn compute_main() {
textureDimensions_022903();
}

View File

@ -1,44 +0,0 @@
// Copyright 2021 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
////////////////////////////////////////////////////////////////////////////////
// File generated by tools/src/cmd/gen
// using the template:
// test/tint/builtins/gen/gen.wgsl.tmpl
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
@group(1) @binding(0) var arg_0: texture_storage_1d<rgba16sint, write>;
// fn textureDimensions(texture: texture_storage_1d<rgba16sint, write>) -> i32
fn textureDimensions_08753d() {
var res: i32 = textureDimensions(arg_0);
}
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
textureDimensions_08753d();
return vec4<f32>();
}
@fragment
fn fragment_main() {
textureDimensions_08753d();
}
@compute @workgroup_size(1)
fn compute_main() {
textureDimensions_08753d();
}

View File

@ -1,34 +0,0 @@
RWTexture1D<int4> arg_0 : register(u0, space1);
void textureDimensions_08753d() {
int tint_tmp;
arg_0.GetDimensions(tint_tmp);
int res = tint_tmp;
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
textureDimensions_08753d();
return (0.0f).xxxx;
}
tint_symbol vertex_main() {
const float4 inner_result = vertex_main_inner();
tint_symbol wrapper_result = (tint_symbol)0;
wrapper_result.value = inner_result;
return wrapper_result;
}
void fragment_main() {
textureDimensions_08753d();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
textureDimensions_08753d();
return;
}

View File

@ -1,34 +0,0 @@
RWTexture1D<int4> arg_0 : register(u0, space1);
void textureDimensions_08753d() {
int tint_tmp;
arg_0.GetDimensions(tint_tmp);
int res = tint_tmp;
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
textureDimensions_08753d();
return (0.0f).xxxx;
}
tint_symbol vertex_main() {
const float4 inner_result = vertex_main_inner();
tint_symbol wrapper_result = (tint_symbol)0;
wrapper_result.value = inner_result;
return wrapper_result;
}
void fragment_main() {
textureDimensions_08753d();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
textureDimensions_08753d();
return;
}

View File

@ -1,78 +0,0 @@
SKIP: FAILED
#version 310 es
layout(rgba16i) uniform highp writeonly iimage1D arg_0;
void textureDimensions_08753d() {
int res = imageSize(arg_0);
}
vec4 vertex_main() {
textureDimensions_08753d();
return vec4(0.0f);
}
void main() {
gl_PointSize = 1.0;
vec4 inner_result = vertex_main();
gl_Position = inner_result;
gl_Position.y = -(gl_Position.y);
gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
return;
}
Error parsing GLSL shader:
ERROR: 0:3: 'iimage1D' : Reserved word.
WARNING: 0:3: 'layout' : useless application of layout qualifier
ERROR: 0:3: '' : compilation terminated
ERROR: 2 compilation errors. No code generated.
#version 310 es
precision mediump float;
layout(rgba16i) uniform highp writeonly iimage1D arg_0;
void textureDimensions_08753d() {
int res = imageSize(arg_0);
}
void fragment_main() {
textureDimensions_08753d();
}
void main() {
fragment_main();
return;
}
Error parsing GLSL shader:
ERROR: 0:4: 'iimage1D' : Reserved word.
WARNING: 0:4: 'layout' : useless application of layout qualifier
ERROR: 0:4: '' : compilation terminated
ERROR: 2 compilation errors. No code generated.
#version 310 es
layout(rgba16i) uniform highp writeonly iimage1D arg_0;
void textureDimensions_08753d() {
int res = imageSize(arg_0);
}
void compute_main() {
textureDimensions_08753d();
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
compute_main();
return;
}
Error parsing GLSL shader:
ERROR: 0:3: 'iimage1D' : Reserved word.
WARNING: 0:3: 'layout' : useless application of layout qualifier
ERROR: 0:3: '' : compilation terminated
ERROR: 2 compilation errors. No code generated.

View File

@ -1,33 +0,0 @@
#include <metal_stdlib>
using namespace metal;
void textureDimensions_08753d(texture1d<int, access::write> tint_symbol_1) {
int res = int(tint_symbol_1.get_width(0));
}
struct tint_symbol {
float4 value [[position]];
};
float4 vertex_main_inner(texture1d<int, access::write> tint_symbol_2) {
textureDimensions_08753d(tint_symbol_2);
return float4(0.0f);
}
vertex tint_symbol vertex_main(texture1d<int, access::write> tint_symbol_3 [[texture(0)]]) {
float4 const inner_result = vertex_main_inner(tint_symbol_3);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
fragment void fragment_main(texture1d<int, access::write> tint_symbol_4 [[texture(0)]]) {
textureDimensions_08753d(tint_symbol_4);
return;
}
kernel void compute_main(texture1d<int, access::write> tint_symbol_5 [[texture(0)]]) {
textureDimensions_08753d(tint_symbol_5);
return;
}

View File

@ -1,76 +0,0 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 36
; Schema: 0
OpCapability Shader
OpCapability Image1D
OpCapability ImageQuery
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
OpEntryPoint GLCompute %compute_main "compute_main"
OpExecutionMode %fragment_main OriginUpperLeft
OpExecutionMode %compute_main LocalSize 1 1 1
OpName %value "value"
OpName %vertex_point_size "vertex_point_size"
OpName %arg_0 "arg_0"
OpName %textureDimensions_08753d "textureDimensions_08753d"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
OpName %vertex_main "vertex_main"
OpName %fragment_main "fragment_main"
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
OpDecorate %arg_0 NonReadable
OpDecorate %arg_0 DescriptorSet 1
OpDecorate %arg_0 Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
%5 = OpConstantNull %v4float
%value = OpVariable %_ptr_Output_v4float Output %5
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
%int = OpTypeInt 32 1
%11 = OpTypeImage %int 1D 0 0 0 2 Rgba16i
%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
%arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant
%void = OpTypeVoid
%13 = OpTypeFunction %void
%_ptr_Function_int = OpTypePointer Function %int
%21 = OpConstantNull %int
%22 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%textureDimensions_08753d = OpFunction %void None %13
%16 = OpLabel
%res = OpVariable %_ptr_Function_int Function %21
%18 = OpLoad %11 %arg_0
%17 = OpImageQuerySize %int %18
OpStore %res %17
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %22
%24 = OpLabel
%25 = OpFunctionCall %void %textureDimensions_08753d
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %13
%27 = OpLabel
%28 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %28
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %13
%31 = OpLabel
%32 = OpFunctionCall %void %textureDimensions_08753d
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %13
%34 = OpLabel
%35 = OpFunctionCall %void %textureDimensions_08753d
OpReturn
OpFunctionEnd

View File

@ -1,21 +0,0 @@
@group(1) @binding(0) var arg_0 : texture_storage_1d<rgba16sint, write>;
fn textureDimensions_08753d() {
var res : i32 = textureDimensions(arg_0);
}
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
textureDimensions_08753d();
return vec4<f32>();
}
@fragment
fn fragment_main() {
textureDimensions_08753d();
}
@compute @workgroup_size(1)
fn compute_main() {
textureDimensions_08753d();
}

View File

@ -0,0 +1,44 @@
// Copyright 2022 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
////////////////////////////////////////////////////////////////////////////////
// File generated by tools/src/cmd/gen
// using the template:
// test/tint/builtins/gen/gen.wgsl.tmpl
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
@group(1) @binding(0) var arg_0: texture_3d<f32>;
// fn textureDimensions(texture: texture_3d<f32>, level: u32) -> vec3<u32>
fn textureDimensions_0890c6() {
var res: vec3<u32> = textureDimensions(arg_0, 1u);
}
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
textureDimensions_0890c6();
return vec4<f32>();
}
@fragment
fn fragment_main() {
textureDimensions_0890c6();
}
@compute @workgroup_size(1)
fn compute_main() {
textureDimensions_0890c6();
}

View File

@ -0,0 +1,34 @@
Texture3D<float4> arg_0 : register(t0, space1);
void textureDimensions_0890c6() {
int4 tint_tmp;
arg_0.GetDimensions(1u, tint_tmp.x, tint_tmp.y, tint_tmp.z, tint_tmp.w);
uint3 res = tint_tmp.xyz;
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
textureDimensions_0890c6();
return (0.0f).xxxx;
}
tint_symbol vertex_main() {
const float4 inner_result = vertex_main_inner();
tint_symbol wrapper_result = (tint_symbol)0;
wrapper_result.value = inner_result;
return wrapper_result;
}
void fragment_main() {
textureDimensions_0890c6();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
textureDimensions_0890c6();
return;
}

View File

@ -0,0 +1,34 @@
Texture3D<float4> arg_0 : register(t0, space1);
void textureDimensions_0890c6() {
int4 tint_tmp;
arg_0.GetDimensions(1u, tint_tmp.x, tint_tmp.y, tint_tmp.z, tint_tmp.w);
uint3 res = tint_tmp.xyz;
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
textureDimensions_0890c6();
return (0.0f).xxxx;
}
tint_symbol vertex_main() {
const float4 inner_result = vertex_main_inner();
tint_symbol wrapper_result = (tint_symbol)0;
wrapper_result.value = inner_result;
return wrapper_result;
}
void fragment_main() {
textureDimensions_0890c6();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
textureDimensions_0890c6();
return;
}

View File

@ -0,0 +1,52 @@
#version 310 es
uniform highp sampler3D arg_0_1;
void textureDimensions_0890c6() {
uvec3 res = uvec3(textureSize(arg_0_1, int(1u)));
}
vec4 vertex_main() {
textureDimensions_0890c6();
return vec4(0.0f);
}
void main() {
gl_PointSize = 1.0;
vec4 inner_result = vertex_main();
gl_Position = inner_result;
gl_Position.y = -(gl_Position.y);
gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
return;
}
#version 310 es
precision mediump float;
uniform highp sampler3D arg_0_1;
void textureDimensions_0890c6() {
uvec3 res = uvec3(textureSize(arg_0_1, int(1u)));
}
void fragment_main() {
textureDimensions_0890c6();
}
void main() {
fragment_main();
return;
}
#version 310 es
uniform highp sampler3D arg_0_1;
void textureDimensions_0890c6() {
uvec3 res = uvec3(textureSize(arg_0_1, int(1u)));
}
void compute_main() {
textureDimensions_0890c6();
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
compute_main();
return;
}

View File

@ -0,0 +1,33 @@
#include <metal_stdlib>
using namespace metal;
void textureDimensions_0890c6(texture3d<float, access::sample> tint_symbol_1) {
uint3 res = uint3(tint_symbol_1.get_width(1u), tint_symbol_1.get_height(1u), tint_symbol_1.get_depth(1u));
}
struct tint_symbol {
float4 value [[position]];
};
float4 vertex_main_inner(texture3d<float, access::sample> tint_symbol_2) {
textureDimensions_0890c6(tint_symbol_2);
return float4(0.0f);
}
vertex tint_symbol vertex_main(texture3d<float, access::sample> tint_symbol_3 [[texture(0)]]) {
float4 const inner_result = vertex_main_inner(tint_symbol_3);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
fragment void fragment_main(texture3d<float, access::sample> tint_symbol_4 [[texture(0)]]) {
textureDimensions_0890c6(tint_symbol_4);
return;
}
kernel void compute_main(texture3d<float, access::sample> tint_symbol_5 [[texture(0)]]) {
textureDimensions_0890c6(tint_symbol_5);
return;
}

View File

@ -0,0 +1,76 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 38
; Schema: 0
OpCapability Shader
OpCapability ImageQuery
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
OpEntryPoint GLCompute %compute_main "compute_main"
OpExecutionMode %fragment_main OriginUpperLeft
OpExecutionMode %compute_main LocalSize 1 1 1
OpName %value "value"
OpName %vertex_point_size "vertex_point_size"
OpName %arg_0 "arg_0"
OpName %textureDimensions_0890c6 "textureDimensions_0890c6"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
OpName %vertex_main "vertex_main"
OpName %fragment_main "fragment_main"
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
OpDecorate %arg_0 DescriptorSet 1
OpDecorate %arg_0 Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
%5 = OpConstantNull %v4float
%value = OpVariable %_ptr_Output_v4float Output %5
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
%11 = OpTypeImage %float 3D 0 0 0 1 Unknown
%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
%arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant
%void = OpTypeVoid
%12 = OpTypeFunction %void
%uint = OpTypeInt 32 0
%v3uint = OpTypeVector %uint 3
%uint_1 = OpConstant %uint 1
%_ptr_Function_v3uint = OpTypePointer Function %v3uint
%23 = OpConstantNull %v3uint
%24 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%textureDimensions_0890c6 = OpFunction %void None %12
%15 = OpLabel
%res = OpVariable %_ptr_Function_v3uint Function %23
%19 = OpLoad %11 %arg_0
%16 = OpImageQuerySizeLod %v3uint %19 %uint_1
OpStore %res %16
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %24
%26 = OpLabel
%27 = OpFunctionCall %void %textureDimensions_0890c6
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %12
%29 = OpLabel
%30 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %30
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %12
%33 = OpLabel
%34 = OpFunctionCall %void %textureDimensions_0890c6
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %12
%36 = OpLabel
%37 = OpFunctionCall %void %textureDimensions_0890c6
OpReturn
OpFunctionEnd

View File

@ -0,0 +1,21 @@
@group(1) @binding(0) var arg_0 : texture_3d<f32>;
fn textureDimensions_0890c6() {
var res : vec3<u32> = textureDimensions(arg_0, 1u);
}
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
textureDimensions_0890c6();
return vec4<f32>();
}
@fragment
fn fragment_main() {
textureDimensions_0890c6();
}
@compute @workgroup_size(1)
fn compute_main() {
textureDimensions_0890c6();
}

View File

@ -0,0 +1,44 @@
// Copyright 2022 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
////////////////////////////////////////////////////////////////////////////////
// File generated by tools/src/cmd/gen
// using the template:
// test/tint/builtins/gen/gen.wgsl.tmpl
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
@group(1) @binding(0) var arg_0: texture_storage_1d<rgba32uint, write>;
// fn textureDimensions(texture: texture_storage_1d<rgba32uint, write>) -> u32
fn textureDimensions_09140b() {
var res: u32 = textureDimensions(arg_0);
}
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
textureDimensions_09140b();
return vec4<f32>();
}
@fragment
fn fragment_main() {
textureDimensions_09140b();
}
@compute @workgroup_size(1)
fn compute_main() {
textureDimensions_09140b();
}

View File

@ -0,0 +1,34 @@
RWTexture1D<uint4> arg_0 : register(u0, space1);
void textureDimensions_09140b() {
int tint_tmp;
arg_0.GetDimensions(tint_tmp);
uint res = tint_tmp;
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
textureDimensions_09140b();
return (0.0f).xxxx;
}
tint_symbol vertex_main() {
const float4 inner_result = vertex_main_inner();
tint_symbol wrapper_result = (tint_symbol)0;
wrapper_result.value = inner_result;
return wrapper_result;
}
void fragment_main() {
textureDimensions_09140b();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
textureDimensions_09140b();
return;
}

View File

@ -0,0 +1,34 @@
RWTexture1D<uint4> arg_0 : register(u0, space1);
void textureDimensions_09140b() {
int tint_tmp;
arg_0.GetDimensions(tint_tmp);
uint res = tint_tmp;
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
textureDimensions_09140b();
return (0.0f).xxxx;
}
tint_symbol vertex_main() {
const float4 inner_result = vertex_main_inner();
tint_symbol wrapper_result = (tint_symbol)0;
wrapper_result.value = inner_result;
return wrapper_result;
}
void fragment_main() {
textureDimensions_09140b();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
textureDimensions_09140b();
return;
}

View File

@ -0,0 +1,78 @@
SKIP: FAILED
#version 310 es
layout(rgba32ui) uniform highp writeonly uimage1D arg_0;
void textureDimensions_09140b() {
uint res = uint(imageSize(arg_0));
}
vec4 vertex_main() {
textureDimensions_09140b();
return vec4(0.0f);
}
void main() {
gl_PointSize = 1.0;
vec4 inner_result = vertex_main();
gl_Position = inner_result;
gl_Position.y = -(gl_Position.y);
gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
return;
}
Error parsing GLSL shader:
ERROR: 0:3: 'uimage1D' : Reserved word.
WARNING: 0:3: 'layout' : useless application of layout qualifier
ERROR: 0:3: '' : compilation terminated
ERROR: 2 compilation errors. No code generated.
#version 310 es
precision mediump float;
layout(rgba32ui) uniform highp writeonly uimage1D arg_0;
void textureDimensions_09140b() {
uint res = uint(imageSize(arg_0));
}
void fragment_main() {
textureDimensions_09140b();
}
void main() {
fragment_main();
return;
}
Error parsing GLSL shader:
ERROR: 0:4: 'uimage1D' : Reserved word.
WARNING: 0:4: 'layout' : useless application of layout qualifier
ERROR: 0:4: '' : compilation terminated
ERROR: 2 compilation errors. No code generated.
#version 310 es
layout(rgba32ui) uniform highp writeonly uimage1D arg_0;
void textureDimensions_09140b() {
uint res = uint(imageSize(arg_0));
}
void compute_main() {
textureDimensions_09140b();
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
compute_main();
return;
}
Error parsing GLSL shader:
ERROR: 0:3: 'uimage1D' : Reserved word.
WARNING: 0:3: 'layout' : useless application of layout qualifier
ERROR: 0:3: '' : compilation terminated
ERROR: 2 compilation errors. No code generated.

View File

@ -0,0 +1,33 @@
#include <metal_stdlib>
using namespace metal;
void textureDimensions_09140b(texture1d<uint, access::write> tint_symbol_1) {
uint res = tint_symbol_1.get_width(0);
}
struct tint_symbol {
float4 value [[position]];
};
float4 vertex_main_inner(texture1d<uint, access::write> tint_symbol_2) {
textureDimensions_09140b(tint_symbol_2);
return float4(0.0f);
}
vertex tint_symbol vertex_main(texture1d<uint, access::write> tint_symbol_3 [[texture(0)]]) {
float4 const inner_result = vertex_main_inner(tint_symbol_3);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
fragment void fragment_main(texture1d<uint, access::write> tint_symbol_4 [[texture(0)]]) {
textureDimensions_09140b(tint_symbol_4);
return;
}
kernel void compute_main(texture1d<uint, access::write> tint_symbol_5 [[texture(0)]]) {
textureDimensions_09140b(tint_symbol_5);
return;
}

View File

@ -0,0 +1,76 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 36
; Schema: 0
OpCapability Shader
OpCapability Image1D
OpCapability ImageQuery
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
OpEntryPoint GLCompute %compute_main "compute_main"
OpExecutionMode %fragment_main OriginUpperLeft
OpExecutionMode %compute_main LocalSize 1 1 1
OpName %value "value"
OpName %vertex_point_size "vertex_point_size"
OpName %arg_0 "arg_0"
OpName %textureDimensions_09140b "textureDimensions_09140b"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
OpName %vertex_main "vertex_main"
OpName %fragment_main "fragment_main"
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
OpDecorate %arg_0 NonReadable
OpDecorate %arg_0 DescriptorSet 1
OpDecorate %arg_0 Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
%5 = OpConstantNull %v4float
%value = OpVariable %_ptr_Output_v4float Output %5
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
%uint = OpTypeInt 32 0
%11 = OpTypeImage %uint 1D 0 0 0 2 Rgba32ui
%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
%arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant
%void = OpTypeVoid
%13 = OpTypeFunction %void
%_ptr_Function_uint = OpTypePointer Function %uint
%21 = OpConstantNull %uint
%22 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%textureDimensions_09140b = OpFunction %void None %13
%16 = OpLabel
%res = OpVariable %_ptr_Function_uint Function %21
%18 = OpLoad %11 %arg_0
%17 = OpImageQuerySize %uint %18
OpStore %res %17
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %22
%24 = OpLabel
%25 = OpFunctionCall %void %textureDimensions_09140b
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %13
%27 = OpLabel
%28 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %28
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %13
%31 = OpLabel
%32 = OpFunctionCall %void %textureDimensions_09140b
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %13
%34 = OpLabel
%35 = OpFunctionCall %void %textureDimensions_09140b
OpReturn
OpFunctionEnd

View File

@ -0,0 +1,21 @@
@group(1) @binding(0) var arg_0 : texture_storage_1d<rgba32uint, write>;
fn textureDimensions_09140b() {
var res : u32 = textureDimensions(arg_0);
}
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
textureDimensions_09140b();
return vec4<f32>();
}
@fragment
fn fragment_main() {
textureDimensions_09140b();
}
@compute @workgroup_size(1)
fn compute_main() {
textureDimensions_09140b();
}

View File

@ -1,44 +0,0 @@
// Copyright 2022 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
////////////////////////////////////////////////////////////////////////////////
// File generated by tools/src/cmd/gen
// using the template:
// test/tint/builtins/gen/gen.wgsl.tmpl
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
@group(1) @binding(0) var arg_0: texture_depth_2d_array;
// fn textureDimensions(texture: texture_depth_2d_array, level: u32) -> vec2<i32>
fn textureDimensions_0a1ce8() {
var res: vec2<i32> = textureDimensions(arg_0, 1u);
}
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
textureDimensions_0a1ce8();
return vec4<f32>();
}
@fragment
fn fragment_main() {
textureDimensions_0a1ce8();
}
@compute @workgroup_size(1)
fn compute_main() {
textureDimensions_0a1ce8();
}

View File

@ -1,34 +0,0 @@
Texture2DArray arg_0 : register(t0, space1);
void textureDimensions_0a1ce8() {
int4 tint_tmp;
arg_0.GetDimensions(1u, tint_tmp.x, tint_tmp.y, tint_tmp.z, tint_tmp.w);
int2 res = tint_tmp.xy;
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
textureDimensions_0a1ce8();
return (0.0f).xxxx;
}
tint_symbol vertex_main() {
const float4 inner_result = vertex_main_inner();
tint_symbol wrapper_result = (tint_symbol)0;
wrapper_result.value = inner_result;
return wrapper_result;
}
void fragment_main() {
textureDimensions_0a1ce8();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
textureDimensions_0a1ce8();
return;
}

View File

@ -1,34 +0,0 @@
Texture2DArray arg_0 : register(t0, space1);
void textureDimensions_0a1ce8() {
int4 tint_tmp;
arg_0.GetDimensions(1u, tint_tmp.x, tint_tmp.y, tint_tmp.z, tint_tmp.w);
int2 res = tint_tmp.xy;
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
textureDimensions_0a1ce8();
return (0.0f).xxxx;
}
tint_symbol vertex_main() {
const float4 inner_result = vertex_main_inner();
tint_symbol wrapper_result = (tint_symbol)0;
wrapper_result.value = inner_result;
return wrapper_result;
}
void fragment_main() {
textureDimensions_0a1ce8();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
textureDimensions_0a1ce8();
return;
}

View File

@ -1,52 +0,0 @@
#version 310 es
uniform highp sampler2DArray arg_0_1;
void textureDimensions_0a1ce8() {
ivec2 res = textureSize(arg_0_1, int(1u)).xy;
}
vec4 vertex_main() {
textureDimensions_0a1ce8();
return vec4(0.0f);
}
void main() {
gl_PointSize = 1.0;
vec4 inner_result = vertex_main();
gl_Position = inner_result;
gl_Position.y = -(gl_Position.y);
gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
return;
}
#version 310 es
precision mediump float;
uniform highp sampler2DArray arg_0_1;
void textureDimensions_0a1ce8() {
ivec2 res = textureSize(arg_0_1, int(1u)).xy;
}
void fragment_main() {
textureDimensions_0a1ce8();
}
void main() {
fragment_main();
return;
}
#version 310 es
uniform highp sampler2DArray arg_0_1;
void textureDimensions_0a1ce8() {
ivec2 res = textureSize(arg_0_1, int(1u)).xy;
}
void compute_main() {
textureDimensions_0a1ce8();
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
compute_main();
return;
}

View File

@ -1,33 +0,0 @@
#include <metal_stdlib>
using namespace metal;
void textureDimensions_0a1ce8(depth2d_array<float, access::sample> tint_symbol_1) {
int2 res = int2(tint_symbol_1.get_width(1u), tint_symbol_1.get_height(1u));
}
struct tint_symbol {
float4 value [[position]];
};
float4 vertex_main_inner(depth2d_array<float, access::sample> tint_symbol_2) {
textureDimensions_0a1ce8(tint_symbol_2);
return float4(0.0f);
}
vertex tint_symbol vertex_main(depth2d_array<float, access::sample> tint_symbol_3 [[texture(0)]]) {
float4 const inner_result = vertex_main_inner(tint_symbol_3);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
fragment void fragment_main(depth2d_array<float, access::sample> tint_symbol_4 [[texture(0)]]) {
textureDimensions_0a1ce8(tint_symbol_4);
return;
}
kernel void compute_main(depth2d_array<float, access::sample> tint_symbol_5 [[texture(0)]]) {
textureDimensions_0a1ce8(tint_symbol_5);
return;
}

View File

@ -1,79 +0,0 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 41
; Schema: 0
OpCapability Shader
OpCapability ImageQuery
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
OpEntryPoint GLCompute %compute_main "compute_main"
OpExecutionMode %fragment_main OriginUpperLeft
OpExecutionMode %compute_main LocalSize 1 1 1
OpName %value "value"
OpName %vertex_point_size "vertex_point_size"
OpName %arg_0 "arg_0"
OpName %textureDimensions_0a1ce8 "textureDimensions_0a1ce8"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
OpName %vertex_main "vertex_main"
OpName %fragment_main "fragment_main"
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
OpDecorate %arg_0 DescriptorSet 1
OpDecorate %arg_0 Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
%5 = OpConstantNull %v4float
%value = OpVariable %_ptr_Output_v4float Output %5
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
%11 = OpTypeImage %float 2D 0 1 0 1 Unknown
%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
%arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant
%void = OpTypeVoid
%12 = OpTypeFunction %void
%int = OpTypeInt 32 1
%v2int = OpTypeVector %int 2
%v3int = OpTypeVector %int 3
%uint = OpTypeInt 32 0
%uint_1 = OpConstant %uint 1
%_ptr_Function_v2int = OpTypePointer Function %v2int
%26 = OpConstantNull %v2int
%27 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%textureDimensions_0a1ce8 = OpFunction %void None %12
%15 = OpLabel
%res = OpVariable %_ptr_Function_v2int Function %26
%21 = OpLoad %11 %arg_0
%19 = OpImageQuerySizeLod %v3int %21 %uint_1
%16 = OpVectorShuffle %v2int %19 %19 0 1
OpStore %res %16
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %27
%29 = OpLabel
%30 = OpFunctionCall %void %textureDimensions_0a1ce8
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %12
%32 = OpLabel
%33 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %33
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %12
%36 = OpLabel
%37 = OpFunctionCall %void %textureDimensions_0a1ce8
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %12
%39 = OpLabel
%40 = OpFunctionCall %void %textureDimensions_0a1ce8
OpReturn
OpFunctionEnd

Some files were not shown because too many files have changed in this diff Show More