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

Resolver: Remove texture_storage read access 

The intrinsics that did anything useful with this were deprecated
several releases ago.

Change-Id: I79e3c901b6a78583853a067ec46cfa98e346517c
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/66262
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
This commit is contained in:
Ben Clayton 2021-10-14 10:10:45 +00:00 committed by Tint LUCI CQ
parent 14fc622161
commit f3f2d0a218
785 changed files with 3985 additions and 37787 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
docs/origin-trial-changes.md
View File

@ -5,6 +5,8 @@
### Breaking Changes ### Breaking Changes
* Deprecated `modf()` and `frexp()` builtin overloads that take a pointer second parameter have been removed. * Deprecated `modf()` and `frexp()` builtin overloads that take a pointer second parameter have been removed.
* Deprecated texture builtin functions that accepted a `read` access controlled storage texture have been removed.
* Storage textures must now only use the `write` access control.
## Changes for M95 ## Changes for M95

5
samples/main.cc
View File

@ -346,9 +346,6 @@ std::string ResourceTypeToString(
return "SampledTexture"; return "SampledTexture";
case tint::inspector::ResourceBinding::ResourceType::kMultisampledTexture: case tint::inspector::ResourceBinding::ResourceType::kMultisampledTexture:
return "MultisampledTexture"; return "MultisampledTexture";
case tint::inspector::ResourceBinding::ResourceType::
kReadOnlyStorageTexture:
return "ReadOnlyStorageTexture";
case tint::inspector::ResourceBinding::ResourceType:: case tint::inspector::ResourceBinding::ResourceType::
kWriteOnlyStorageTexture: kWriteOnlyStorageTexture:
return "WriteOnlyStorageTexture"; return "WriteOnlyStorageTexture";
@ -719,8 +716,6 @@ bool GenerateMsl(const tint::Program* program, const Options& options) {
case tint::inspector::ResourceBinding::ResourceType::kSampledTexture: case tint::inspector::ResourceBinding::ResourceType::kSampledTexture:
case tint::inspector::ResourceBinding::ResourceType:: case tint::inspector::ResourceBinding::ResourceType::
kMultisampledTexture: kMultisampledTexture:
case tint::inspector::ResourceBinding::ResourceType::
kReadOnlyStorageTexture:
case tint::inspector::ResourceBinding::ResourceType:: case tint::inspector::ResourceBinding::ResourceType::
kWriteOnlyStorageTexture: kWriteOnlyStorageTexture:
case tint::inspector::ResourceBinding::ResourceType::kDepthTexture: case tint::inspector::ResourceBinding::ResourceType::kDepthTexture:

312
src/ast/intrinsic_texture_helper_test.cc
View File

@ -418,49 +418,6 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
"textureDimensions", "textureDimensions",
[](ProgramBuilder* b) { return b->ExprList("texture"); }, [](ProgramBuilder* b) { return b->ExprList("texture"); },
}, },
{
ValidTextureOverload::kDimensionsStorageRO1d,
"textureDimensions(t : texture_storage_1d<rgba32float>) -> i32",
ast::Access::kRead,
ast::ImageFormat::kRgba32Float,
ast::TextureDimension::k1d,
TextureDataType::kF32,
"textureDimensions",
[](ProgramBuilder* b) { return b->ExprList("texture"); },
},
{
ValidTextureOverload::kDimensionsStorageRO2d,
"textureDimensions(t : texture_storage_2d<rgba32float>) -> "
"vec2<i32>",
ast::Access::kRead,
ast::ImageFormat::kRgba32Float,
ast::TextureDimension::k2d,
TextureDataType::kF32,
"textureDimensions",
[](ProgramBuilder* b) { return b->ExprList("texture"); },
},
{
ValidTextureOverload::kDimensionsStorageRO2dArray,
"textureDimensions(t : texture_storage_2d_array<rgba32float>) -> "
"vec2<i32>",
ast::Access::kRead,
ast::ImageFormat::kRgba32Float,
ast::TextureDimension::k2dArray,
TextureDataType::kF32,
"textureDimensions",
[](ProgramBuilder* b) { return b->ExprList("texture"); },
},
{
ValidTextureOverload::kDimensionsStorageRO3d,
"textureDimensions(t : texture_storage_3d<rgba32float>) -> "
"vec3<i32>",
ast::Access::kRead,
ast::ImageFormat::kRgba32Float,
ast::TextureDimension::k3d,
TextureDataType::kF32,
"textureDimensions",
[](ProgramBuilder* b) { return b->ExprList("texture"); },
},
{ {
ValidTextureOverload::kDimensionsStorageWO1d, ValidTextureOverload::kDimensionsStorageWO1d,
"textureDimensions(t : texture_storage_1d<rgba32float>) -> i32", "textureDimensions(t : texture_storage_1d<rgba32float>) -> i32",
@ -1899,275 +1856,6 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
4); // level 4); // level
}, },
}, },
{
ValidTextureOverload::kLoadStorageRO1dRgba32float,
"textureLoad(t : texture_storage_1d<rgba32float>,\n"
" coords : i32) -> vec4<f32>",
ast::Access::kRead,
ast::ImageFormat::kRgba32Float,
ast::TextureDimension::k1d,
TextureDataType::kF32,
"textureLoad",
[](ProgramBuilder* b) {
return b->ExprList("texture", // t
1); // coords
},
},
{
ValidTextureOverload::kLoadStorageRO2dRgba8unorm,
"textureLoad(t : texture_storage_2d<rgba8unorm>,\n"
" coords : vec2<i32>) -> vec4<f32>",
ast::Access::kRead,
ast::ImageFormat::kRgba8Unorm,
ast::TextureDimension::k2d,
TextureDataType::kF32,
"textureLoad",
[](ProgramBuilder* b) {
return b->ExprList("texture", // t
b->vec2<i32>(1, 2)); // array_index
},
},
{
ValidTextureOverload::kLoadStorageRO2dRgba8snorm,
"textureLoad(t : texture_storage_2d<rgba8snorm>,\n"
" coords : vec2<i32>) -> vec4<f32>",
ast::Access::kRead,
ast::ImageFormat::kRgba8Snorm,
ast::TextureDimension::k2d,
TextureDataType::kF32,
"textureLoad",
[](ProgramBuilder* b) {
return b->ExprList("texture", // t
b->vec2<i32>(1, 2)); // array_index
},
},
{
ValidTextureOverload::kLoadStorageRO2dRgba8uint,
"textureLoad(t : texture_storage_2d<rgba8uint>,\n"
" coords : vec2<i32>) -> vec4<u32>",
ast::Access::kRead,
ast::ImageFormat::kRgba8Uint,
ast::TextureDimension::k2d,
TextureDataType::kU32,
"textureLoad",
[](ProgramBuilder* b) {
return b->ExprList("texture", // t
b->vec2<i32>(1, 2)); // array_index
},
},
{
ValidTextureOverload::kLoadStorageRO2dRgba8sint,
"textureLoad(t : texture_storage_2d<rgba8sint>,\n"
" coords : vec2<i32>) -> vec4<i32>",
ast::Access::kRead,
ast::ImageFormat::kRgba8Sint,
ast::TextureDimension::k2d,
TextureDataType::kI32,
"textureLoad",
[](ProgramBuilder* b) {
return b->ExprList("texture", // t
b->vec2<i32>(1, 2)); // array_index
},
},
{
ValidTextureOverload::kLoadStorageRO2dRgba16uint,
"textureLoad(t : texture_storage_2d<rgba16uint>,\n"
" coords : vec2<i32>) -> vec4<u32>",
ast::Access::kRead,
ast::ImageFormat::kRgba16Uint,
ast::TextureDimension::k2d,
TextureDataType::kU32,
"textureLoad",
[](ProgramBuilder* b) {
return b->ExprList("texture", // t
b->vec2<i32>(1, 2)); // array_index
},
},
{
ValidTextureOverload::kLoadStorageRO2dRgba16sint,
"textureLoad(t : texture_storage_2d<rgba16sint>,\n"
" coords : vec2<i32>) -> vec4<i32>",
ast::Access::kRead,
ast::ImageFormat::kRgba16Sint,
ast::TextureDimension::k2d,
TextureDataType::kI32,
"textureLoad",
[](ProgramBuilder* b) {
return b->ExprList("texture", // t
b->vec2<i32>(1, 2)); // array_index
},
},
{
ValidTextureOverload::kLoadStorageRO2dRgba16float,
"textureLoad(t : texture_storage_2d<rgba16float>,\n"
" coords : vec2<i32>) -> vec4<f32>",
ast::Access::kRead,
ast::ImageFormat::kRgba16Float,
ast::TextureDimension::k2d,
TextureDataType::kF32,
"textureLoad",
[](ProgramBuilder* b) {
return b->ExprList("texture", // t
b->vec2<i32>(1, 2)); // array_index
},
},
{
ValidTextureOverload::kLoadStorageRO2dR32uint,
"textureLoad(t : texture_storage_2d<r32uint>,\n"
" coords : vec2<i32>) -> vec4<u32>",
ast::Access::kRead,
ast::ImageFormat::kR32Uint,
ast::TextureDimension::k2d,
TextureDataType::kU32,
"textureLoad",
[](ProgramBuilder* b) {
return b->ExprList("texture", // t
b->vec2<i32>(1, 2)); // array_index
},
},
{
ValidTextureOverload::kLoadStorageRO2dR32sint,
"textureLoad(t : texture_storage_2d<r32sint>,\n"
" coords : vec2<i32>) -> vec4<i32>",
ast::Access::kRead,
ast::ImageFormat::kR32Sint,
ast::TextureDimension::k2d,
TextureDataType::kI32,
"textureLoad",
[](ProgramBuilder* b) {
return b->ExprList("texture", // t
b->vec2<i32>(1, 2)); // array_index
},
},
{
ValidTextureOverload::kLoadStorageRO2dR32float,
"textureLoad(t : texture_storage_2d<r32float>,\n"
" coords : vec2<i32>) -> vec4<f32>",
ast::Access::kRead,
ast::ImageFormat::kR32Float,
ast::TextureDimension::k2d,
TextureDataType::kF32,
"textureLoad",
[](ProgramBuilder* b) {
return b->ExprList("texture", // t
b->vec2<i32>(1, 2)); // array_index
},
},
{
ValidTextureOverload::kLoadStorageRO2dRg32uint,
"textureLoad(t : texture_storage_2d<rg32uint>,\n"
" coords : vec2<i32>) -> vec4<u32>",
ast::Access::kRead,
ast::ImageFormat::kRg32Uint,
ast::TextureDimension::k2d,
TextureDataType::kU32,
"textureLoad",
[](ProgramBuilder* b) {
return b->ExprList("texture", // t
b->vec2<i32>(1, 2)); // array_index
},
},
{
ValidTextureOverload::kLoadStorageRO2dRg32sint,
"textureLoad(t : texture_storage_2d<rg32sint>,\n"
" coords : vec2<i32>) -> vec4<i32>",
ast::Access::kRead,
ast::ImageFormat::kRg32Sint,
ast::TextureDimension::k2d,
TextureDataType::kI32,
"textureLoad",
[](ProgramBuilder* b) {
return b->ExprList("texture", // t
b->vec2<i32>(1, 2)); // array_index
},
},
{
ValidTextureOverload::kLoadStorageRO2dRg32float,
"textureLoad(t : texture_storage_2d<rg32float>,\n"
" coords : vec2<i32>) -> vec4<f32>",
ast::Access::kRead,
ast::ImageFormat::kRg32Float,
ast::TextureDimension::k2d,
TextureDataType::kF32,
"textureLoad",
[](ProgramBuilder* b) {
return b->ExprList("texture", // t
b->vec2<i32>(1, 2)); // array_index
},
},
{
ValidTextureOverload::kLoadStorageRO2dRgba32uint,
"textureLoad(t : texture_storage_2d<rgba32uint>,\n"
" coords : vec2<i32>) -> vec4<u32>",
ast::Access::kRead,
ast::ImageFormat::kRgba32Uint,
ast::TextureDimension::k2d,
TextureDataType::kU32,
"textureLoad",
[](ProgramBuilder* b) {
return b->ExprList("texture", // t
b->vec2<i32>(1, 2)); // array_index
},
},
{
ValidTextureOverload::kLoadStorageRO2dRgba32sint,
"textureLoad(t : texture_storage_2d<rgba32sint>,\n"
" coords : vec2<i32>) -> vec4<i32>",
ast::Access::kRead,
ast::ImageFormat::kRgba32Sint,
ast::TextureDimension::k2d,
TextureDataType::kI32,
"textureLoad",
[](ProgramBuilder* b) {
return b->ExprList("texture", // t
b->vec2<i32>(1, 2)); // array_index
},
},
{
ValidTextureOverload::kLoadStorageRO2dRgba32float,
"textureLoad(t : texture_storage_2d<rgba32float>,\n"
" coords : vec2<i32>) -> vec4<f32>",
ast::Access::kRead,
ast::ImageFormat::kRgba32Float,
ast::TextureDimension::k2d,
TextureDataType::kF32,
"textureLoad",
[](ProgramBuilder* b) {
return b->ExprList("texture", // t
b->vec2<i32>(1, 2)); // array_index
},
},
{
ValidTextureOverload::kLoadStorageRO2dArrayRgba32float,
"textureLoad(t : "
"texture_storage_2d_array<rgba32float>,\n"
" coords : vec2<i32>,\n"
" array_index : i32) -> vec4<f32>",
ast::Access::kRead,
ast::ImageFormat::kRgba32Float,
ast::TextureDimension::k2dArray,
TextureDataType::kF32,
"textureLoad",
[](ProgramBuilder* b) {
return b->ExprList("texture", // t
b->vec2<i32>(1, 2), // coords
3); // array_index
},
},
{
ValidTextureOverload::kLoadStorageRO3dRgba32float,
"textureLoad(t : texture_storage_3d<rgba32float>,\n"
" coords : vec3<i32>) -> vec4<f32>",
ast::Access::kRead,
ast::ImageFormat::kRgba32Float,
ast::TextureDimension::k3d,
TextureDataType::kF32,
"textureLoad",
[](ProgramBuilder* b) {
return b->ExprList("texture", // t
b->vec3<i32>(1, 2, 3)); // coords
},
},
{ {
ValidTextureOverload::kStoreWO1dRgba32float, ValidTextureOverload::kStoreWO1dRgba32float,
"textureStore(t : texture_storage_1d<rgba32float>,\n" "textureStore(t : texture_storage_1d<rgba32float>,\n"

23
src/ast/intrinsic_texture_helper_test.h
View File

@ -61,10 +61,6 @@ enum class ValidTextureOverload {
kDimensionsDepthCubeArray, kDimensionsDepthCubeArray,
kDimensionsDepthCubeArrayLevel, kDimensionsDepthCubeArrayLevel,
kDimensionsDepthMultisampled2d, kDimensionsDepthMultisampled2d,
kDimensionsStorageRO1d,
kDimensionsStorageRO2d,
kDimensionsStorageRO2dArray,
kDimensionsStorageRO3d,
kDimensionsStorageWO1d, kDimensionsStorageWO1d,
kDimensionsStorageWO2d, kDimensionsStorageWO2d,
kDimensionsStorageWO2dArray, kDimensionsStorageWO2dArray,
@ -160,25 +156,6 @@ enum class ValidTextureOverload {
kLoadDepth2dLevelF32, kLoadDepth2dLevelF32,
kLoadDepth2dArrayLevelF32, kLoadDepth2dArrayLevelF32,
kLoadDepthMultisampled2dF32, kLoadDepthMultisampled2dF32,
kLoadStorageRO1dRgba32float, // Not permutated for all texel formats
kLoadStorageRO2dRgba8unorm,
kLoadStorageRO2dRgba8snorm,
kLoadStorageRO2dRgba8uint,
kLoadStorageRO2dRgba8sint,
kLoadStorageRO2dRgba16uint,
kLoadStorageRO2dRgba16sint,
kLoadStorageRO2dRgba16float,
kLoadStorageRO2dR32uint,
kLoadStorageRO2dR32sint,
kLoadStorageRO2dR32float,
kLoadStorageRO2dRg32uint,
kLoadStorageRO2dRg32sint,
kLoadStorageRO2dRg32float,
kLoadStorageRO2dRgba32uint,
kLoadStorageRO2dRgba32sint,
kLoadStorageRO2dRgba32float,
kLoadStorageRO2dArrayRgba32float, // Not permutated for all texel formats
kLoadStorageRO3dRgba32float, // Not permutated for all texel formats
kStoreWO1dRgba32float, // Not permutated for all texel formats kStoreWO1dRgba32float, // Not permutated for all texel formats
kStoreWO2dRgba32float, // Not permutated for all texel formats kStoreWO2dRgba32float, // Not permutated for all texel formats
kStoreWO2dArrayRgba32float, // Not permutated for all texel formats kStoreWO2dArrayRgba32float, // Not permutated for all texel formats

6
src/ast/module_clone_test.cc
View File

@ -48,10 +48,10 @@ type t1 = array<vec4<f32>>;
var<private> g0 : u32 = 20u; var<private> g0 : u32 = 20u;
var<private> g1 : f32 = 123.0; var<private> g1 : f32 = 123.0;
[[group(0), binding(0)]] var g2 : texture_2d<f32>; [[group(0), binding(0)]] var g2 : texture_2d<f32>;
[[group(1), binding(0)]] var g3 : texture_storage_2d<r32uint, read>; [[group(1), binding(0)]] var g3 : texture_depth_2d;
[[group(2), binding(0)]] var g4 : texture_storage_2d<rg32float, write>; [[group(2), binding(0)]] var g4 : texture_storage_2d<rg32float, write>;
[[group(3), binding(0)]] var g5 : texture_storage_2d<r32uint, read>; [[group(3), binding(0)]] var g5 : texture_depth_cube_array;
[[group(4), binding(0)]] var g6 : texture_storage_2d<rg32float, write>; [[group(4), binding(0)]] var g6 : texture_external;
var<private> g7 : vec3<f32>; var<private> g7 : vec3<f32>;
[[group(0), binding(1)]] var<storage, write> g8 : S0; [[group(0), binding(1)]] var<storage, write> g8 : S0;

19
src/inspector/inspector.cc
View File

@ -364,7 +364,6 @@ std::vector<ResourceBinding> Inspector::GetResourceBindings(
&Inspector::GetComparisonSamplerResourceBindings, &Inspector::GetComparisonSamplerResourceBindings,
&Inspector::GetSampledTextureResourceBindings, &Inspector::GetSampledTextureResourceBindings,
&Inspector::GetMultisampledTextureResourceBindings, &Inspector::GetMultisampledTextureResourceBindings,
&Inspector::GetReadOnlyStorageTextureResourceBindings,
&Inspector::GetWriteOnlyStorageTextureResourceBindings, &Inspector::GetWriteOnlyStorageTextureResourceBindings,
&Inspector::GetDepthTextureResourceBindings, &Inspector::GetDepthTextureResourceBindings,
&Inspector::GetDepthMultisampledTextureResourceBindings, &Inspector::GetDepthMultisampledTextureResourceBindings,
@ -481,16 +480,10 @@ std::vector<ResourceBinding> Inspector::GetMultisampledTextureResourceBindings(
return GetSampledTextureResourceBindingsImpl(entry_point, true); return GetSampledTextureResourceBindingsImpl(entry_point, true);
} }
std::vector<ResourceBinding>
Inspector::GetReadOnlyStorageTextureResourceBindings(
const std::string& entry_point) {
return GetStorageTextureResourceBindingsImpl(entry_point, true);
}
std::vector<ResourceBinding> std::vector<ResourceBinding>
Inspector::GetWriteOnlyStorageTextureResourceBindings( Inspector::GetWriteOnlyStorageTextureResourceBindings(
const std::string& entry_point) { const std::string& entry_point) {
return GetStorageTextureResourceBindingsImpl(entry_point, false); return GetStorageTextureResourceBindingsImpl(entry_point);
} }
std::vector<ResourceBinding> Inspector::GetTextureResourceBindings( std::vector<ResourceBinding> Inspector::GetTextureResourceBindings(
@ -751,8 +744,7 @@ std::vector<ResourceBinding> Inspector::GetSampledTextureResourceBindingsImpl(
} }
std::vector<ResourceBinding> Inspector::GetStorageTextureResourceBindingsImpl( std::vector<ResourceBinding> Inspector::GetStorageTextureResourceBindingsImpl(
const std::string& entry_point, const std::string& entry_point) {
bool read_only) {
auto* func = FindEntryPointByName(entry_point); auto* func = FindEntryPointByName(entry_point);
if (!func) { if (!func) {
return {}; return {};
@ -766,14 +758,9 @@ std::vector<ResourceBinding> Inspector::GetStorageTextureResourceBindingsImpl(
auto* texture_type = var->Type()->UnwrapRef()->As<sem::StorageTexture>(); auto* texture_type = var->Type()->UnwrapRef()->As<sem::StorageTexture>();
if (read_only != (texture_type->access() == ast::Access::kRead)) {
continue;
}
ResourceBinding entry; ResourceBinding entry;
entry.resource_type = entry.resource_type =
read_only ? ResourceBinding::ResourceType::kReadOnlyStorageTexture ResourceBinding::ResourceType::kWriteOnlyStorageTexture;
: ResourceBinding::ResourceType::kWriteOnlyStorageTexture;
entry.bind_group = binding_info.group->value(); entry.bind_group = binding_info.group->value();
entry.binding = binding_info.binding->value(); entry.binding = binding_info.binding->value();

12
src/inspector/inspector.h
View File

@ -106,11 +106,6 @@ class Inspector {
std::vector<ResourceBinding> GetMultisampledTextureResourceBindings( std::vector<ResourceBinding> GetMultisampledTextureResourceBindings(
const std::string& entry_point); const std::string& entry_point);
/// @param entry_point name of the entry point to get information about.
/// @returns vector of all of the bindings for read-only storage textures.
std::vector<ResourceBinding> GetReadOnlyStorageTextureResourceBindings(
const std::string& entry_point);
/// @param entry_point name of the entry point to get information about. /// @param entry_point name of the entry point to get information about.
/// @returns vector of all of the bindings for write-only storage textures. /// @returns vector of all of the bindings for write-only storage textures.
std::vector<ResourceBinding> GetWriteOnlyStorageTextureResourceBindings( std::vector<ResourceBinding> GetWriteOnlyStorageTextureResourceBindings(
@ -202,14 +197,11 @@ class Inspector {
bool multisampled_only); bool multisampled_only);
/// @param entry_point name of the entry point to get information about. /// @param entry_point name of the entry point to get information about.
/// @param read_only if true get only read-only bindings, otherwise get
/// write-only bindings.
/// @returns vector of all of the bindings for the requested storage textures. /// @returns vector of all of the bindings for the requested storage textures.
std::vector<ResourceBinding> GetStorageTextureResourceBindingsImpl( std::vector<ResourceBinding> GetStorageTextureResourceBindingsImpl(
const std::string& entry_point, const std::string& entry_point);
bool read_only);
/// Constructes |sampler_targets_| if it hasn't already been instantiated. /// Constructs |sampler_targets_| if it hasn't already been instantiated.
void GenerateSamplerTargets(); void GenerateSamplerTargets();
/// For a N-uple of expressions, resolve to the appropriate global resources /// For a N-uple of expressions, resolve to the appropriate global resources

59
src/inspector/inspector_test.cc
View File

@ -134,7 +134,7 @@ typedef std::tuple<ast::ImageFormat,
ResourceBinding::ImageFormat, ResourceBinding::ImageFormat,
ResourceBinding::SampledKind> ResourceBinding::SampledKind>
ImageFormatParams; ImageFormatParams;
typedef std::tuple<bool, DimensionParams, ImageFormatParams> typedef std::tuple<DimensionParams, ImageFormatParams>
GetStorageTextureTestParams; GetStorageTextureTestParams;
class InspectorGetStorageTextureResourceBindingsTestWithParam class InspectorGetStorageTextureResourceBindingsTestWithParam
: public InspectorBuilder, : public InspectorBuilder,
@ -1267,18 +1267,13 @@ TEST_F(InspectorGetResourceBindingsTest, Simple) {
Func("depth_ms_func", {}, ty.void_(), {Ignore("depth_ms_texture")}); Func("depth_ms_func", {}, ty.void_(), {Ignore("depth_ms_texture")});
auto* st_type = MakeStorageTextureTypes(ast::TextureDimension::k2d, auto* st_type = MakeStorageTextureTypes(ast::TextureDimension::k2d,
ast::ImageFormat::kR32Uint, false); ast::ImageFormat::kR32Uint);
AddStorageTexture("st_var", st_type, 4, 0); AddStorageTexture("st_var", st_type, 4, 0);
MakeStorageTextureBodyFunction("st_func", "st_var", ty.vec2<i32>(), {}); MakeStorageTextureBodyFunction("st_func", "st_var", ty.vec2<i32>(), {});
auto* rost_type = MakeStorageTextureTypes(ast::TextureDimension::k2d,
ast::ImageFormat::kR32Uint, true);
AddStorageTexture("rost_var", rost_type, 4, 1);
MakeStorageTextureBodyFunction("rost_func", "rost_var", ty.vec2<i32>(), {});
MakeCallerBodyFunction("ep_func", MakeCallerBodyFunction("ep_func",
{"ub_func", "sb_func", "rosb_func", "s_func", {"ub_func", "sb_func", "rosb_func", "s_func",
"cs_func", "depth_ms_func", "st_func", "rost_func"}, "cs_func", "depth_ms_func", "st_func"},
ast::DecorationList{ ast::DecorationList{
Stage(ast::PipelineStage::kFragment), Stage(ast::PipelineStage::kFragment),
}); });
@ -1287,7 +1282,7 @@ TEST_F(InspectorGetResourceBindingsTest, Simple) {
auto result = inspector.GetResourceBindings("ep_func"); auto result = inspector.GetResourceBindings("ep_func");
ASSERT_FALSE(inspector.has_error()) << inspector.error(); ASSERT_FALSE(inspector.has_error()) << inspector.error();
ASSERT_EQ(10u, result.size()); ASSERT_EQ(9u, result.size());
EXPECT_EQ(ResourceBinding::ResourceType::kUniformBuffer, EXPECT_EQ(ResourceBinding::ResourceType::kUniformBuffer,
result[0].resource_type); result[0].resource_type);
@ -1318,25 +1313,20 @@ TEST_F(InspectorGetResourceBindingsTest, Simple) {
EXPECT_EQ(2u, result[5].bind_group); EXPECT_EQ(2u, result[5].bind_group);
EXPECT_EQ(0u, result[5].binding); EXPECT_EQ(0u, result[5].binding);
EXPECT_EQ(ResourceBinding::ResourceType::kReadOnlyStorageTexture, EXPECT_EQ(ResourceBinding::ResourceType::kWriteOnlyStorageTexture,
result[6].resource_type); result[6].resource_type);
EXPECT_EQ(4u, result[6].bind_group); EXPECT_EQ(4u, result[6].bind_group);
EXPECT_EQ(1u, result[6].binding); EXPECT_EQ(0u, result[6].binding);
EXPECT_EQ(ResourceBinding::ResourceType::kWriteOnlyStorageTexture,
result[7].resource_type);
EXPECT_EQ(4u, result[7].bind_group);
EXPECT_EQ(0u, result[7].binding);
EXPECT_EQ(ResourceBinding::ResourceType::kDepthTexture, EXPECT_EQ(ResourceBinding::ResourceType::kDepthTexture,
result[8].resource_type); result[7].resource_type);
EXPECT_EQ(3u, result[8].bind_group); EXPECT_EQ(3u, result[7].bind_group);
EXPECT_EQ(1u, result[8].binding); EXPECT_EQ(1u, result[7].binding);
EXPECT_EQ(ResourceBinding::ResourceType::kDepthMultisampledTexture, EXPECT_EQ(ResourceBinding::ResourceType::kDepthMultisampledTexture,
result[9].resource_type); result[8].resource_type);
EXPECT_EQ(3u, result[9].bind_group); EXPECT_EQ(3u, result[8].bind_group);
EXPECT_EQ(3u, result[9].binding); EXPECT_EQ(3u, result[8].binding);
} }
TEST_F(InspectorGetUniformBufferResourceBindingsTest, MissingEntryPoint) { TEST_F(InspectorGetUniformBufferResourceBindingsTest, MissingEntryPoint) {
@ -2358,20 +2348,15 @@ TEST_F(InspectorGetStorageTextureResourceBindingsTest, Empty) {
Inspector& inspector = Build(); Inspector& inspector = Build();
auto result = inspector.GetReadOnlyStorageTextureResourceBindings("ep"); auto result = inspector.GetWriteOnlyStorageTextureResourceBindings("ep");
ASSERT_FALSE(inspector.has_error()) << inspector.error();
EXPECT_EQ(0u, result.size());
result = inspector.GetWriteOnlyStorageTextureResourceBindings("ep");
ASSERT_FALSE(inspector.has_error()) << inspector.error(); ASSERT_FALSE(inspector.has_error()) << inspector.error();
EXPECT_EQ(0u, result.size()); EXPECT_EQ(0u, result.size());
} }
TEST_P(InspectorGetStorageTextureResourceBindingsTestWithParam, Simple) { TEST_P(InspectorGetStorageTextureResourceBindingsTestWithParam, Simple) {
bool read_only;
DimensionParams dim_params; DimensionParams dim_params;
ImageFormatParams format_params; ImageFormatParams format_params;
std::tie(read_only, dim_params, format_params) = GetParam(); std::tie(dim_params, format_params) = GetParam();
ast::TextureDimension dim; ast::TextureDimension dim;
ResourceBinding::TextureDimension expected_dim; ResourceBinding::TextureDimension expected_dim;
@ -2382,7 +2367,7 @@ TEST_P(InspectorGetStorageTextureResourceBindingsTestWithParam, Simple) {
ResourceBinding::SampledKind expected_kind; ResourceBinding::SampledKind expected_kind;
std::tie(format, expected_format, expected_kind) = format_params; std::tie(format, expected_format, expected_kind) = format_params;
auto* st_type = MakeStorageTextureTypes(dim, format, read_only); auto* st_type = MakeStorageTextureTypes(dim, format);
AddStorageTexture("st_var", st_type, 0, 0); AddStorageTexture("st_var", st_type, 0, 0);
ast::Type* dim_type = nullptr; ast::Type* dim_type = nullptr;
@ -2409,33 +2394,23 @@ TEST_P(InspectorGetStorageTextureResourceBindingsTestWithParam, Simple) {
Inspector& inspector = Build(); Inspector& inspector = Build();
auto result = auto result = inspector.GetWriteOnlyStorageTextureResourceBindings("ep");
read_only ? inspector.GetReadOnlyStorageTextureResourceBindings("ep")
: inspector.GetWriteOnlyStorageTextureResourceBindings("ep");
ASSERT_FALSE(inspector.has_error()) << inspector.error(); ASSERT_FALSE(inspector.has_error()) << inspector.error();
ASSERT_EQ(1u, result.size()); ASSERT_EQ(1u, result.size());
EXPECT_EQ(read_only ? ResourceBinding::ResourceType::kReadOnlyStorageTexture EXPECT_EQ(ResourceBinding::ResourceType::kWriteOnlyStorageTexture,
: ResourceBinding::ResourceType::kWriteOnlyStorageTexture,
result[0].resource_type); result[0].resource_type);
EXPECT_EQ(0u, result[0].bind_group); EXPECT_EQ(0u, result[0].bind_group);
EXPECT_EQ(0u, result[0].binding); EXPECT_EQ(0u, result[0].binding);
EXPECT_EQ(expected_dim, result[0].dim); EXPECT_EQ(expected_dim, result[0].dim);
EXPECT_EQ(expected_format, result[0].image_format); EXPECT_EQ(expected_format, result[0].image_format);
EXPECT_EQ(expected_kind, result[0].sampled_kind); EXPECT_EQ(expected_kind, result[0].sampled_kind);
result = read_only
? inspector.GetWriteOnlyStorageTextureResourceBindings("ep")
: inspector.GetReadOnlyStorageTextureResourceBindings("ep");
ASSERT_FALSE(inspector.has_error()) << inspector.error();
ASSERT_EQ(0u, result.size());
} }
INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P(
InspectorGetStorageTextureResourceBindingsTest, InspectorGetStorageTextureResourceBindingsTest,
InspectorGetStorageTextureResourceBindingsTestWithParam, InspectorGetStorageTextureResourceBindingsTestWithParam,
testing::Combine( testing::Combine(
testing::Bool(),
testing::Values( testing::Values(
std::make_tuple(ast::TextureDimension::k1d, std::make_tuple(ast::TextureDimension::k1d,
ResourceBinding::TextureDimension::k1d), ResourceBinding::TextureDimension::k1d),

1
src/inspector/resource_binding.h
View File

@ -96,7 +96,6 @@ struct ResourceBinding {
kComparisonSampler, kComparisonSampler,
kSampledTexture, kSampledTexture,
kMultisampledTexture, kMultisampledTexture,
kReadOnlyStorageTexture,
kWriteOnlyStorageTexture, kWriteOnlyStorageTexture,
kDepthTexture, kDepthTexture,
kDepthMultisampledTexture, kDepthMultisampledTexture,

6
src/inspector/test_inspector_builder.cc
View File

@ -318,10 +318,8 @@ ast::Type* InspectorBuilder::GetCoordsType(ast::TextureDimension dim,
} }
ast::Type* InspectorBuilder::MakeStorageTextureTypes(ast::TextureDimension dim, ast::Type* InspectorBuilder::MakeStorageTextureTypes(ast::TextureDimension dim,
ast::ImageFormat format, ast::ImageFormat format) {
bool read_only) { return ty.storage_texture(dim, format, ast::Access::kWrite);
auto access = read_only ? ast::Access::kRead : ast::Access::kWrite;
return ty.storage_texture(dim, format, access);
} }
void InspectorBuilder::AddStorageTexture(const std::string& name, void InspectorBuilder::AddStorageTexture(const std::string& name,

4
src/inspector/test_inspector_builder.h
View File

@ -329,11 +329,9 @@ class InspectorBuilder : public ProgramBuilder {
/// Generates appropriate types for a Read-Only StorageTexture /// Generates appropriate types for a Read-Only StorageTexture
/// @param dim the texture dimension of the storage texture /// @param dim the texture dimension of the storage texture
/// @param format the image format of the storage texture /// @param format the image format of the storage texture
/// @param read_only should the access type be read only, otherwise write only
/// @returns the storage texture type /// @returns the storage texture type
ast::Type* MakeStorageTextureTypes(ast::TextureDimension dim, ast::Type* MakeStorageTextureTypes(ast::TextureDimension dim,
ast::ImageFormat format, ast::ImageFormat format);
bool read_only);
/// Adds a storage texture variable to the program /// Adds a storage texture variable to the program
/// @param name the name of the variable /// @param name the name of the variable

6343
src/intrinsic_table.inl
View File

File diff suppressed because it is too large Load Diff

40
src/intrinsic_table_test.cc
View File

@ -389,30 +389,6 @@ TEST_F(IntrinsicTableTest, MatchExternalTexture) {
EXPECT_EQ(result->Parameters()[1]->Usage(), ParameterUsage::kCoords); EXPECT_EQ(result->Parameters()[1]->Usage(), ParameterUsage::kCoords);
} }
TEST_F(IntrinsicTableTest, MatchROStorageTexture) {
auto* f32 = create<sem::F32>();
auto* i32 = create<sem::I32>();
auto* vec2_i32 = create<sem::Vector>(i32, 2);
auto* vec4_f32 = create<sem::Vector>(f32, 4);
auto* subtype =
sem::StorageTexture::SubtypeFor(ast::ImageFormat::kR32Float, Types());
auto* tex = create<sem::StorageTexture>(ast::TextureDimension::k2d,
ast::ImageFormat::kR32Float,
ast::Access::kRead, subtype);
auto* result =
table->Lookup(IntrinsicType::kTextureLoad, {tex, vec2_i32}, Source{});
ASSERT_NE(result, nullptr) << Diagnostics().str();
ASSERT_EQ(Diagnostics().str(), "");
EXPECT_THAT(result->Type(), IntrinsicType::kTextureLoad);
EXPECT_THAT(result->ReturnType(), vec4_f32);
ASSERT_EQ(result->Parameters().size(), 2u);
EXPECT_EQ(result->Parameters()[0]->Type(), tex);
EXPECT_EQ(result->Parameters()[0]->Usage(), ParameterUsage::kTexture);
EXPECT_EQ(result->Parameters()[1]->Type(), vec2_i32);
EXPECT_EQ(result->Parameters()[1]->Usage(), ParameterUsage::kCoords);
}
TEST_F(IntrinsicTableTest, MatchWOStorageTexture) { TEST_F(IntrinsicTableTest, MatchWOStorageTexture) {
auto* f32 = create<sem::F32>(); auto* f32 = create<sem::F32>();
auto* i32 = create<sem::I32>(); auto* i32 = create<sem::I32>();
@ -566,10 +542,10 @@ TEST_F(IntrinsicTableTest, OverloadOrderByNumberOfParameters) {
textureDimensions(texture: texture_depth_cube) -> vec2<i32> textureDimensions(texture: texture_depth_cube) -> vec2<i32>
textureDimensions(texture: texture_depth_cube_array) -> vec2<i32> textureDimensions(texture: texture_depth_cube_array) -> vec2<i32>
textureDimensions(texture: texture_depth_multisampled_2d) -> vec2<i32> textureDimensions(texture: texture_depth_multisampled_2d) -> vec2<i32>
textureDimensions(texture: texture_storage_1d<F, A>) -> i32 where: A is read or write textureDimensions(texture: texture_storage_1d<F, A>) -> i32 where: A is write
textureDimensions(texture: texture_storage_2d<F, A>) -> vec2<i32> where: A is read or 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 read or 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 read or write textureDimensions(texture: texture_storage_3d<F, A>) -> vec3<i32> where: A is write
textureDimensions(texture: texture_external) -> vec2<i32> textureDimensions(texture: texture_external) -> vec2<i32>
)"); )");
} }
@ -605,10 +581,10 @@ TEST_F(IntrinsicTableTest, OverloadOrderByMatchingParameter) {
textureDimensions(texture: texture_depth_cube) -> vec2<i32> textureDimensions(texture: texture_depth_cube) -> vec2<i32>
textureDimensions(texture: texture_depth_cube_array) -> vec2<i32> textureDimensions(texture: texture_depth_cube_array) -> vec2<i32>
textureDimensions(texture: texture_depth_multisampled_2d) -> vec2<i32> textureDimensions(texture: texture_depth_multisampled_2d) -> vec2<i32>
textureDimensions(texture: texture_storage_1d<F, A>) -> i32 where: A is read or write textureDimensions(texture: texture_storage_1d<F, A>) -> i32 where: A is write
textureDimensions(texture: texture_storage_2d<F, A>) -> vec2<i32> where: A is read or 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 read or 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 read or write textureDimensions(texture: texture_storage_3d<F, A>) -> vec3<i32> where: A is write
textureDimensions(texture: texture_external) -> vec2<i32> textureDimensions(texture: texture_external) -> vec2<i32>
)"); )");
} }

24
src/intrinsics.def
View File

@ -128,7 +128,7 @@ match i32_texel_format:
match u32_texel_format: match u32_texel_format:
rgba8uint | rgba16uint | r32uint | rg32uint | rgba32uint rgba8uint | rgba16uint | r32uint | rg32uint | rgba32uint
match read_or_write: read | write match write_only: write
match function_private_workgroup: function | private | workgroup match function_private_workgroup: function | private | workgroup
match workgroup_or_storage: workgroup | storage match workgroup_or_storage: workgroup | storage
@ -416,16 +416,16 @@ fn textureDimensions(texture: texture_depth_cube, level: i32) -> vec2<i32>
fn textureDimensions(texture: texture_depth_cube_array) -> vec2<i32> fn textureDimensions(texture: texture_depth_cube_array) -> vec2<i32>
fn textureDimensions(texture: texture_depth_cube_array, level: i32) -> vec2<i32> fn textureDimensions(texture: texture_depth_cube_array, level: i32) -> vec2<i32>
fn textureDimensions(texture: texture_depth_multisampled_2d) -> vec2<i32> fn textureDimensions(texture: texture_depth_multisampled_2d) -> vec2<i32>
fn textureDimensions<F: texel_format, A: read_or_write>(texture: texture_storage_1d<F, A>) -> i32 fn textureDimensions<F: texel_format, A: write_only>(texture: texture_storage_1d<F, A>) -> i32
fn textureDimensions<F: texel_format, A: read_or_write>(texture: texture_storage_2d<F, A>) -> vec2<i32> fn textureDimensions<F: texel_format, A: write_only>(texture: texture_storage_2d<F, A>) -> vec2<i32>
fn textureDimensions<F: texel_format, A: read_or_write>(texture: texture_storage_2d_array<F, A>) -> vec2<i32> fn textureDimensions<F: texel_format, A: write_only>(texture: texture_storage_2d_array<F, A>) -> vec2<i32>
fn textureDimensions<F: texel_format, A: read_or_write>(texture: texture_storage_3d<F, A>) -> vec3<i32> fn textureDimensions<F: texel_format, A: write_only>(texture: texture_storage_3d<F, A>) -> vec3<i32>
fn textureDimensions(texture: texture_external) -> vec2<i32> fn textureDimensions(texture: texture_external) -> vec2<i32>
fn textureNumLayers<T: fiu32>(texture: texture_2d_array<T>) -> i32 fn textureNumLayers<T: fiu32>(texture: texture_2d_array<T>) -> i32
fn textureNumLayers<T: fiu32>(texture: texture_cube_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_2d_array) -> i32
fn textureNumLayers(texture: texture_depth_cube_array) -> i32 fn textureNumLayers(texture: texture_depth_cube_array) -> i32
fn textureNumLayers<F: texel_format, A: read_or_write>(texture: texture_storage_2d_array<F, A>) -> i32 fn textureNumLayers<F: texel_format, A: write_only>(texture: texture_storage_2d_array<F, A>) -> i32
fn textureNumLevels<T: fiu32>(texture: texture_1d<T>) -> 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<T>) -> i32
fn textureNumLevels<T: fiu32>(texture: texture_2d_array<T>) -> i32 fn textureNumLevels<T: fiu32>(texture: texture_2d_array<T>) -> i32
@ -516,18 +516,6 @@ fn textureLoad<T: fiu32>(texture: texture_multisampled_2d<T>, coords: vec2<i32>,
fn textureLoad(texture: texture_depth_2d, coords: vec2<i32>, level: i32) -> f32 fn textureLoad(texture: texture_depth_2d, coords: vec2<i32>, level: i32) -> f32
fn textureLoad(texture: texture_depth_2d_array, coords: vec2<i32>, array_index: i32, level: i32) -> f32 fn textureLoad(texture: texture_depth_2d_array, coords: vec2<i32>, array_index: i32, level: i32) -> f32
fn textureLoad(texture: texture_depth_multisampled_2d, coords: vec2<i32>, sample_index: i32) -> f32 fn textureLoad(texture: texture_depth_multisampled_2d, coords: vec2<i32>, sample_index: i32) -> f32
[[deprecated]] fn textureLoad(texture: texture_storage_1d<f32_texel_format, read>, coords: i32) -> vec4<f32>
[[deprecated]] fn textureLoad(texture: texture_storage_2d<f32_texel_format, read>, coords: vec2<i32>) -> vec4<f32>
[[deprecated]] fn textureLoad(texture: texture_storage_2d_array<f32_texel_format, read>, coords: vec2<i32>, array_index: i32) -> vec4<f32>
[[deprecated]] fn textureLoad(texture: texture_storage_3d<f32_texel_format, read>, coords: vec3<i32>) -> vec4<f32>
[[deprecated]] fn textureLoad(texture: texture_storage_1d<i32_texel_format, read>, coords: i32) -> vec4<i32>
[[deprecated]] fn textureLoad(texture: texture_storage_2d<i32_texel_format, read>, coords: vec2<i32>) -> vec4<i32>
[[deprecated]] fn textureLoad(texture: texture_storage_2d_array<i32_texel_format, read>, coords: vec2<i32>, array_index: i32) -> vec4<i32>
[[deprecated]] fn textureLoad(texture: texture_storage_3d<i32_texel_format, read>, coords: vec3<i32>) -> vec4<i32>
[[deprecated]] fn textureLoad(texture: texture_storage_1d<u32_texel_format, read>, coords: i32) -> vec4<u32>
[[deprecated]] fn textureLoad(texture: texture_storage_2d<u32_texel_format, read>, coords: vec2<i32>) -> vec4<u32>
[[deprecated]] fn textureLoad(texture: texture_storage_2d_array<u32_texel_format, read>, coords: vec2<i32>, array_index: i32) -> vec4<u32>
[[deprecated]] fn textureLoad(texture: texture_storage_3d<u32_texel_format, read>, coords: vec3<i32>) -> vec4<u32>
fn textureLoad(texture: texture_external, coords: vec2<i32>) -> vec4<f32> fn textureLoad(texture: texture_external, coords: vec2<i32>) -> vec4<f32>
[[stage("fragment", "compute")]] fn atomicLoad<T: iu32, S: workgroup_or_storage>(ptr<S, atomic<T>, read_write>) -> T [[stage("fragment", "compute")]] fn atomicLoad<T: iu32, S: workgroup_or_storage>(ptr<S, atomic<T>, read_write>) -> T

16
src/reader/spirv/function.cc
View File

@ -5288,12 +5288,8 @@ bool FunctionEmitter::EmitImageAccess(const spvtools::opt::Instruction& inst) {
case SpvOpImageDrefGather: case SpvOpImageDrefGather:
return Fail() << " image gather is not yet supported"; return Fail() << " image gather is not yet supported";
case SpvOpImageFetch: case SpvOpImageFetch:
// Read a single texel from a sampled image.
builtin_name = "textureLoad";
use_level_of_detail_suffix = false;
break;
case SpvOpImageRead: case SpvOpImageRead:
// Read a single texel from a storage image. // Read a single texel from a sampled or storage image.
builtin_name = "textureLoad"; builtin_name = "textureLoad";
use_level_of_detail_suffix = false; use_level_of_detail_suffix = false;
break; break;
@ -5365,11 +5361,11 @@ bool FunctionEmitter::EmitImageAccess(const spvtools::opt::Instruction& inst) {
image_operands_mask ^= SpvImageOperandsLodMask; image_operands_mask ^= SpvImageOperandsLodMask;
arg_index++; arg_index++;
} else if ((opcode == SpvOpImageFetch) && } else if ((opcode == SpvOpImageFetch || opcode == SpvOpImageRead) &&
(texture_type->Is<SampledTexture>() || !texture_type
texture_type->Is<DepthTexture>())) { ->IsAnyOf<DepthMultisampledTexture, MultisampledTexture>()) {
// textureLoad on sampled texture and depth texture requires an explicit // textureLoad requires an explicit level-of-detail parameter for
// level-of-detail parameter. // non-multisampled texture types.
params.push_back(parser_impl_.MakeNullValue(ty_.I32())); params.push_back(parser_impl_.MakeNullValue(ty_.I32()));
} }
if (arg_index + 1 < num_args && if (arg_index + 1 < num_args &&

5
src/reader/spirv/parser_impl.cc
View File

@ -2537,7 +2537,7 @@ const Pointer* ParserImpl::GetTypeForHandleVar(
// WGSL textures are always formatted. Unformatted textures are always // WGSL textures are always formatted. Unformatted textures are always
// sampled. // sampled.
if (usage.IsSampledTexture() || if (usage.IsSampledTexture() || usage.IsStorageReadTexture() ||
(image_type->format() == SpvImageFormatUnknown)) { (image_type->format() == SpvImageFormatUnknown)) {
// Make a sampled texture type. // Make a sampled texture type.
auto* ast_sampled_component_type = auto* ast_sampled_component_type =
@ -2566,8 +2566,7 @@ const Pointer* ParserImpl::GetTypeForHandleVar(
ast_store_type = ty_.SampledTexture(dim, ast_sampled_component_type); ast_store_type = ty_.SampledTexture(dim, ast_sampled_component_type);
} }
} else { } else {
const auto access = usage.IsStorageReadTexture() ? ast::Access::kRead const auto access = ast::Access::kWrite;
: ast::Access::kWrite;
const auto format = enum_converter_.ToImageFormat(image_type->format()); const auto format = enum_converter_.ToImageFormat(image_type->format());
if (format == ast::ImageFormat::kNone) { if (format == ast::ImageFormat::kNone) {
return nullptr; return nullptr;

2
src/reader/spirv/parser_impl_convert_type_test.cc
View File

@ -695,7 +695,7 @@ TEST_F(SpvParserTest, ConvertType_Array_NoDeduplication) {
TEST_F(SpvParserTest, ConvertType_RuntimeArray_NoDeduplication) { TEST_F(SpvParserTest, ConvertType_RuntimeArray_NoDeduplication) {
// Prove that distinct SPIR-V runtime arrays map to distinct WGSL types. // Prove that distinct SPIR-V runtime arrays map to distinct WGSL types.
// The implementation already deduplciates them because it knows // The implementation already de-duplicates them because it knows
// runtime-arrays normally have stride decorations. // runtime-arrays normally have stride decorations.
auto assembly = Preamble() + R"( auto assembly = Preamble() + R"(
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0

22
src/reader/spirv/parser_impl_handle_test.cc
View File

@ -1309,7 +1309,7 @@ INSTANTIATE_TEST_SUITE_P(Images,
x_10 x_10
none none
undefined undefined
__storage_texture_1d_rg32float_read __sampled_texture_1d__f32
})"}, })"},
DeclUnderspecifiedHandleCase{R"( DeclUnderspecifiedHandleCase{R"(
OpDecorate %10 NonReadable OpDecorate %10 NonReadable
@ -4615,7 +4615,7 @@ INSTANTIATE_TEST_SUITE_P(ImageRead_OptionalParams,
x_20 x_20
none none
undefined undefined
__storage_texture_2d_rgba32float_read __sampled_texture_2d__f32
})", })",
R"(VariableDeclStatement{ R"(VariableDeclStatement{
VariableConst{ VariableConst{
@ -4629,6 +4629,7 @@ INSTANTIATE_TEST_SUITE_P(ImageRead_OptionalParams,
( (
Identifier[not set]{x_20} Identifier[not set]{x_20}
Identifier[not set]{vi12} Identifier[not set]{vi12}
ScalarConstructor[not set]{0}
) )
} }
} }
@ -5079,7 +5080,7 @@ INSTANTIATE_TEST_SUITE_P(
x_20 x_20
none none
undefined undefined
__storage_texture_2d_rgba32float_read __sampled_texture_2d__f32
})", })",
R"(VariableDeclStatement{ R"(VariableDeclStatement{
VariableConst{ VariableConst{
@ -5093,6 +5094,7 @@ INSTANTIATE_TEST_SUITE_P(
( (
Identifier[not set]{x_20} Identifier[not set]{x_20}
Identifier[not set]{vi12} Identifier[not set]{vi12}
ScalarConstructor[not set]{0}
) )
} }
} }
@ -5108,7 +5110,7 @@ INSTANTIATE_TEST_SUITE_P(
x_20 x_20
none none
undefined undefined
__storage_texture_2d_rgba32uint_read __sampled_texture_2d__u32
})", })",
R"(VariableDeclStatement{ R"(VariableDeclStatement{
VariableConst{ VariableConst{
@ -5122,6 +5124,7 @@ INSTANTIATE_TEST_SUITE_P(
( (
Identifier[not set]{x_20} Identifier[not set]{x_20}
Identifier[not set]{vi12} Identifier[not set]{vi12}
ScalarConstructor[not set]{0}
) )
} }
} }
@ -5143,7 +5146,7 @@ INSTANTIATE_TEST_SUITE_P(
x_20 x_20
none none
undefined undefined
__storage_texture_2d_rgba32sint_read __sampled_texture_2d__i32
})", })",
R"(VariableDeclStatement{ R"(VariableDeclStatement{
VariableConst{ VariableConst{
@ -5157,6 +5160,7 @@ INSTANTIATE_TEST_SUITE_P(
( (
Identifier[not set]{x_20} Identifier[not set]{x_20}
Identifier[not set]{vi12} Identifier[not set]{vi12}
ScalarConstructor[not set]{0}
) )
} }
} }
@ -5237,7 +5241,7 @@ INSTANTIATE_TEST_SUITE_P(
x_20 x_20
none none
undefined undefined
__storage_texture_1d_rgba32float_read __sampled_texture_1d__f32
})", })",
R"(VariableDeclStatement{ R"(VariableDeclStatement{
VariableConst{ VariableConst{
@ -5271,7 +5275,7 @@ INSTANTIATE_TEST_SUITE_P(
x_20 x_20
none none
undefined undefined
__storage_texture_2d_rgba32float_read __sampled_texture_2d__f32
})", })",
R"(VariableDeclStatement{ R"(VariableDeclStatement{
VariableConst{ VariableConst{
@ -5305,7 +5309,7 @@ INSTANTIATE_TEST_SUITE_P(
x_20 x_20
none none
undefined undefined
__storage_texture_3d_rgba32float_read __sampled_texture_3d__f32
})", })",
R"(VariableDeclStatement{ R"(VariableDeclStatement{
VariableConst{ VariableConst{
@ -5380,7 +5384,7 @@ INSTANTIATE_TEST_SUITE_P(
x_20 x_20
none none
undefined undefined
__storage_texture_2d_array_rgba32float_read __sampled_texture_2d_array__f32
})", })",
R"(VariableDeclStatement{ R"(VariableDeclStatement{
VariableConst{ VariableConst{

6
src/resolver/assignment_validation_test.cc
View File

@ -221,14 +221,14 @@ TEST_F(ResolverAssignmentValidationTest, AssignToConstant_Fail) {
} }
TEST_F(ResolverAssignmentValidationTest, AssignNonConstructible_Handle) { TEST_F(ResolverAssignmentValidationTest, AssignNonConstructible_Handle) {
// var a : texture_storage_1d<rgba8unorm, read>; // var a : texture_storage_1d<rgba8unorm, write>;
// var b : texture_storage_1d<rgba8unorm, read>; // var b : texture_storage_1d<rgba8unorm, write>;
// a = b; // a = b;
auto make_type = [&] { auto make_type = [&] {
return ty.storage_texture(ast::TextureDimension::k1d, return ty.storage_texture(ast::TextureDimension::k1d,
ast::ImageFormat::kRgba8Unorm, ast::ImageFormat::kRgba8Unorm,
ast::Access::kRead); ast::Access::kWrite);
}; };
Global("a", make_type(), ast::StorageClass::kNone, Global("a", make_type(), ast::StorageClass::kNone,

91
src/resolver/intrinsic_test.cc
View File

@ -268,71 +268,6 @@ class ResolverIntrinsicTest_TextureOperation
} }
}; };
using ResolverIntrinsicTest_StorageTextureOperation =
ResolverIntrinsicTest_TextureOperation;
TEST_P(ResolverIntrinsicTest_StorageTextureOperation, TextureLoadRo) {
auto dim = GetParam().dim;
auto type = GetParam().type;
auto format = GetParam().format;
auto* coords_type = GetCoordsType(dim, ty.i32());
auto* texture_type = ty.storage_texture(dim, format, ast::Access::kRead);
ast::ExpressionList call_params;
add_call_param("texture", texture_type, &call_params);
add_call_param("coords", coords_type, &call_params);
if (ast::IsTextureArray(dim)) {
add_call_param("array_index", ty.i32(), &call_params);
}
auto* expr = Call("textureLoad", call_params);
WrapInFunction(expr);
EXPECT_TRUE(r()->Resolve()) << r()->error();
ASSERT_NE(TypeOf(expr), nullptr);
ASSERT_TRUE(TypeOf(expr)->Is<sem::Vector>());
if (type == Texture::kF32) {
EXPECT_TRUE(TypeOf(expr)->As<sem::Vector>()->type()->Is<sem::F32>());
} else if (type == Texture::kI32) {
EXPECT_TRUE(TypeOf(expr)->As<sem::Vector>()->type()->Is<sem::I32>());
} else {
EXPECT_TRUE(TypeOf(expr)->As<sem::Vector>()->type()->Is<sem::U32>());
}
EXPECT_EQ(TypeOf(expr)->As<sem::Vector>()->Width(), 4u);
}
INSTANTIATE_TEST_SUITE_P(
ResolverTest,
ResolverIntrinsicTest_StorageTextureOperation,
testing::Values(
TextureTestParams{ast::TextureDimension::k1d, Texture::kF32,
ast::ImageFormat::kR32Float},
TextureTestParams{ast::TextureDimension::k1d, Texture::kI32,
ast::ImageFormat::kR32Sint},
TextureTestParams{ast::TextureDimension::k1d, Texture::kF32,
ast::ImageFormat::kRgba8Unorm},
TextureTestParams{ast::TextureDimension::k2d, Texture::kF32,
ast::ImageFormat::kR32Float},
TextureTestParams{ast::TextureDimension::k2d, Texture::kI32,
ast::ImageFormat::kR32Sint},
TextureTestParams{ast::TextureDimension::k2d, Texture::kF32,
ast::ImageFormat::kRgba8Unorm},
TextureTestParams{ast::TextureDimension::k2dArray, Texture::kF32,
ast::ImageFormat::kR32Float},
TextureTestParams{ast::TextureDimension::k2dArray, Texture::kI32,
ast::ImageFormat::kR32Sint},
TextureTestParams{ast::TextureDimension::k2dArray, Texture::kF32,
ast::ImageFormat::kRgba8Unorm},
TextureTestParams{ast::TextureDimension::k3d, Texture::kF32,
ast::ImageFormat::kR32Float},
TextureTestParams{ast::TextureDimension::k3d, Texture::kI32,
ast::ImageFormat::kR32Sint},
TextureTestParams{ast::TextureDimension::k3d, Texture::kF32,
ast::ImageFormat::kRgba8Unorm}));
using ResolverIntrinsicTest_SampledTextureOperation = using ResolverIntrinsicTest_SampledTextureOperation =
ResolverIntrinsicTest_TextureOperation; ResolverIntrinsicTest_TextureOperation;
TEST_P(ResolverIntrinsicTest_SampledTextureOperation, TextureLoadSampled) { TEST_P(ResolverIntrinsicTest_SampledTextureOperation, TextureLoadSampled) {
@ -1806,10 +1741,6 @@ const char* expected_texture_overload(
case ValidTextureOverload::kDimensionsDepthCube: case ValidTextureOverload::kDimensionsDepthCube:
case ValidTextureOverload::kDimensionsDepthCubeArray: case ValidTextureOverload::kDimensionsDepthCubeArray:
case ValidTextureOverload::kDimensionsDepthMultisampled2d: case ValidTextureOverload::kDimensionsDepthMultisampled2d:
case ValidTextureOverload::kDimensionsStorageRO1d:
case ValidTextureOverload::kDimensionsStorageRO2d:
case ValidTextureOverload::kDimensionsStorageRO2dArray:
case ValidTextureOverload::kDimensionsStorageRO3d:
case ValidTextureOverload::kDimensionsStorageWO1d: case ValidTextureOverload::kDimensionsStorageWO1d:
case ValidTextureOverload::kDimensionsStorageWO2d: case ValidTextureOverload::kDimensionsStorageWO2d:
case ValidTextureOverload::kDimensionsStorageWO2dArray: case ValidTextureOverload::kDimensionsStorageWO2dArray:
@ -1981,28 +1912,6 @@ const char* expected_texture_overload(
return R"(textureLoad(texture, coords, sample_index))"; return R"(textureLoad(texture, coords, sample_index))";
case ValidTextureOverload::kLoadDepth2dArrayLevelF32: case ValidTextureOverload::kLoadDepth2dArrayLevelF32:
return R"(textureLoad(texture, coords, array_index, level))"; return R"(textureLoad(texture, coords, array_index, level))";
case ValidTextureOverload::kLoadStorageRO1dRgba32float:
case ValidTextureOverload::kLoadStorageRO2dRgba8unorm:
case ValidTextureOverload::kLoadStorageRO2dRgba8snorm:
case ValidTextureOverload::kLoadStorageRO2dRgba8uint:
case ValidTextureOverload::kLoadStorageRO2dRgba8sint:
case ValidTextureOverload::kLoadStorageRO2dRgba16uint:
case ValidTextureOverload::kLoadStorageRO2dRgba16sint:
case ValidTextureOverload::kLoadStorageRO2dRgba16float:
case ValidTextureOverload::kLoadStorageRO2dR32uint:
case ValidTextureOverload::kLoadStorageRO2dR32sint:
case ValidTextureOverload::kLoadStorageRO2dR32float:
case ValidTextureOverload::kLoadStorageRO2dRg32uint:
case ValidTextureOverload::kLoadStorageRO2dRg32sint:
case ValidTextureOverload::kLoadStorageRO2dRg32float:
case ValidTextureOverload::kLoadStorageRO2dRgba32uint:
case ValidTextureOverload::kLoadStorageRO2dRgba32sint:
case ValidTextureOverload::kLoadStorageRO2dRgba32float:
return R"(textureLoad(texture, coords))";
case ValidTextureOverload::kLoadStorageRO2dArrayRgba32float:
return R"(textureLoad(texture, coords, array_index))";
case ValidTextureOverload::kLoadStorageRO3dRgba32float:
return R"(textureLoad(texture, coords))";
case ValidTextureOverload::kStoreWO1dRgba32float: case ValidTextureOverload::kStoreWO1dRgba32float:
case ValidTextureOverload::kStoreWO2dRgba32float: case ValidTextureOverload::kStoreWO2dRgba32float:
case ValidTextureOverload::kStoreWO3dRgba32float: case ValidTextureOverload::kStoreWO3dRgba32float:

16
src/resolver/resolver.cc
View File

@ -424,17 +424,15 @@ bool Resolver::ValidateAtomic(const ast::Atomic* a, const sem::Atomic* s) {
bool Resolver::ValidateStorageTexture(const ast::StorageTexture* t) { bool Resolver::ValidateStorageTexture(const ast::StorageTexture* t) {
switch (t->access()) { switch (t->access()) {
case ast::Access::kUndefined:
AddError("storage textures must have access control", t->source());
return false;
case ast::Access::kReadWrite:
AddError("storage textures only support read-only and write-only access",
t->source());
return false;
case ast::Access::kRead:
case ast::Access::kWrite: case ast::Access::kWrite:
break; break;
case ast::Access::kUndefined:
AddError("storage texture missing access control", t->source());
return false;
default:
AddError("storage textures currently only support 'write' access control",
t->source());
return false;
} }
if (!IsValidStorageTextureDimension(t->dim())) { if (!IsValidStorageTextureDimension(t->dim())) {

93
src/resolver/type_validation_test.cc
View File

@ -698,10 +698,7 @@ TEST_P(MultisampledTextureDimensionTest, All) {
auto& params = GetParam(); auto& params = GetParam();
Global(Source{{12, 34}}, "a", ty.multisampled_texture(params.dim, ty.i32()), Global(Source{{12, 34}}, "a", ty.multisampled_texture(params.dim, ty.i32()),
ast::StorageClass::kNone, nullptr, ast::StorageClass::kNone, nullptr,
ast::DecorationList{ ast::DecorationList{GroupAndBinding(0, 0)});
create<ast::BindingDecoration>(0),
create<ast::GroupDecoration>(0),
});
if (params.is_valid) { if (params.is_valid) {
EXPECT_TRUE(r()->Resolve()) << r()->error(); EXPECT_TRUE(r()->Resolve()) << r()->error();
@ -750,10 +747,7 @@ TEST_P(MultisampledTextureTypeTest, All) {
ty.multisampled_texture(ast::TextureDimension::k2d, ty.multisampled_texture(ast::TextureDimension::k2d,
params.type_func(*this)), params.type_func(*this)),
ast::StorageClass::kNone, nullptr, ast::StorageClass::kNone, nullptr,
ast::DecorationList{ ast::DecorationList{GroupAndBinding(0, 0)});
create<ast::BindingDecoration>(0),
create<ast::GroupDecoration>(0),
});
if (params.is_valid) { if (params.is_valid) {
EXPECT_TRUE(r()->Resolve()) << r()->error(); EXPECT_TRUE(r()->Resolve()) << r()->error();
@ -787,17 +781,15 @@ static constexpr DimensionParams Dimension_cases[] = {
using StorageTextureDimensionTest = ResolverTestWithParam<DimensionParams>; using StorageTextureDimensionTest = ResolverTestWithParam<DimensionParams>;
TEST_P(StorageTextureDimensionTest, All) { TEST_P(StorageTextureDimensionTest, All) {
// [[group(0), binding(0)]] // [[group(0), binding(0)]]
// var a : texture_storage_*<ru32int, read>; // var a : texture_storage_*<ru32int, write>;
auto& params = GetParam(); auto& params = GetParam();
auto* st = ty.storage_texture(Source{{12, 34}}, params.dim, auto* st =
ast::ImageFormat::kR32Uint, ast::Access::kRead); ty.storage_texture(Source{{12, 34}}, params.dim,
ast::ImageFormat::kR32Uint, ast::Access::kWrite);
Global("a", st, ast::StorageClass::kNone, Global("a", st, ast::StorageClass::kNone,
ast::DecorationList{ ast::DecorationList{GroupAndBinding(0, 0)});
create<ast::BindingDecoration>(0),
create<ast::GroupDecoration>(0),
});
if (params.is_valid) { if (params.is_valid) {
EXPECT_TRUE(r()->Resolve()) << r()->error(); EXPECT_TRUE(r()->Resolve()) << r()->error();
@ -858,45 +850,33 @@ using StorageTextureFormatTest = ResolverTestWithParam<FormatParams>;
TEST_P(StorageTextureFormatTest, All) { TEST_P(StorageTextureFormatTest, All) {
auto& params = GetParam(); auto& params = GetParam();
// [[group(0), binding(0)]] // [[group(0), binding(0)]]
// var a : texture_storage_1d<*, read>; // var a : texture_storage_1d<*, write>;
// [[group(0), binding(1)]] // [[group(0), binding(1)]]
// var b : texture_storage_2d<*, read>; // var b : texture_storage_2d<*, write>;
// [[group(0), binding(2)]] // [[group(0), binding(2)]]
// var c : texture_storage_2d_array<*, read>; // var c : texture_storage_2d_array<*, write>;
// [[group(0), binding(3)]] // [[group(0), binding(3)]]
// var d : texture_storage_3d<*, read>; // var d : texture_storage_3d<*, write>;
auto* st_a = ty.storage_texture(Source{{12, 34}}, ast::TextureDimension::k1d, auto* st_a = ty.storage_texture(Source{{12, 34}}, ast::TextureDimension::k1d,
params.format, ast::Access::kRead); params.format, ast::Access::kWrite);
Global("a", st_a, ast::StorageClass::kNone, Global("a", st_a, ast::StorageClass::kNone,
ast::DecorationList{ ast::DecorationList{GroupAndBinding(0, 0)});
create<ast::BindingDecoration>(0),
create<ast::GroupDecoration>(0),
});
auto* st_b = ty.storage_texture(ast::TextureDimension::k2d, params.format, auto* st_b = ty.storage_texture(ast::TextureDimension::k2d, params.format,
ast::Access::kRead); ast::Access::kWrite);
Global("b", st_b, ast::StorageClass::kNone, Global("b", st_b, ast::StorageClass::kNone,
ast::DecorationList{ ast::DecorationList{GroupAndBinding(0, 1)});
create<ast::BindingDecoration>(0),
create<ast::GroupDecoration>(1),
});
auto* st_c = ty.storage_texture(ast::TextureDimension::k2dArray, auto* st_c = ty.storage_texture(ast::TextureDimension::k2dArray,
params.format, ast::Access::kRead); params.format, ast::Access::kWrite);
Global("c", st_c, ast::StorageClass::kNone, Global("c", st_c, ast::StorageClass::kNone,
ast::DecorationList{ ast::DecorationList{GroupAndBinding(0, 2)});
create<ast::BindingDecoration>(0),
create<ast::GroupDecoration>(2),
});
auto* st_d = ty.storage_texture(ast::TextureDimension::k3d, params.format, auto* st_d = ty.storage_texture(ast::TextureDimension::k3d, params.format,
ast::Access::kRead); ast::Access::kWrite);
Global("d", st_d, ast::StorageClass::kNone, Global("d", st_d, ast::StorageClass::kNone,
ast::DecorationList{ ast::DecorationList{GroupAndBinding(0, 3)});
create<ast::BindingDecoration>(0),
create<ast::GroupDecoration>(3),
});
if (params.is_valid) { if (params.is_valid) {
EXPECT_TRUE(r()->Resolve()) << r()->error(); EXPECT_TRUE(r()->Resolve()) << r()->error();
@ -923,14 +903,11 @@ TEST_F(StorageTextureAccessTest, MissingAccess_Fail) {
ast::ImageFormat::kR32Uint, ast::Access::kUndefined); ast::ImageFormat::kR32Uint, ast::Access::kUndefined);
Global("a", st, ast::StorageClass::kNone, Global("a", st, ast::StorageClass::kNone,
ast::DecorationList{ ast::DecorationList{GroupAndBinding(0, 0)});
create<ast::BindingDecoration>(0),
create<ast::GroupDecoration>(0),
});
EXPECT_FALSE(r()->Resolve()); EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(), EXPECT_EQ(r()->error(),
"12:34 error: storage textures must have access control"); "12:34 error: storage texture missing access control");
} }
TEST_F(StorageTextureAccessTest, RWAccess_Fail) { TEST_F(StorageTextureAccessTest, RWAccess_Fail) {
@ -942,31 +919,28 @@ TEST_F(StorageTextureAccessTest, RWAccess_Fail) {
ast::ImageFormat::kR32Uint, ast::Access::kReadWrite); ast::ImageFormat::kR32Uint, ast::Access::kReadWrite);
Global("a", st, ast::StorageClass::kNone, nullptr, Global("a", st, ast::StorageClass::kNone, nullptr,
ast::DecorationList{ ast::DecorationList{GroupAndBinding(0, 0)});
create<ast::BindingDecoration>(0),
create<ast::GroupDecoration>(0),
});
EXPECT_FALSE(r()->Resolve()); EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(), EXPECT_EQ(r()->error(),
"12:34 error: storage textures only support read-only and " "12:34 error: storage textures currently only support 'write' "
"write-only access"); "access control");
} }
TEST_F(StorageTextureAccessTest, ReadOnlyAccess_Pass) { TEST_F(StorageTextureAccessTest, ReadOnlyAccess_Fail) {
// [[group(0), binding(0)]] // [[group(0), binding(0)]]
// var a : texture_storage_1d<ru32int, read>; // var a : texture_storage_1d<ru32int, read>;
auto* st = ty.storage_texture(ast::TextureDimension::k1d, auto* st = ty.storage_texture(Source{{12, 34}}, ast::TextureDimension::k1d,
ast::ImageFormat::kR32Uint, ast::Access::kRead); ast::ImageFormat::kR32Uint, ast::Access::kRead);
Global("a", st, ast::StorageClass::kNone, nullptr, Global("a", st, ast::StorageClass::kNone, nullptr,
ast::DecorationList{ ast::DecorationList{GroupAndBinding(0, 0)});
create<ast::BindingDecoration>(0),
create<ast::GroupDecoration>(0),
});
EXPECT_TRUE(r()->Resolve()) << r()->error(); EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
"12:34 error: storage textures currently only support 'write' "
"access control");
} }
TEST_F(StorageTextureAccessTest, WriteOnlyAccess_Pass) { TEST_F(StorageTextureAccessTest, WriteOnlyAccess_Pass) {
@ -978,10 +952,7 @@ TEST_F(StorageTextureAccessTest, WriteOnlyAccess_Pass) {
ast::Access::kWrite); ast::Access::kWrite);
Global("a", st, ast::StorageClass::kNone, nullptr, Global("a", st, ast::StorageClass::kNone, nullptr,
ast::DecorationList{ ast::DecorationList{GroupAndBinding(0, 0)});
create<ast::BindingDecoration>(0),
create<ast::GroupDecoration>(0),
});
EXPECT_TRUE(r()->Resolve()) << r()->error(); EXPECT_TRUE(r()->Resolve()) << r()->error();
} }

23
src/transform/robustness_test.cc
View File

@ -663,10 +663,6 @@ TEST_F(RobustnessTest, TextureLoad_Clamp) {
[[group(0), binding(0)]] var tex_ms_2d : texture_multisampled_2d<f32>; [[group(0), binding(0)]] var tex_ms_2d : texture_multisampled_2d<f32>;
[[group(0), binding(0)]] var tex_depth_2d : texture_depth_2d; [[group(0), binding(0)]] var tex_depth_2d : texture_depth_2d;
[[group(0), binding(0)]] var tex_depth_2d_arr : texture_depth_2d_array; [[group(0), binding(0)]] var tex_depth_2d_arr : texture_depth_2d_array;
[[group(0), binding(0)]] var tex_storage_1d : texture_storage_1d<rgba8sint, read>;
[[group(0), binding(0)]] var tex_storage_2d : texture_storage_2d<rgba8sint, read>;
[[group(0), binding(0)]] var tex_storage_2d_arr : texture_storage_2d_array<rgba8sint, read>;
[[group(0), binding(0)]] var tex_storage_3d : texture_storage_3d<rgba8sint, read>;
[[group(0), binding(0)]] var tex_external : texture_external; [[group(0), binding(0)]] var tex_external : texture_external;
fn f() { fn f() {
@ -681,15 +677,12 @@ fn f() {
ignore(textureLoad(tex_ms_2d, vec2<i32>(1, 2), sample_idx)); ignore(textureLoad(tex_ms_2d, vec2<i32>(1, 2), sample_idx));
ignore(textureLoad(tex_depth_2d, vec2<i32>(1, 2), level_idx)); ignore(textureLoad(tex_depth_2d, vec2<i32>(1, 2), level_idx));
ignore(textureLoad(tex_depth_2d_arr, vec2<i32>(1, 2), array_idx, level_idx)); ignore(textureLoad(tex_depth_2d_arr, vec2<i32>(1, 2), array_idx, level_idx));
ignore(textureLoad(tex_storage_1d, 1));
ignore(textureLoad(tex_storage_2d, vec2<i32>(1, 2)));
ignore(textureLoad(tex_storage_2d_arr, vec2<i32>(1, 2), array_idx));
ignore(textureLoad(tex_storage_3d, vec3<i32>(1, 2, 3)));
ignore(textureLoad(tex_external, vec2<i32>(1, 2))); ignore(textureLoad(tex_external, vec2<i32>(1, 2)));
} }
)"; )";
auto* expect = R"( auto* expect =
R"(
[[group(0), binding(0)]] var tex_1d : texture_1d<f32>; [[group(0), binding(0)]] var tex_1d : texture_1d<f32>;
[[group(0), binding(0)]] var tex_2d : texture_2d<f32>; [[group(0), binding(0)]] var tex_2d : texture_2d<f32>;
@ -704,14 +697,6 @@ fn f() {
[[group(0), binding(0)]] var tex_depth_2d_arr : texture_depth_2d_array; [[group(0), binding(0)]] var tex_depth_2d_arr : texture_depth_2d_array;
[[group(0), binding(0)]] var tex_storage_1d : texture_storage_1d<rgba8sint, read>;
[[group(0), binding(0)]] var tex_storage_2d : texture_storage_2d<rgba8sint, read>;
[[group(0), binding(0)]] var tex_storage_2d_arr : texture_storage_2d_array<rgba8sint, read>;
[[group(0), binding(0)]] var tex_storage_3d : texture_storage_3d<rgba8sint, read>;
[[group(0), binding(0)]] var tex_external : texture_external; [[group(0), binding(0)]] var tex_external : texture_external;
fn f() { fn f() {
@ -725,10 +710,6 @@ fn f() {
ignore(textureLoad(tex_ms_2d, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_ms_2d) - vec2<i32>(1))), sample_idx)); ignore(textureLoad(tex_ms_2d, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_ms_2d) - vec2<i32>(1))), sample_idx));
ignore(textureLoad(tex_depth_2d, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_depth_2d, clamp(level_idx, 0, (textureNumLevels(tex_depth_2d) - 1))) - vec2<i32>(1))), clamp(level_idx, 0, (textureNumLevels(tex_depth_2d) - 1)))); ignore(textureLoad(tex_depth_2d, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_depth_2d, clamp(level_idx, 0, (textureNumLevels(tex_depth_2d) - 1))) - vec2<i32>(1))), clamp(level_idx, 0, (textureNumLevels(tex_depth_2d) - 1))));
ignore(textureLoad(tex_depth_2d_arr, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_depth_2d_arr, clamp(level_idx, 0, (textureNumLevels(tex_depth_2d_arr) - 1))) - vec2<i32>(1))), clamp(array_idx, 0, (textureNumLayers(tex_depth_2d_arr) - 1)), clamp(level_idx, 0, (textureNumLevels(tex_depth_2d_arr) - 1)))); ignore(textureLoad(tex_depth_2d_arr, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_depth_2d_arr, clamp(level_idx, 0, (textureNumLevels(tex_depth_2d_arr) - 1))) - vec2<i32>(1))), clamp(array_idx, 0, (textureNumLayers(tex_depth_2d_arr) - 1)), clamp(level_idx, 0, (textureNumLevels(tex_depth_2d_arr) - 1))));
ignore(textureLoad(tex_storage_1d, clamp(1, i32(), (textureDimensions(tex_storage_1d) - i32(1)))));
ignore(textureLoad(tex_storage_2d, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_storage_2d) - vec2<i32>(1)))));
ignore(textureLoad(tex_storage_2d_arr, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_storage_2d_arr) - vec2<i32>(1))), clamp(array_idx, 0, (textureNumLayers(tex_storage_2d_arr) - 1))));
ignore(textureLoad(tex_storage_3d, clamp(vec3<i32>(1, 2, 3), vec3<i32>(), (textureDimensions(tex_storage_3d) - vec3<i32>(1)))));
ignore(textureLoad(tex_external, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_external) - vec2<i32>(1))))); ignore(textureLoad(tex_external, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_external) - vec2<i32>(1)))));
} }
)"; )";

26
src/writer/glsl/generator_impl_intrinsic_texture_test.cc
View File

@ -37,20 +37,16 @@ ExpectedResult expected_texture_overload(
using ValidTextureOverload = ast::intrinsic::test::ValidTextureOverload; using ValidTextureOverload = ast::intrinsic::test::ValidTextureOverload;
switch (overload) { switch (overload) {
case ValidTextureOverload::kDimensions1d: case ValidTextureOverload::kDimensions1d:
case ValidTextureOverload::kDimensionsStorageRO1d:
case ValidTextureOverload::kDimensionsStorageWO1d: case ValidTextureOverload::kDimensionsStorageWO1d:
case ValidTextureOverload::kDimensions2d: case ValidTextureOverload::kDimensions2d:
case ValidTextureOverload::kDimensionsDepth2d: case ValidTextureOverload::kDimensionsDepth2d:
case ValidTextureOverload::kDimensionsStorageRO2d:
case ValidTextureOverload::kDimensionsStorageWO2d: case ValidTextureOverload::kDimensionsStorageWO2d:
case ValidTextureOverload::kDimensionsDepthMultisampled2d: case ValidTextureOverload::kDimensionsDepthMultisampled2d:
case ValidTextureOverload::kDimensionsMultisampled2d: case ValidTextureOverload::kDimensionsMultisampled2d:
case ValidTextureOverload::kDimensions2dArray: case ValidTextureOverload::kDimensions2dArray:
case ValidTextureOverload::kDimensionsDepth2dArray: case ValidTextureOverload::kDimensionsDepth2dArray:
case ValidTextureOverload::kDimensionsStorageRO2dArray:
case ValidTextureOverload::kDimensionsStorageWO2dArray: case ValidTextureOverload::kDimensionsStorageWO2dArray:
case ValidTextureOverload::kDimensions3d: case ValidTextureOverload::kDimensions3d:
case ValidTextureOverload::kDimensionsStorageRO3d:
case ValidTextureOverload::kDimensionsStorageWO3d: case ValidTextureOverload::kDimensionsStorageWO3d:
case ValidTextureOverload::kDimensionsCube: case ValidTextureOverload::kDimensionsCube:
case ValidTextureOverload::kDimensionsDepthCube: case ValidTextureOverload::kDimensionsDepthCube:
@ -222,28 +218,6 @@ ExpectedResult expected_texture_overload(
return R"(texture.Load(ivec3(1, 2, 3)).x;)"; return R"(texture.Load(ivec3(1, 2, 3)).x;)";
case ValidTextureOverload::kLoadDepth2dArrayLevelF32: case ValidTextureOverload::kLoadDepth2dArrayLevelF32:
return R"(texture.Load(ivec4(1, 2, 3, 4)).x;)"; return R"(texture.Load(ivec4(1, 2, 3, 4)).x;)";
case ValidTextureOverload::kLoadStorageRO1dRgba32float:
return R"(texture.Load(ivec2(1, 0));)";
case ValidTextureOverload::kLoadStorageRO2dRgba8unorm:
case ValidTextureOverload::kLoadStorageRO2dRgba8snorm:
case ValidTextureOverload::kLoadStorageRO2dRgba8uint:
case ValidTextureOverload::kLoadStorageRO2dRgba8sint:
case ValidTextureOverload::kLoadStorageRO2dRgba16uint:
case ValidTextureOverload::kLoadStorageRO2dRgba16sint:
case ValidTextureOverload::kLoadStorageRO2dRgba16float:
case ValidTextureOverload::kLoadStorageRO2dR32uint:
case ValidTextureOverload::kLoadStorageRO2dR32sint:
case ValidTextureOverload::kLoadStorageRO2dR32float:
case ValidTextureOverload::kLoadStorageRO2dRg32uint:
case ValidTextureOverload::kLoadStorageRO2dRg32sint:
case ValidTextureOverload::kLoadStorageRO2dRg32float:
case ValidTextureOverload::kLoadStorageRO2dRgba32uint:
case ValidTextureOverload::kLoadStorageRO2dRgba32sint:
case ValidTextureOverload::kLoadStorageRO2dRgba32float:
return R"(texture.Load(ivec3(1, 2, 0));)";
case ValidTextureOverload::kLoadStorageRO2dArrayRgba32float:
case ValidTextureOverload::kLoadStorageRO3dRgba32float:
return R"(texture.Load(ivec4(1, 2, 3, 0));)";
case ValidTextureOverload::kStoreWO1dRgba32float: case ValidTextureOverload::kStoreWO1dRgba32float:
return R"(texture[1] = vec4(2.0f, 3.0f, 4.0f, 5.0f);)"; return R"(texture[1] = vec4(2.0f, 3.0f, 4.0f, 5.0f);)";
case ValidTextureOverload::kStoreWO2dRgba32float: case ValidTextureOverload::kStoreWO2dRgba32float:

60
src/writer/glsl/generator_impl_type_test.cc
View File

@ -535,21 +535,17 @@ TEST_F(GlslGeneratorImplTest_Type, EmitMultisampledTexture) {
struct GlslStorageTextureData { struct GlslStorageTextureData {
ast::TextureDimension dim; ast::TextureDimension dim;
ast::ImageFormat imgfmt; ast::ImageFormat imgfmt;
bool ro;
std::string result; std::string result;
}; };
inline std::ostream& operator<<(std::ostream& out, inline std::ostream& operator<<(std::ostream& out,
GlslStorageTextureData data) { GlslStorageTextureData data) {
out << data.dim << (data.ro ? "ReadOnly" : "WriteOnly"); return out << data.dim;
return out;
} }
using GlslStorageTexturesTest = TestParamHelper<GlslStorageTextureData>; using GlslStorageTexturesTest = TestParamHelper<GlslStorageTextureData>;
TEST_P(GlslStorageTexturesTest, Emit) { TEST_P(GlslStorageTexturesTest, Emit) {
auto params = GetParam(); auto params = GetParam();
auto* t = auto* t = ty.storage_texture(params.dim, params.imgfmt, ast::Access::kWrite);
ty.storage_texture(params.dim, params.imgfmt,
params.ro ? ast::Access::kRead : ast::Access::kWrite);
Global("tex", t, Global("tex", t,
ast::DecorationList{ ast::DecorationList{
@ -569,44 +565,44 @@ INSTANTIATE_TEST_SUITE_P(
GlslGeneratorImplTest_Type, GlslGeneratorImplTest_Type,
GlslStorageTexturesTest, GlslStorageTexturesTest,
testing::Values( testing::Values(
GlslStorageTextureData{ast::TextureDimension::k1d,
ast::ImageFormat::kRgba8Unorm, true,
"Texture1D<float4> tex : register(t1, space2);"},
GlslStorageTextureData{ast::TextureDimension::k2d,
ast::ImageFormat::kRgba16Float, true,
"Texture2D<float4> tex : register(t1, space2);"},
GlslStorageTextureData{ GlslStorageTextureData{
ast::TextureDimension::k2dArray, ast::ImageFormat::kR32Float, true, ast::TextureDimension::k1d, ast::ImageFormat::kRgba8Unorm,
"Texture2DArray<float4> tex : register(t1, space2);"},
GlslStorageTextureData{ast::TextureDimension::k3d,
ast::ImageFormat::kRg32Float, true,
"Texture3D<float4> tex : register(t1, space2);"},
GlslStorageTextureData{
ast::TextureDimension::k1d, ast::ImageFormat::kRgba32Float, false,
"RWTexture1D<float4> tex : register(u1, space2);"}, "RWTexture1D<float4> tex : register(u1, space2);"},
GlslStorageTextureData{ GlslStorageTextureData{
ast::TextureDimension::k2d, ast::ImageFormat::kRgba16Uint, false, ast::TextureDimension::k2d, ast::ImageFormat::kRgba16Float,
"RWTexture2D<float4> tex : register(u1, space2);"},
GlslStorageTextureData{
ast::TextureDimension::k2dArray, ast::ImageFormat::kR32Float,
"RWTexture2DArray<float4> tex : register(u1, space2);"},
GlslStorageTextureData{
ast::TextureDimension::k3d, ast::ImageFormat::kRg32Float,
"RWTexture3D<float4> tex : register(u1, space2);"},
GlslStorageTextureData{
ast::TextureDimension::k1d, ast::ImageFormat::kRgba32Float,
"RWTexture1D<float4> tex : register(u1, space2);"},
GlslStorageTextureData{
ast::TextureDimension::k2d, ast::ImageFormat::kRgba16Uint,
"RWTexture2D<uint4> tex : register(u1, space2);"}, "RWTexture2D<uint4> tex : register(u1, space2);"},
GlslStorageTextureData{ GlslStorageTextureData{
ast::TextureDimension::k2dArray, ast::ImageFormat::kR32Uint, false, ast::TextureDimension::k2dArray, ast::ImageFormat::kR32Uint,
"RWTexture2DArray<uint4> tex : register(u1, space2);"}, "RWTexture2DArray<uint4> tex : register(u1, space2);"},
GlslStorageTextureData{ GlslStorageTextureData{
ast::TextureDimension::k3d, ast::ImageFormat::kRg32Uint, false, ast::TextureDimension::k3d, ast::ImageFormat::kRg32Uint,
"RWTexture3D<uint4> tex : register(u1, space2);"}, "RWTexture3D<uint4> tex : register(u1, space2);"},
GlslStorageTextureData{ast::TextureDimension::k1d, GlslStorageTextureData{
ast::ImageFormat::kRgba32Uint, true, ast::TextureDimension::k1d, ast::ImageFormat::kRgba32Uint,
"Texture1D<uint4> tex : register(t1, space2);"}, "RWTexture1D<uint4> tex : register(u1, space2);"},
GlslStorageTextureData{ast::TextureDimension::k2d, GlslStorageTextureData{ast::TextureDimension::k2d,
ast::ImageFormat::kRgba16Sint, true, ast::ImageFormat::kRgba16Sint,
"Texture2D<int4> tex : register(t1, space2);"}, "RWTexture2D<int4> tex : register(u1, space2);"},
GlslStorageTextureData{ GlslStorageTextureData{
ast::TextureDimension::k2dArray, ast::ImageFormat::kR32Sint, true, ast::TextureDimension::k2dArray, ast::ImageFormat::kR32Sint,
"Texture2DArray<int4> tex : register(t1, space2);"}, "RWTexture2DArray<int4> tex : register(u1, space2);"},
GlslStorageTextureData{ast::TextureDimension::k3d, GlslStorageTextureData{ast::TextureDimension::k3d,
ast::ImageFormat::kRg32Sint, true, ast::ImageFormat::kRg32Sint,
"Texture3D<int4> tex : register(t1, space2);"}, "RWTexture3D<int4> tex : register(u1, space2);"},
GlslStorageTextureData{ GlslStorageTextureData{
ast::TextureDimension::k1d, ast::ImageFormat::kRgba32Sint, false, ast::TextureDimension::k1d, ast::ImageFormat::kRgba32Sint,
"RWTexture1D<int4> tex : register(u1, space2);"})); "RWTexture1D<int4> tex : register(u1, space2);"}));
} // namespace } // namespace

4
src/writer/hlsl/generator_impl.cc
View File

@ -2436,11 +2436,9 @@ bool GeneratorImpl::EmitHandleVariable(const sem::Variable* var) {
if (unwrapped_type->Is<sem::Texture>()) { if (unwrapped_type->Is<sem::Texture>()) {
register_space = "t"; register_space = "t";
if (auto* storage_tex = unwrapped_type->As<sem::StorageTexture>()) { if (unwrapped_type->Is<sem::StorageTexture>()) {
if (storage_tex->access() != ast::Access::kRead) {
register_space = "u"; register_space = "u";
} }
}
} else if (unwrapped_type->Is<sem::Sampler>()) { } else if (unwrapped_type->Is<sem::Sampler>()) {
register_space = "s"; register_space = "s";
} }

26
src/writer/hlsl/generator_impl_intrinsic_texture_test.cc
View File

@ -38,7 +38,6 @@ ExpectedResult expected_texture_overload(
using ValidTextureOverload = ast::intrinsic::test::ValidTextureOverload; using ValidTextureOverload = ast::intrinsic::test::ValidTextureOverload;
switch (overload) { switch (overload) {
case ValidTextureOverload::kDimensions1d: case ValidTextureOverload::kDimensions1d:
case ValidTextureOverload::kDimensionsStorageRO1d:
case ValidTextureOverload::kDimensionsStorageWO1d: case ValidTextureOverload::kDimensionsStorageWO1d:
return { return {
R"(int tint_tmp; R"(int tint_tmp;
@ -48,7 +47,6 @@ ExpectedResult expected_texture_overload(
}; };
case ValidTextureOverload::kDimensions2d: case ValidTextureOverload::kDimensions2d:
case ValidTextureOverload::kDimensionsDepth2d: case ValidTextureOverload::kDimensionsDepth2d:
case ValidTextureOverload::kDimensionsStorageRO2d:
case ValidTextureOverload::kDimensionsStorageWO2d: case ValidTextureOverload::kDimensionsStorageWO2d:
return { return {
R"(int2 tint_tmp; R"(int2 tint_tmp;
@ -67,7 +65,6 @@ ExpectedResult expected_texture_overload(
case ValidTextureOverload::kDimensions2dArray: case ValidTextureOverload::kDimensions2dArray:
case ValidTextureOverload::kDimensionsDepth2dArray: case ValidTextureOverload::kDimensionsDepth2dArray:
case ValidTextureOverload::kDimensionsStorageRO2dArray:
case ValidTextureOverload::kDimensionsStorageWO2dArray: case ValidTextureOverload::kDimensionsStorageWO2dArray:
return { return {
R"(int3 tint_tmp; R"(int3 tint_tmp;
@ -76,7 +73,6 @@ ExpectedResult expected_texture_overload(
"tint_tmp.xy;", "tint_tmp.xy;",
}; };
case ValidTextureOverload::kDimensions3d: case ValidTextureOverload::kDimensions3d:
case ValidTextureOverload::kDimensionsStorageRO3d:
case ValidTextureOverload::kDimensionsStorageWO3d: case ValidTextureOverload::kDimensionsStorageWO3d:
return { return {
R"(int3 tint_tmp; R"(int3 tint_tmp;
@ -317,28 +313,6 @@ ExpectedResult expected_texture_overload(
return R"(tint_symbol.Load(int3(1, 2, 3)).x;)"; return R"(tint_symbol.Load(int3(1, 2, 3)).x;)";
case ValidTextureOverload::kLoadDepth2dArrayLevelF32: case ValidTextureOverload::kLoadDepth2dArrayLevelF32:
return R"(tint_symbol.Load(int4(1, 2, 3, 4)).x;)"; return R"(tint_symbol.Load(int4(1, 2, 3, 4)).x;)";
case ValidTextureOverload::kLoadStorageRO1dRgba32float:
return R"(tint_symbol.Load(int2(1, 0));)";
case ValidTextureOverload::kLoadStorageRO2dRgba8unorm:
case ValidTextureOverload::kLoadStorageRO2dRgba8snorm:
case ValidTextureOverload::kLoadStorageRO2dRgba8uint:
case ValidTextureOverload::kLoadStorageRO2dRgba8sint:
case ValidTextureOverload::kLoadStorageRO2dRgba16uint:
case ValidTextureOverload::kLoadStorageRO2dRgba16sint:
case ValidTextureOverload::kLoadStorageRO2dRgba16float:
case ValidTextureOverload::kLoadStorageRO2dR32uint:
case ValidTextureOverload::kLoadStorageRO2dR32sint:
case ValidTextureOverload::kLoadStorageRO2dR32float:
case ValidTextureOverload::kLoadStorageRO2dRg32uint:
case ValidTextureOverload::kLoadStorageRO2dRg32sint:
case ValidTextureOverload::kLoadStorageRO2dRg32float:
case ValidTextureOverload::kLoadStorageRO2dRgba32uint:
case ValidTextureOverload::kLoadStorageRO2dRgba32sint:
case ValidTextureOverload::kLoadStorageRO2dRgba32float:
return R"(tint_symbol.Load(int3(1, 2, 0));)";
case ValidTextureOverload::kLoadStorageRO2dArrayRgba32float:
case ValidTextureOverload::kLoadStorageRO3dRgba32float:
return R"(tint_symbol.Load(int4(1, 2, 3, 0));)";
case ValidTextureOverload::kStoreWO1dRgba32float: case ValidTextureOverload::kStoreWO1dRgba32float:
return R"(tint_symbol[1] = float4(2.0f, 3.0f, 4.0f, 5.0f);)"; return R"(tint_symbol[1] = float4(2.0f, 3.0f, 4.0f, 5.0f);)";
case ValidTextureOverload::kStoreWO2dRgba32float: case ValidTextureOverload::kStoreWO2dRgba32float:

65
src/writer/hlsl/generator_impl_type_test.cc
View File

@ -533,27 +533,20 @@ TEST_F(HlslGeneratorImplTest_Type, EmitMultisampledTexture) {
struct HlslStorageTextureData { struct HlslStorageTextureData {
ast::TextureDimension dim; ast::TextureDimension dim;
ast::ImageFormat imgfmt; ast::ImageFormat imgfmt;
bool ro;
std::string result; std::string result;
}; };
inline std::ostream& operator<<(std::ostream& out, inline std::ostream& operator<<(std::ostream& out,
HlslStorageTextureData data) { HlslStorageTextureData data) {
out << data.dim << (data.ro ? "ReadOnly" : "WriteOnly"); out << data.dim;
return out; return out;
} }
using HlslStorageTexturesTest = TestParamHelper<HlslStorageTextureData>; using HlslStorageTexturesTest = TestParamHelper<HlslStorageTextureData>;
TEST_P(HlslStorageTexturesTest, Emit) { TEST_P(HlslStorageTexturesTest, Emit) {
auto params = GetParam(); auto params = GetParam();
auto* t = auto* t = ty.storage_texture(params.dim, params.imgfmt, ast::Access::kWrite);
ty.storage_texture(params.dim, params.imgfmt,
params.ro ? ast::Access::kRead : ast::Access::kWrite);
Global("tex", t, Global("tex", t, ast::DecorationList{GroupAndBinding(2, 1)});
ast::DecorationList{
create<ast::BindingDecoration>(1),
create<ast::GroupDecoration>(2),
});
Func("main", {}, ty.void_(), {Ignore(Call("textureDimensions", "tex"))}, Func("main", {}, ty.void_(), {Ignore(Call("textureDimensions", "tex"))},
{Stage(ast::PipelineStage::kFragment)}); {Stage(ast::PipelineStage::kFragment)});
@ -567,44 +560,44 @@ INSTANTIATE_TEST_SUITE_P(
HlslGeneratorImplTest_Type, HlslGeneratorImplTest_Type,
HlslStorageTexturesTest, HlslStorageTexturesTest,
testing::Values( testing::Values(
HlslStorageTextureData{ast::TextureDimension::k1d,
ast::ImageFormat::kRgba8Unorm, true,
"Texture1D<float4> tex : register(t1, space2);"},
HlslStorageTextureData{ast::TextureDimension::k2d,
ast::ImageFormat::kRgba16Float, true,
"Texture2D<float4> tex : register(t1, space2);"},
HlslStorageTextureData{ HlslStorageTextureData{
ast::TextureDimension::k2dArray, ast::ImageFormat::kR32Float, true, ast::TextureDimension::k1d, ast::ImageFormat::kRgba8Unorm,
"Texture2DArray<float4> tex : register(t1, space2);"},
HlslStorageTextureData{ast::TextureDimension::k3d,
ast::ImageFormat::kRg32Float, true,
"Texture3D<float4> tex : register(t1, space2);"},
HlslStorageTextureData{
ast::TextureDimension::k1d, ast::ImageFormat::kRgba32Float, false,
"RWTexture1D<float4> tex : register(u1, space2);"}, "RWTexture1D<float4> tex : register(u1, space2);"},
HlslStorageTextureData{ HlslStorageTextureData{
ast::TextureDimension::k2d, ast::ImageFormat::kRgba16Uint, false, ast::TextureDimension::k2d, ast::ImageFormat::kRgba16Float,
"RWTexture2D<float4> tex : register(u1, space2);"},
HlslStorageTextureData{
ast::TextureDimension::k2dArray, ast::ImageFormat::kR32Float,
"RWTexture2DArray<float4> tex : register(u1, space2);"},
HlslStorageTextureData{
ast::TextureDimension::k3d, ast::ImageFormat::kRg32Float,
"RWTexture3D<float4> tex : register(u1, space2);"},
HlslStorageTextureData{
ast::TextureDimension::k1d, ast::ImageFormat::kRgba32Float,
"RWTexture1D<float4> tex : register(u1, space2);"},
HlslStorageTextureData{
ast::TextureDimension::k2d, ast::ImageFormat::kRgba16Uint,
"RWTexture2D<uint4> tex : register(u1, space2);"}, "RWTexture2D<uint4> tex : register(u1, space2);"},
HlslStorageTextureData{ HlslStorageTextureData{
ast::TextureDimension::k2dArray, ast::ImageFormat::kR32Uint, false, ast::TextureDimension::k2dArray, ast::ImageFormat::kR32Uint,
"RWTexture2DArray<uint4> tex : register(u1, space2);"}, "RWTexture2DArray<uint4> tex : register(u1, space2);"},
HlslStorageTextureData{ HlslStorageTextureData{
ast::TextureDimension::k3d, ast::ImageFormat::kRg32Uint, false, ast::TextureDimension::k3d, ast::ImageFormat::kRg32Uint,
"RWTexture3D<uint4> tex : register(u1, space2);"}, "RWTexture3D<uint4> tex : register(u1, space2);"},
HlslStorageTextureData{ast::TextureDimension::k1d, HlslStorageTextureData{
ast::ImageFormat::kRgba32Uint, true, ast::TextureDimension::k1d, ast::ImageFormat::kRgba32Uint,
"Texture1D<uint4> tex : register(t1, space2);"}, "RWTexture1D<uint4> tex : register(u1, space2);"},
HlslStorageTextureData{ast::TextureDimension::k2d, HlslStorageTextureData{ast::TextureDimension::k2d,
ast::ImageFormat::kRgba16Sint, true, ast::ImageFormat::kRgba16Sint,
"Texture2D<int4> tex : register(t1, space2);"}, "RWTexture2D<int4> tex : register(u1, space2);"},
HlslStorageTextureData{ HlslStorageTextureData{
ast::TextureDimension::k2dArray, ast::ImageFormat::kR32Sint, true, ast::TextureDimension::k2dArray, ast::ImageFormat::kR32Sint,
"Texture2DArray<int4> tex : register(t1, space2);"}, "RWTexture2DArray<int4> tex : register(u1, space2);"},
HlslStorageTextureData{ast::TextureDimension::k3d, HlslStorageTextureData{ast::TextureDimension::k3d,
ast::ImageFormat::kRg32Sint, true, ast::ImageFormat::kRg32Sint,
"Texture3D<int4> tex : register(t1, space2);"}, "RWTexture3D<int4> tex : register(u1, space2);"},
HlslStorageTextureData{ HlslStorageTextureData{
ast::TextureDimension::k1d, ast::ImageFormat::kRgba32Sint, false, ast::TextureDimension::k1d, ast::ImageFormat::kRgba32Sint,
"RWTexture1D<int4> tex : register(u1, space2);"})); "RWTexture1D<int4> tex : register(u1, space2);"}));
} // namespace } // namespace

27
src/writer/msl/generator_impl_intrinsic_texture_test.cc
View File

@ -26,7 +26,6 @@ std::string expected_texture_overload(
using ValidTextureOverload = ast::intrinsic::test::ValidTextureOverload; using ValidTextureOverload = ast::intrinsic::test::ValidTextureOverload;
switch (overload) { switch (overload) {
case ValidTextureOverload::kDimensions1d: case ValidTextureOverload::kDimensions1d:
case ValidTextureOverload::kDimensionsStorageRO1d:
case ValidTextureOverload::kDimensionsStorageWO1d: case ValidTextureOverload::kDimensionsStorageWO1d:
return R"(int(texture.get_width()))"; return R"(int(texture.get_width()))";
case ValidTextureOverload::kDimensions2d: case ValidTextureOverload::kDimensions2d:
@ -39,13 +38,10 @@ std::string expected_texture_overload(
case ValidTextureOverload::kDimensionsDepthCube: case ValidTextureOverload::kDimensionsDepthCube:
case ValidTextureOverload::kDimensionsDepthCubeArray: case ValidTextureOverload::kDimensionsDepthCubeArray:
case ValidTextureOverload::kDimensionsDepthMultisampled2d: case ValidTextureOverload::kDimensionsDepthMultisampled2d:
case ValidTextureOverload::kDimensionsStorageRO2d:
case ValidTextureOverload::kDimensionsStorageRO2dArray:
case ValidTextureOverload::kDimensionsStorageWO2d: case ValidTextureOverload::kDimensionsStorageWO2d:
case ValidTextureOverload::kDimensionsStorageWO2dArray: case ValidTextureOverload::kDimensionsStorageWO2dArray:
return R"(int2(texture.get_width(), texture.get_height()))"; return R"(int2(texture.get_width(), texture.get_height()))";
case ValidTextureOverload::kDimensions3d: case ValidTextureOverload::kDimensions3d:
case ValidTextureOverload::kDimensionsStorageRO3d:
case ValidTextureOverload::kDimensionsStorageWO3d: case ValidTextureOverload::kDimensionsStorageWO3d:
return R"(int3(texture.get_width(), texture.get_height(), texture.get_depth()))"; return R"(int3(texture.get_width(), texture.get_height(), texture.get_depth()))";
case ValidTextureOverload::kDimensions2dLevel: case ValidTextureOverload::kDimensions2dLevel:
@ -227,29 +223,6 @@ std::string expected_texture_overload(
return R"(texture.read(uint2(int2(1, 2)), 3))"; return R"(texture.read(uint2(int2(1, 2)), 3))";
case ValidTextureOverload::kLoadDepth2dArrayLevelF32: case ValidTextureOverload::kLoadDepth2dArrayLevelF32:
return R"(texture.read(uint2(int2(1, 2)), 3, 4))"; return R"(texture.read(uint2(int2(1, 2)), 3, 4))";
case ValidTextureOverload::kLoadStorageRO1dRgba32float:
return R"(texture.read(uint(1)))";
case ValidTextureOverload::kLoadStorageRO2dRgba8unorm:
case ValidTextureOverload::kLoadStorageRO2dRgba8snorm:
case ValidTextureOverload::kLoadStorageRO2dRgba8uint:
case ValidTextureOverload::kLoadStorageRO2dRgba8sint:
case ValidTextureOverload::kLoadStorageRO2dRgba16uint:
case ValidTextureOverload::kLoadStorageRO2dRgba16sint:
case ValidTextureOverload::kLoadStorageRO2dRgba16float:
case ValidTextureOverload::kLoadStorageRO2dR32uint:
case ValidTextureOverload::kLoadStorageRO2dR32sint:
case ValidTextureOverload::kLoadStorageRO2dR32float:
case ValidTextureOverload::kLoadStorageRO2dRg32uint:
case ValidTextureOverload::kLoadStorageRO2dRg32sint:
case ValidTextureOverload::kLoadStorageRO2dRg32float:
case ValidTextureOverload::kLoadStorageRO2dRgba32uint:
case ValidTextureOverload::kLoadStorageRO2dRgba32sint:
case ValidTextureOverload::kLoadStorageRO2dRgba32float:
return R"(texture.read(uint2(int2(1, 2))))";
case ValidTextureOverload::kLoadStorageRO2dArrayRgba32float:
return R"(texture.read(uint2(int2(1, 2)), 3))";
case ValidTextureOverload::kLoadStorageRO3dRgba32float:
return R"(texture.read(uint3(int3(1, 2, 3))))";
case ValidTextureOverload::kStoreWO1dRgba32float: case ValidTextureOverload::kStoreWO1dRgba32float:
return R"(texture.write(float4(2.0f, 3.0f, 4.0f, 5.0f), uint(1)))"; return R"(texture.write(float4(2.0f, 3.0f, 4.0f, 5.0f), uint(1)))";
case ValidTextureOverload::kStoreWO2dRgba32float: case ValidTextureOverload::kStoreWO2dRgba32float:

27
src/writer/msl/generator_impl_type_test.cc
View File

@ -843,20 +843,17 @@ TEST_F(MslGeneratorImplTest, Emit_TypeMultisampledTexture) {
struct MslStorageTextureData { struct MslStorageTextureData {
ast::TextureDimension dim; ast::TextureDimension dim;
bool ro;
std::string result; std::string result;
}; };
inline std::ostream& operator<<(std::ostream& out, MslStorageTextureData data) { inline std::ostream& operator<<(std::ostream& out, MslStorageTextureData data) {
out << data.dim << (data.ro ? "ReadOnly" : "WriteOnly"); return out << data.dim;
return out;
} }
using MslStorageTexturesTest = TestParamHelper<MslStorageTextureData>; using MslStorageTexturesTest = TestParamHelper<MslStorageTextureData>;
TEST_P(MslStorageTexturesTest, Emit) { TEST_P(MslStorageTexturesTest, Emit) {
auto params = GetParam(); auto params = GetParam();
auto* s = auto* s = ty.storage_texture(params.dim, ast::ImageFormat::kR32Float,
ty.storage_texture(params.dim, ast::ImageFormat::kR32Float, ast::Access::kWrite);
params.ro ? ast::Access::kRead : ast::Access::kWrite);
Global("test_var", s, Global("test_var", s,
ast::DecorationList{ ast::DecorationList{
create<ast::BindingDecoration>(0), create<ast::BindingDecoration>(0),
@ -872,22 +869,14 @@ TEST_P(MslStorageTexturesTest, Emit) {
INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P(
MslGeneratorImplTest, MslGeneratorImplTest,
MslStorageTexturesTest, MslStorageTexturesTest,
testing::Values( testing::Values(MslStorageTextureData{ast::TextureDimension::k1d,
MslStorageTextureData{ast::TextureDimension::k1d, true,
"texture1d<float, access::read>"},
MslStorageTextureData{ast::TextureDimension::k2d, true,
"texture2d<float, access::read>"},
MslStorageTextureData{ast::TextureDimension::k2dArray, true,
"texture2d_array<float, access::read>"},
MslStorageTextureData{ast::TextureDimension::k3d, true,
"texture3d<float, access::read>"},
MslStorageTextureData{ast::TextureDimension::k1d, false,
"texture1d<float, access::write>"}, "texture1d<float, access::write>"},
MslStorageTextureData{ast::TextureDimension::k2d, false, MslStorageTextureData{ast::TextureDimension::k2d,
"texture2d<float, access::write>"}, "texture2d<float, access::write>"},
MslStorageTextureData{ast::TextureDimension::k2dArray, false, MslStorageTextureData{
ast::TextureDimension::k2dArray,
"texture2d_array<float, access::write>"}, "texture2d_array<float, access::write>"},
MslStorageTextureData{ast::TextureDimension::k3d, false, MslStorageTextureData{ast::TextureDimension::k3d,
"texture3d<float, access::write>"})); "texture3d<float, access::write>"}));
} // namespace } // namespace

36
src/writer/spirv/builder_global_variable_test.cc
View File

@ -506,34 +506,6 @@ OpName %5 "c"
)"); )");
} }
TEST_F(BuilderTest, GlobalVar_TextureStorageReadOnly) {
// var<uniform_constant> a : texture_storage_2d<r32uint, read>;
auto* type =
ty.storage_texture(ast::TextureDimension::k2d, ast::ImageFormat::kR32Uint,
ast::Access::kRead);
auto* var_a = Global("a", type,
ast::DecorationList{
create<ast::BindingDecoration>(0),
create<ast::GroupDecoration>(0),
});
spirv::Builder& b = Build();
EXPECT_TRUE(b.GenerateGlobalVariable(var_a)) << b.error();
EXPECT_EQ(DumpInstructions(b.annots()), R"(OpDecorate %1 NonWritable
OpDecorate %1 Binding 0
OpDecorate %1 DescriptorSet 0
)");
EXPECT_EQ(DumpInstructions(b.types()), R"(%4 = OpTypeInt 32 0
%3 = OpTypeImage %4 2D 0 0 0 2 R32ui
%2 = OpTypePointer UniformConstant %3
%1 = OpVariable %2 UniformConstant
)");
}
TEST_F(BuilderTest, GlobalVar_TextureStorageWriteOnly) { TEST_F(BuilderTest, GlobalVar_TextureStorageWriteOnly) {
// var<uniform_constant> a : texture_storage_2d<r32uint, write>; // var<uniform_constant> a : texture_storage_2d<r32uint, write>;
@ -564,13 +536,15 @@ OpDecorate %1 DescriptorSet 0
// Check that multiple texture_storage types with different access modifiers // Check that multiple texture_storage types with different access modifiers
// only produces a single OpTypeImage. // only produces a single OpTypeImage.
TEST_F(BuilderTest, GlobalVar_TextureStorageWithDifferentAccess) { // Test disabled as storage textures currently only support 'write' access. In
// var<uniform_constant> a : texture_storage_2d<r32uint, read>; // the future we'll likely support read_write.
TEST_F(BuilderTest, DISABLED_GlobalVar_TextureStorageWithDifferentAccess) {
// var<uniform_constant> a : texture_storage_2d<r32uint, read_write>;
// var<uniform_constant> b : texture_storage_2d<r32uint, write>; // var<uniform_constant> b : texture_storage_2d<r32uint, write>;
auto* type_a = auto* type_a =
ty.storage_texture(ast::TextureDimension::k2d, ast::ImageFormat::kR32Uint, ty.storage_texture(ast::TextureDimension::k2d, ast::ImageFormat::kR32Uint,
ast::Access::kRead); ast::Access::kReadWrite);
auto* var_a = Global("a", type_a, ast::StorageClass::kNone, auto* var_a = Global("a", type_a, ast::StorageClass::kNone,
ast::DecorationList{ ast::DecorationList{
create<ast::BindingDecoration>(0), create<ast::BindingDecoration>(0),

518
src/writer/spirv/builder_intrinsic_texture_test.cc
View File

@ -492,88 +492,6 @@ OpCapability ImageQuery
)", )",
R"( R"(
OpCapability ImageQuery OpCapability ImageQuery
)"};
case ValidTextureOverload::kDimensionsStorageRO1d:
return {
R"(
%4 = OpTypeFloat 32
%3 = OpTypeImage %4 1D 0 0 0 2 Rgba32f
%2 = OpTypePointer UniformConstant %3
%1 = OpVariable %2 UniformConstant
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeInt 32 1
)",
R"(
%10 = OpLoad %3 %1
%8 = OpImageQuerySize %9 %10
)",
R"(
OpCapability Image1D
OpCapability ImageQuery
)"};
case ValidTextureOverload::kDimensionsStorageRO2d:
return {
R"(
%4 = OpTypeFloat 32
%3 = OpTypeImage %4 2D 0 0 0 2 Rgba32f
%2 = OpTypePointer UniformConstant %3
%1 = OpVariable %2 UniformConstant
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%10 = OpTypeInt 32 1
%9 = OpTypeVector %10 2
)",
R"(
%11 = OpLoad %3 %1
%8 = OpImageQuerySize %9 %11
)",
R"(
OpCapability ImageQuery
)"};
case ValidTextureOverload::kDimensionsStorageRO2dArray:
return {
R"(
%4 = OpTypeFloat 32
%3 = OpTypeImage %4 2D 0 1 0 2 Rgba32f
%2 = OpTypePointer UniformConstant %3
%1 = OpVariable %2 UniformConstant
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%10 = OpTypeInt 32 1
%9 = OpTypeVector %10 2
%12 = OpTypeVector %10 3
)",
R"(
%13 = OpLoad %3 %1
%11 = OpImageQuerySize %12 %13
%8 = OpVectorShuffle %9 %11 %11 0 1
)",
R"(
OpCapability ImageQuery
)"};
case ValidTextureOverload::kDimensionsStorageRO3d:
return {
R"(
%4 = OpTypeFloat 32
%3 = OpTypeImage %4 3D 0 0 0 2 Rgba32f
%2 = OpTypePointer UniformConstant %3
%1 = OpVariable %2 UniformConstant
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%10 = OpTypeInt 32 1
%9 = OpTypeVector %10 3
)",
R"(
%11 = OpLoad %3 %1
%8 = OpImageQuerySize %9 %11
)",
R"(
OpCapability ImageQuery
)"}; )"};
case ValidTextureOverload::kDimensionsStorageWO1d: case ValidTextureOverload::kDimensionsStorageWO1d:
return { return {
@ -3115,442 +3033,6 @@ OpCapability Sampled1D
%11 = OpLoad %3 %1 %11 = OpLoad %3 %1
%9 = OpImageFetch %10 %11 %17 Sample %18 %9 = OpImageFetch %10 %11 %17 Sample %18
%8 = OpCompositeExtract %4 %9 0 %8 = OpCompositeExtract %4 %9 0
)",
R"(
)"};
case ValidTextureOverload::kLoadStorageRO1dRgba32float:
return {
R"(
%4 = OpTypeFloat 32
%3 = OpTypeImage %4 1D 0 0 0 2 Rgba32f
%2 = OpTypePointer UniformConstant %3
%1 = OpVariable %2 UniformConstant
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeVector %4 4
%11 = OpTypeInt 32 1
%12 = OpConstant %11 1
)",
R"(
%10 = OpLoad %3 %1
%8 = OpImageRead %9 %10 %12
)",
R"(
OpCapability Image1D
)"};
case ValidTextureOverload::kLoadStorageRO2dRgba8unorm:
return {
R"(
%4 = OpTypeFloat 32
%3 = OpTypeImage %4 2D 0 0 0 2 Rgba8
%2 = OpTypePointer UniformConstant %3
%1 = OpVariable %2 UniformConstant
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeVector %4 4
%12 = OpTypeInt 32 1
%11 = OpTypeVector %12 2
%13 = OpConstant %12 1
%14 = OpConstant %12 2
%15 = OpConstantComposite %11 %13 %14
)",
R"(
%10 = OpLoad %3 %1
%8 = OpImageRead %9 %10 %15
)",
R"(
)"};
case ValidTextureOverload::kLoadStorageRO2dRgba8snorm:
return {
R"(
%4 = OpTypeFloat 32
%3 = OpTypeImage %4 2D 0 0 0 2 Rgba8Snorm
%2 = OpTypePointer UniformConstant %3
%1 = OpVariable %2 UniformConstant
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeVector %4 4
%12 = OpTypeInt 32 1
%11 = OpTypeVector %12 2
%13 = OpConstant %12 1
%14 = OpConstant %12 2
%15 = OpConstantComposite %11 %13 %14
)",
R"(
%10 = OpLoad %3 %1
%8 = OpImageRead %9 %10 %15
)",
R"(
)"};
case ValidTextureOverload::kLoadStorageRO2dRgba8uint:
return {
R"(
%4 = OpTypeInt 32 0
%3 = OpTypeImage %4 2D 0 0 0 2 Rgba8ui
%2 = OpTypePointer UniformConstant %3
%1 = OpVariable %2 UniformConstant
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeVector %4 4
%12 = OpTypeInt 32 1
%11 = OpTypeVector %12 2
%13 = OpConstant %12 1
%14 = OpConstant %12 2
%15 = OpConstantComposite %11 %13 %14
)",
R"(
%10 = OpLoad %3 %1
%8 = OpImageRead %9 %10 %15
)",
R"(
)"};
case ValidTextureOverload::kLoadStorageRO2dRgba8sint:
return {
R"(
%4 = OpTypeInt 32 1
%3 = OpTypeImage %4 2D 0 0 0 2 Rgba8i
%2 = OpTypePointer UniformConstant %3
%1 = OpVariable %2 UniformConstant
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeVector %4 4
%11 = OpTypeVector %4 2
%12 = OpConstant %4 1
%13 = OpConstant %4 2
%14 = OpConstantComposite %11 %12 %13
)",
R"(
%10 = OpLoad %3 %1
%8 = OpImageRead %9 %10 %14
)",
R"(
)"};
case ValidTextureOverload::kLoadStorageRO2dRgba16uint:
return {
R"(
%4 = OpTypeInt 32 0
%3 = OpTypeImage %4 2D 0 0 0 2 Rgba16ui
%2 = OpTypePointer UniformConstant %3
%1 = OpVariable %2 UniformConstant
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeVector %4 4
%12 = OpTypeInt 32 1
%11 = OpTypeVector %12 2
%13 = OpConstant %12 1
%14 = OpConstant %12 2
%15 = OpConstantComposite %11 %13 %14
)",
R"(
%10 = OpLoad %3 %1
%8 = OpImageRead %9 %10 %15
)",
R"(
)"};
case ValidTextureOverload::kLoadStorageRO2dRgba16sint:
return {
R"(
%4 = OpTypeInt 32 1
%3 = OpTypeImage %4 2D 0 0 0 2 Rgba16i
%2 = OpTypePointer UniformConstant %3
%1 = OpVariable %2 UniformConstant
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeVector %4 4
%11 = OpTypeVector %4 2
%12 = OpConstant %4 1
%13 = OpConstant %4 2
%14 = OpConstantComposite %11 %12 %13
)",
R"(
%10 = OpLoad %3 %1
%8 = OpImageRead %9 %10 %14
)",
R"(
)"};
case ValidTextureOverload::kLoadStorageRO2dRgba16float:
return {
R"(
%4 = OpTypeFloat 32
%3 = OpTypeImage %4 2D 0 0 0 2 Rgba16f
%2 = OpTypePointer UniformConstant %3
%1 = OpVariable %2 UniformConstant
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeVector %4 4
%12 = OpTypeInt 32 1
%11 = OpTypeVector %12 2
%13 = OpConstant %12 1
%14 = OpConstant %12 2
%15 = OpConstantComposite %11 %13 %14
)",
R"(
%10 = OpLoad %3 %1
%8 = OpImageRead %9 %10 %15
)",
R"(
)"};
case ValidTextureOverload::kLoadStorageRO2dR32uint:
return {
R"(
%4 = OpTypeInt 32 0
%3 = OpTypeImage %4 2D 0 0 0 2 R32ui
%2 = OpTypePointer UniformConstant %3
%1 = OpVariable %2 UniformConstant
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeVector %4 4
%12 = OpTypeInt 32 1
%11 = OpTypeVector %12 2
%13 = OpConstant %12 1
%14 = OpConstant %12 2
%15 = OpConstantComposite %11 %13 %14
)",
R"(
%10 = OpLoad %3 %1
%8 = OpImageRead %9 %10 %15
)",
R"(
)"};
case ValidTextureOverload::kLoadStorageRO2dR32sint:
return {
R"(
%4 = OpTypeInt 32 1
%3 = OpTypeImage %4 2D 0 0 0 2 R32i
%2 = OpTypePointer UniformConstant %3
%1 = OpVariable %2 UniformConstant
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeVector %4 4
%11 = OpTypeVector %4 2
%12 = OpConstant %4 1
%13 = OpConstant %4 2
%14 = OpConstantComposite %11 %12 %13
)",
R"(
%10 = OpLoad %3 %1
%8 = OpImageRead %9 %10 %14
)",
R"(
)"};
case ValidTextureOverload::kLoadStorageRO2dR32float:
return {
R"(
%4 = OpTypeFloat 32
%3 = OpTypeImage %4 2D 0 0 0 2 R32f
%2 = OpTypePointer UniformConstant %3
%1 = OpVariable %2 UniformConstant
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeVector %4 4
%12 = OpTypeInt 32 1
%11 = OpTypeVector %12 2
%13 = OpConstant %12 1
%14 = OpConstant %12 2
%15 = OpConstantComposite %11 %13 %14
)",
R"(
%10 = OpLoad %3 %1
%8 = OpImageRead %9 %10 %15
)",
R"(
)"};
case ValidTextureOverload::kLoadStorageRO2dRg32uint:
return {
R"(
%4 = OpTypeInt 32 0
%3 = OpTypeImage %4 2D 0 0 0 2 Rg32ui
%2 = OpTypePointer UniformConstant %3
%1 = OpVariable %2 UniformConstant
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeVector %4 4
%12 = OpTypeInt 32 1
%11 = OpTypeVector %12 2
%13 = OpConstant %12 1
%14 = OpConstant %12 2
%15 = OpConstantComposite %11 %13 %14
)",
R"(
%10 = OpLoad %3 %1
%8 = OpImageRead %9 %10 %15
)",
R"(
OpCapability StorageImageExtendedFormats
)"};
case ValidTextureOverload::kLoadStorageRO2dRg32sint:
return {
R"(
%4 = OpTypeInt 32 1
%3 = OpTypeImage %4 2D 0 0 0 2 Rg32i
%2 = OpTypePointer UniformConstant %3
%1 = OpVariable %2 UniformConstant
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeVector %4 4
%11 = OpTypeVector %4 2
%12 = OpConstant %4 1
%13 = OpConstant %4 2
%14 = OpConstantComposite %11 %12 %13
)",
R"(
%10 = OpLoad %3 %1
%8 = OpImageRead %9 %10 %14
)",
R"(
OpCapability StorageImageExtendedFormats
)"};
case ValidTextureOverload::kLoadStorageRO2dRg32float:
return {
R"(
%4 = OpTypeFloat 32
%3 = OpTypeImage %4 2D 0 0 0 2 Rg32f
%2 = OpTypePointer UniformConstant %3
%1 = OpVariable %2 UniformConstant
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeVector %4 4
%12 = OpTypeInt 32 1
%11 = OpTypeVector %12 2
%13 = OpConstant %12 1
%14 = OpConstant %12 2
%15 = OpConstantComposite %11 %13 %14
)",
R"(
%10 = OpLoad %3 %1
%8 = OpImageRead %9 %10 %15
)",
R"(
OpCapability StorageImageExtendedFormats
)"};
case ValidTextureOverload::kLoadStorageRO2dRgba32uint:
return {
R"(
%4 = OpTypeInt 32 0
%3 = OpTypeImage %4 2D 0 0 0 2 Rgba32ui
%2 = OpTypePointer UniformConstant %3
%1 = OpVariable %2 UniformConstant
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeVector %4 4
%12 = OpTypeInt 32 1
%11 = OpTypeVector %12 2
%13 = OpConstant %12 1
%14 = OpConstant %12 2
%15 = OpConstantComposite %11 %13 %14
)",
R"(
%10 = OpLoad %3 %1
%8 = OpImageRead %9 %10 %15
)",
R"(
)"};
case ValidTextureOverload::kLoadStorageRO2dRgba32sint:
return {
R"(
%4 = OpTypeInt 32 1
%3 = OpTypeImage %4 2D 0 0 0 2 Rgba32i
%2 = OpTypePointer UniformConstant %3
%1 = OpVariable %2 UniformConstant
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeVector %4 4
%11 = OpTypeVector %4 2
%12 = OpConstant %4 1
%13 = OpConstant %4 2
%14 = OpConstantComposite %11 %12 %13
)",
R"(
%10 = OpLoad %3 %1
%8 = OpImageRead %9 %10 %14
)",
R"(
)"};
case ValidTextureOverload::kLoadStorageRO2dRgba32float:
return {
R"(
%4 = OpTypeFloat 32
%3 = OpTypeImage %4 2D 0 0 0 2 Rgba32f
%2 = OpTypePointer UniformConstant %3
%1 = OpVariable %2 UniformConstant
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeVector %4 4
%12 = OpTypeInt 32 1
%11 = OpTypeVector %12 2
%13 = OpConstant %12 1
%14 = OpConstant %12 2
%15 = OpConstantComposite %11 %13 %14
)",
R"(
%10 = OpLoad %3 %1
%8 = OpImageRead %9 %10 %15
)",
R"(
)"};
case ValidTextureOverload::kLoadStorageRO2dArrayRgba32float:
return {
R"(
%4 = OpTypeFloat 32
%3 = OpTypeImage %4 2D 0 1 0 2 Rgba32f
%2 = OpTypePointer UniformConstant %3
%1 = OpVariable %2 UniformConstant
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeVector %4 4
%12 = OpTypeInt 32 1
%11 = OpTypeVector %12 3
%13 = OpConstant %12 1
%14 = OpConstant %12 2
%15 = OpConstant %12 3
%16 = OpConstantComposite %11 %13 %14 %15
)",
R"(
%10 = OpLoad %3 %1
%8 = OpImageRead %9 %10 %16
)",
R"(
)"};
case ValidTextureOverload::kLoadStorageRO3dRgba32float:
return {
R"(
%4 = OpTypeFloat 32
%3 = OpTypeImage %4 3D 0 0 0 2 Rgba32f
%2 = OpTypePointer UniformConstant %3
%1 = OpVariable %2 UniformConstant
%7 = OpTypeSampler
%6 = OpTypePointer UniformConstant %7
%5 = OpVariable %6 UniformConstant
%9 = OpTypeVector %4 4
%12 = OpTypeInt 32 1
%11 = OpTypeVector %12 3
%13 = OpConstant %12 1
%14 = OpConstant %12 2
%15 = OpConstant %12 3
%16 = OpConstantComposite %11 %13 %14 %15
)",
R"(
%10 = OpLoad %3 %1
%8 = OpImageRead %9 %10 %16
)", )",
R"( R"(
)"}; )"};

29
src/writer/spirv/builder_type_test.cc
View File

@ -814,8 +814,9 @@ TEST_F(BuilderTest_Type, SampledTexture_Generate_CubeArray) {
} }
TEST_F(BuilderTest_Type, StorageTexture_Generate_1d) { TEST_F(BuilderTest_Type, StorageTexture_Generate_1d) {
auto* s = ty.storage_texture(ast::TextureDimension::k1d, auto* s =
ast::ImageFormat::kR32Float, ast::Access::kRead); ty.storage_texture(ast::TextureDimension::k1d,
ast::ImageFormat::kR32Float, ast::Access::kWrite);
Global("test_var", s, Global("test_var", s,
ast::DecorationList{ ast::DecorationList{
@ -833,8 +834,9 @@ TEST_F(BuilderTest_Type, StorageTexture_Generate_1d) {
} }
TEST_F(BuilderTest_Type, StorageTexture_Generate_2d) { TEST_F(BuilderTest_Type, StorageTexture_Generate_2d) {
auto* s = ty.storage_texture(ast::TextureDimension::k2d, auto* s =
ast::ImageFormat::kR32Float, ast::Access::kRead); ty.storage_texture(ast::TextureDimension::k2d,
ast::ImageFormat::kR32Float, ast::Access::kWrite);
Global("test_var", s, Global("test_var", s,
ast::DecorationList{ ast::DecorationList{
@ -852,8 +854,9 @@ TEST_F(BuilderTest_Type, StorageTexture_Generate_2d) {
} }
TEST_F(BuilderTest_Type, StorageTexture_Generate_2dArray) { TEST_F(BuilderTest_Type, StorageTexture_Generate_2dArray) {
auto* s = ty.storage_texture(ast::TextureDimension::k2dArray, auto* s =
ast::ImageFormat::kR32Float, ast::Access::kRead); ty.storage_texture(ast::TextureDimension::k2dArray,
ast::ImageFormat::kR32Float, ast::Access::kWrite);
Global("test_var", s, Global("test_var", s,
ast::DecorationList{ ast::DecorationList{
@ -871,8 +874,9 @@ TEST_F(BuilderTest_Type, StorageTexture_Generate_2dArray) {
} }
TEST_F(BuilderTest_Type, StorageTexture_Generate_3d) { TEST_F(BuilderTest_Type, StorageTexture_Generate_3d) {
auto* s = ty.storage_texture(ast::TextureDimension::k3d, auto* s =
ast::ImageFormat::kR32Float, ast::Access::kRead); ty.storage_texture(ast::TextureDimension::k3d,
ast::ImageFormat::kR32Float, ast::Access::kWrite);
Global("test_var", s, Global("test_var", s,
ast::DecorationList{ ast::DecorationList{
@ -891,8 +895,9 @@ TEST_F(BuilderTest_Type, StorageTexture_Generate_3d) {
TEST_F(BuilderTest_Type, TEST_F(BuilderTest_Type,
StorageTexture_Generate_SampledTypeFloat_Format_r32float) { StorageTexture_Generate_SampledTypeFloat_Format_r32float) {
auto* s = ty.storage_texture(ast::TextureDimension::k2d, auto* s =
ast::ImageFormat::kR32Float, ast::Access::kRead); ty.storage_texture(ast::TextureDimension::k2d,
ast::ImageFormat::kR32Float, ast::Access::kWrite);
Global("test_var", s, Global("test_var", s,
ast::DecorationList{ ast::DecorationList{
@ -912,7 +917,7 @@ TEST_F(BuilderTest_Type,
TEST_F(BuilderTest_Type, TEST_F(BuilderTest_Type,
StorageTexture_Generate_SampledTypeSint_Format_r32sint) { StorageTexture_Generate_SampledTypeSint_Format_r32sint) {
auto* s = ty.storage_texture(ast::TextureDimension::k2d, auto* s = ty.storage_texture(ast::TextureDimension::k2d,
ast::ImageFormat::kR32Sint, ast::Access::kRead); ast::ImageFormat::kR32Sint, ast::Access::kWrite);
Global("test_var", s, Global("test_var", s,
ast::DecorationList{ ast::DecorationList{
@ -932,7 +937,7 @@ TEST_F(BuilderTest_Type,
TEST_F(BuilderTest_Type, TEST_F(BuilderTest_Type,
StorageTexture_Generate_SampledTypeUint_Format_r32uint) { StorageTexture_Generate_SampledTypeUint_Format_r32uint) {
auto* s = ty.storage_texture(ast::TextureDimension::k2d, auto* s = ty.storage_texture(ast::TextureDimension::k2d,
ast::ImageFormat::kR32Uint, ast::Access::kRead); ast::ImageFormat::kR32Uint, ast::Access::kWrite);
Global("test_var", s, Global("test_var", s,
ast::DecorationList{ ast::DecorationList{

12
src/writer/wgsl/generator_impl_type_test.cc
View File

@ -453,18 +453,6 @@ INSTANTIATE_TEST_SUITE_P(
WgslGeneratorImplTest, WgslGeneratorImplTest,
WgslGenerator_StorageTextureTest, WgslGenerator_StorageTextureTest,
testing::Values( testing::Values(
StorageTextureData{ast::ImageFormat::kRgba8Sint,
ast::TextureDimension::k1d, ast::Access::kRead,
"texture_storage_1d<rgba8sint, read>"},
StorageTextureData{ast::ImageFormat::kRgba8Sint,
ast::TextureDimension::k2d, ast::Access::kRead,
"texture_storage_2d<rgba8sint, read>"},
StorageTextureData{ast::ImageFormat::kRgba8Sint,
ast::TextureDimension::k2dArray, ast::Access::kRead,
"texture_storage_2d_array<rgba8sint, read>"},
StorageTextureData{ast::ImageFormat::kRgba8Sint,
ast::TextureDimension::k3d, ast::Access::kRead,
"texture_storage_3d<rgba8sint, read>"},
StorageTextureData{ast::ImageFormat::kRgba8Sint, StorageTextureData{ast::ImageFormat::kRgba8Sint,
ast::TextureDimension::k1d, ast::Access::kWrite, ast::TextureDimension::k1d, ast::Access::kWrite,
"texture_storage_1d<rgba8sint, write>"}, "texture_storage_1d<rgba8sint, write>"},

6
test/bug/tint/1083.wgsl.expected.hlsl
View File

@ -1,9 +1,5 @@
SKIP: FAILED
[numthreads(1, 1, 1)] [numthreads(1, 1, 1)]
void f() { void f() {
const int c = (1 / 0); const int c = (1 / 1);
return; return;
} }
O:\src\chrome\src\third_party\dawn\third_party\tint\test\Shader@0x000002477B82F5C0(3,18-22): error X4010: Unsigned integer divide by zero

1
test/bug/tint/413.spvasm.expected.hlsl
View File

@ -1,4 +1,3 @@
warning: use of deprecated intrinsic
Texture2D<uint4> Src : register(t0, space0); Texture2D<uint4> Src : register(t0, space0);
RWTexture2D<uint4> Dst : register(u1, space0); RWTexture2D<uint4> Dst : register(u1, space0);

7
test/bug/tint/413.spvasm.expected.msl
View File

@ -1,10 +1,9 @@
warning: use of deprecated intrinsic
#include <metal_stdlib> #include <metal_stdlib>
using namespace metal; using namespace metal;
void main_1(texture2d<uint, access::read> tint_symbol_1, texture2d<uint, access::write> tint_symbol_2) { void main_1(texture2d<uint, access::sample> tint_symbol_1, texture2d<uint, access::write> tint_symbol_2) {
uint4 srcValue = 0u; uint4 srcValue = 0u;
uint4 const x_18 = tint_symbol_1.read(uint2(int2(0, 0))); uint4 const x_18 = tint_symbol_1.read(uint2(int2(0, 0)), 0);
srcValue = x_18; srcValue = x_18;
uint const x_22 = srcValue.x; uint const x_22 = srcValue.x;
srcValue.x = (x_22 + as_type<uint>(1)); srcValue.x = (x_22 + as_type<uint>(1));
@ -13,7 +12,7 @@ void main_1(texture2d<uint, access::read> tint_symbol_1, texture2d<uint, access:
return; return;
} }
kernel void tint_symbol(texture2d<uint, access::read> tint_symbol_3 [[texture(0)]], texture2d<uint, access::write> tint_symbol_4 [[texture(1)]]) { kernel void tint_symbol(texture2d<uint, access::sample> tint_symbol_3 [[texture(0)]], texture2d<uint, access::write> tint_symbol_4 [[texture(1)]]) {
main_1(tint_symbol_3, tint_symbol_4); main_1(tint_symbol_3, tint_symbol_4);
return; return;
} }

53
test/bug/tint/413.spvasm.expected.spvasm
View File

@ -1,8 +1,7 @@
warning: use of deprecated intrinsic
; SPIR-V ; SPIR-V
; Version: 1.3 ; Version: 1.3
; Generator: Google Tint Compiler; 0 ; Generator: Google Tint Compiler; 0
; Bound: 35 ; Bound: 36
; Schema: 0 ; Schema: 0
OpCapability Shader OpCapability Shader
OpMemoryModel Logical GLSL450 OpMemoryModel Logical GLSL450
@ -13,49 +12,49 @@ warning: use of deprecated intrinsic
OpName %main_1 "main_1" OpName %main_1 "main_1"
OpName %srcValue "srcValue" OpName %srcValue "srcValue"
OpName %main "main" OpName %main "main"
OpDecorate %Src NonWritable
OpDecorate %Src DescriptorSet 0 OpDecorate %Src DescriptorSet 0
OpDecorate %Src Binding 0 OpDecorate %Src Binding 0
OpDecorate %Dst NonReadable OpDecorate %Dst NonReadable
OpDecorate %Dst DescriptorSet 0 OpDecorate %Dst DescriptorSet 0
OpDecorate %Dst Binding 1 OpDecorate %Dst Binding 1
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%3 = OpTypeImage %uint 2D 0 0 0 2 R32ui %3 = OpTypeImage %uint 2D 0 0 0 1 Unknown
%_ptr_UniformConstant_3 = OpTypePointer UniformConstant %3 %_ptr_UniformConstant_3 = OpTypePointer UniformConstant %3
%Src = OpVariable %_ptr_UniformConstant_3 UniformConstant %Src = OpVariable %_ptr_UniformConstant_3 UniformConstant
%_ptr_UniformConstant_3_0 = OpTypePointer UniformConstant %3 %7 = OpTypeImage %uint 2D 0 0 0 2 R32ui
%Dst = OpVariable %_ptr_UniformConstant_3_0 UniformConstant %_ptr_UniformConstant_7 = OpTypePointer UniformConstant %7
%Dst = OpVariable %_ptr_UniformConstant_7 UniformConstant
%void = OpTypeVoid %void = OpTypeVoid
%7 = OpTypeFunction %void %8 = OpTypeFunction %void
%v4uint = OpTypeVector %uint 4 %v4uint = OpTypeVector %uint 4
%_ptr_Function_v4uint = OpTypePointer Function %v4uint %_ptr_Function_v4uint = OpTypePointer Function %v4uint
%14 = OpConstantNull %v4uint %15 = OpConstantNull %v4uint
%int = OpTypeInt 32 1 %int = OpTypeInt 32 1
%v2int = OpTypeVector %int 2 %v2int = OpTypeVector %int 2
%int_0 = OpConstant %int 0 %int_0 = OpConstant %int 0
%20 = OpConstantComposite %v2int %int_0 %int_0 %21 = OpConstantComposite %v2int %int_0 %int_0
%uint_0 = OpConstant %uint 0 %uint_0 = OpConstant %uint 0
%_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Function_uint = OpTypePointer Function %uint
%int_1 = OpConstant %int 1 %int_1 = OpConstant %int 1
%main_1 = OpFunction %void None %7 %main_1 = OpFunction %void None %8
%10 = OpLabel %11 = OpLabel
%srcValue = OpVariable %_ptr_Function_v4uint Function %14 %srcValue = OpVariable %_ptr_Function_v4uint Function %15
%16 = OpLoad %3 %Src %17 = OpLoad %3 %Src
%15 = OpImageRead %v4uint %16 %20 %16 = OpImageFetch %v4uint %17 %21 Lod %int_0
OpStore %srcValue %15 OpStore %srcValue %16
%23 = OpAccessChain %_ptr_Function_uint %srcValue %uint_0 %24 = OpAccessChain %_ptr_Function_uint %srcValue %uint_0
%24 = OpLoad %uint %23 %25 = OpLoad %uint %24
%25 = OpAccessChain %_ptr_Function_uint %srcValue %uint_0 %26 = OpAccessChain %_ptr_Function_uint %srcValue %uint_0
%26 = OpBitcast %uint %int_1 %27 = OpBitcast %uint %int_1
%28 = OpIAdd %uint %24 %26 %29 = OpIAdd %uint %25 %27
OpStore %25 %28 OpStore %26 %29
%29 = OpLoad %v4uint %srcValue %30 = OpLoad %v4uint %srcValue
%31 = OpLoad %3 %Dst %32 = OpLoad %7 %Dst
OpImageWrite %31 %20 %29 OpImageWrite %32 %21 %30
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd
%main = OpFunction %void None %7 %main = OpFunction %void None %8
%33 = OpLabel %34 = OpLabel
%34 = OpFunctionCall %void %main_1 %35 = OpFunctionCall %void %main_1
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

5
test/bug/tint/413.spvasm.expected.wgsl
View File

@ -1,11 +1,10 @@
warning: use of deprecated intrinsic [[group(0), binding(0)]] var Src : texture_2d<u32>;
[[group(0), binding(0)]] var Src : texture_storage_2d<r32uint, read>;
[[group(0), binding(1)]] var Dst : texture_storage_2d<r32uint, write>; [[group(0), binding(1)]] var Dst : texture_storage_2d<r32uint, write>;
fn main_1() { fn main_1() {
var srcValue : vec4<u32>; var srcValue : vec4<u32>;
let x_18 : vec4<u32> = textureLoad(Src, vec2<i32>(0, 0)); let x_18 : vec4<u32> = textureLoad(Src, vec2<i32>(0, 0), 0);
srcValue = x_18; srcValue = x_18;
let x_21 : ptr<function, u32> = &(srcValue.x); let x_21 : ptr<function, u32> = &(srcValue.x);
let x_22 : u32 = *(x_21); let x_22 : u32 = *(x_21);

4
test/bug/tint/453.wgsl
View File

@ -1,10 +1,10 @@
[[group(0), binding(0)]] var Src : texture_storage_2d<r32uint, read>; [[group(0), binding(0)]] var Src : texture_2d<u32>;
[[group(0), binding(1)]] var Dst : texture_storage_2d<r32uint, write>; [[group(0), binding(1)]] var Dst : texture_storage_2d<r32uint, write>;
[[stage(compute), workgroup_size(1)]] [[stage(compute), workgroup_size(1)]]
fn main() { fn main() {
var srcValue : vec4<u32>; var srcValue : vec4<u32>;
let x_22 : vec4<u32> = textureLoad(Src, vec2<i32>(0, 0)); let x_22 : vec4<u32> = textureLoad(Src, vec2<i32>(0, 0), 0);
srcValue = x_22; srcValue = x_22;
let x_23 : ptr<function, u32> = &srcValue.x; let x_23 : ptr<function, u32> = &srcValue.x;
let x_24 : u32 = *x_23; let x_24 : u32 = *x_23;

4
test/bug/tint/453.wgsl.expected.hlsl
View File

@ -1,7 +1,3 @@
bug/tint/453.wgsl:7:26 warning: use of deprecated intrinsic
let x_22 : vec4<u32> = textureLoad(Src, vec2<i32>(0, 0));
^^^^^^^^^^^
Texture2D<uint4> Src : register(t0, space0); Texture2D<uint4> Src : register(t0, space0);
RWTexture2D<uint4> Dst : register(u1, space0); RWTexture2D<uint4> Dst : register(u1, space0);

8
test/bug/tint/453.wgsl.expected.msl
View File

@ -1,13 +1,9 @@
bug/tint/453.wgsl:7:26 warning: use of deprecated intrinsic
let x_22 : vec4<u32> = textureLoad(Src, vec2<i32>(0, 0));
^^^^^^^^^^^
#include <metal_stdlib> #include <metal_stdlib>
using namespace metal; using namespace metal;
kernel void tint_symbol(texture2d<uint, access::read> tint_symbol_1 [[texture(0)]], texture2d<uint, access::write> tint_symbol_2 [[texture(1)]]) { kernel void tint_symbol(texture2d<uint, access::sample> tint_symbol_1 [[texture(0)]], texture2d<uint, access::write> tint_symbol_2 [[texture(1)]]) {
uint4 srcValue = 0u; uint4 srcValue = 0u;
uint4 const x_22 = tint_symbol_1.read(uint2(int2(0, 0))); uint4 const x_22 = tint_symbol_1.read(uint2(int2(0, 0)), 0);
srcValue = x_22; srcValue = x_22;
uint const x_24 = srcValue.x; uint const x_24 = srcValue.x;
uint const x_25 = (x_24 + 1u); uint const x_25 = (x_24 + 1u);

46
test/bug/tint/453.wgsl.expected.spvasm
View File

@ -1,11 +1,7 @@
bug/tint/453.wgsl:7:26 warning: use of deprecated intrinsic
let x_22 : vec4<u32> = textureLoad(Src, vec2<i32>(0, 0));
^^^^^^^^^^^
; SPIR-V ; SPIR-V
; Version: 1.3 ; Version: 1.3
; Generator: Google Tint Compiler; 0 ; Generator: Google Tint Compiler; 0
; Bound: 31 ; Bound: 32
; Schema: 0 ; Schema: 0
OpCapability Shader OpCapability Shader
OpMemoryModel Logical GLSL450 OpMemoryModel Logical GLSL450
@ -15,42 +11,42 @@ bug/tint/453.wgsl:7:26 warning: use of deprecated intrinsic
OpName %Dst "Dst" OpName %Dst "Dst"
OpName %main "main" OpName %main "main"
OpName %srcValue "srcValue" OpName %srcValue "srcValue"
OpDecorate %Src NonWritable
OpDecorate %Src DescriptorSet 0 OpDecorate %Src DescriptorSet 0
OpDecorate %Src Binding 0 OpDecorate %Src Binding 0
OpDecorate %Dst NonReadable OpDecorate %Dst NonReadable
OpDecorate %Dst DescriptorSet 0 OpDecorate %Dst DescriptorSet 0
OpDecorate %Dst Binding 1 OpDecorate %Dst Binding 1
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%3 = OpTypeImage %uint 2D 0 0 0 2 R32ui %3 = OpTypeImage %uint 2D 0 0 0 1 Unknown
%_ptr_UniformConstant_3 = OpTypePointer UniformConstant %3 %_ptr_UniformConstant_3 = OpTypePointer UniformConstant %3
%Src = OpVariable %_ptr_UniformConstant_3 UniformConstant %Src = OpVariable %_ptr_UniformConstant_3 UniformConstant
%_ptr_UniformConstant_3_0 = OpTypePointer UniformConstant %3 %7 = OpTypeImage %uint 2D 0 0 0 2 R32ui
%Dst = OpVariable %_ptr_UniformConstant_3_0 UniformConstant %_ptr_UniformConstant_7 = OpTypePointer UniformConstant %7
%Dst = OpVariable %_ptr_UniformConstant_7 UniformConstant
%void = OpTypeVoid %void = OpTypeVoid
%7 = OpTypeFunction %void %8 = OpTypeFunction %void
%v4uint = OpTypeVector %uint 4 %v4uint = OpTypeVector %uint 4
%_ptr_Function_v4uint = OpTypePointer Function %v4uint %_ptr_Function_v4uint = OpTypePointer Function %v4uint
%14 = OpConstantNull %v4uint %15 = OpConstantNull %v4uint
%int = OpTypeInt 32 1 %int = OpTypeInt 32 1
%v2int = OpTypeVector %int 2 %v2int = OpTypeVector %int 2
%int_0 = OpConstant %int 0 %int_0 = OpConstant %int 0
%20 = OpConstantComposite %v2int %int_0 %int_0 %21 = OpConstantComposite %v2int %int_0 %int_0
%uint_0 = OpConstant %uint 0 %uint_0 = OpConstant %uint 0
%_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Function_uint = OpTypePointer Function %uint
%uint_1 = OpConstant %uint 1 %uint_1 = OpConstant %uint 1
%main = OpFunction %void None %7 %main = OpFunction %void None %8
%10 = OpLabel %11 = OpLabel
%srcValue = OpVariable %_ptr_Function_v4uint Function %14 %srcValue = OpVariable %_ptr_Function_v4uint Function %15
%16 = OpLoad %3 %Src %17 = OpLoad %3 %Src
%15 = OpImageRead %v4uint %16 %20 %16 = OpImageFetch %v4uint %17 %21 Lod %int_0
OpStore %srcValue %15 OpStore %srcValue %16
%23 = OpAccessChain %_ptr_Function_uint %srcValue %uint_0 %24 = OpAccessChain %_ptr_Function_uint %srcValue %uint_0
%24 = OpLoad %uint %23 %25 = OpLoad %uint %24
%26 = OpIAdd %uint %24 %uint_1 %27 = OpIAdd %uint %25 %uint_1
%27 = OpLoad %v4uint %srcValue %28 = OpLoad %v4uint %srcValue
%29 = OpLoad %3 %Dst %30 = OpLoad %7 %Dst
%30 = OpVectorShuffle %v4uint %27 %27 0 0 0 0 %31 = OpVectorShuffle %v4uint %28 %28 0 0 0 0
OpImageWrite %29 %20 %30 OpImageWrite %30 %21 %31
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

8
test/bug/tint/453.wgsl.expected.wgsl
View File

@ -1,15 +1,11 @@
bug/tint/453.wgsl:7:26 warning: use of deprecated intrinsic [[group(0), binding(0)]] var Src : texture_2d<u32>;
let x_22 : vec4<u32> = textureLoad(Src, vec2<i32>(0, 0));
^^^^^^^^^^^
[[group(0), binding(0)]] var Src : texture_storage_2d<r32uint, read>;
[[group(0), binding(1)]] var Dst : texture_storage_2d<r32uint, write>; [[group(0), binding(1)]] var Dst : texture_storage_2d<r32uint, write>;
[[stage(compute), workgroup_size(1)]] [[stage(compute), workgroup_size(1)]]
fn main() { fn main() {
var srcValue : vec4<u32>; var srcValue : vec4<u32>;
let x_22 : vec4<u32> = textureLoad(Src, vec2<i32>(0, 0)); let x_22 : vec4<u32> = textureLoad(Src, vec2<i32>(0, 0), 0);
srcValue = x_22; srcValue = x_22;
let x_23 : ptr<function, u32> = &(srcValue.x); let x_23 : ptr<function, u32> = &(srcValue.x);
let x_24 : u32 = *(x_23); let x_24 : u32 = *(x_23);

46
test/intrinsics/gen/textureDimensions/08a62e.wgsl
View File

@ -1,46 +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/intrinsic-gen
// using the template:
// test/intrinsics/intrinsics.wgsl.tmpl
// and the intrinsic defintion file:
// src/intrinsics.def
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
[[group(1), binding(0)]] var arg_0: texture_storage_1d<rgba32sint, read>;
// fn textureDimensions(texture: texture_storage_1d<rgba32sint, read>) -> i32
fn textureDimensions_08a62e() {
var res: i32 = textureDimensions(arg_0);
}
[[stage(vertex)]]
fn vertex_main() -> [[builtin(position)]] vec4<f32> {
textureDimensions_08a62e();
return vec4<f32>();
}
[[stage(fragment)]]
fn fragment_main() {
textureDimensions_08a62e();
}
[[stage(compute), workgroup_size(1)]]
fn compute_main() {
textureDimensions_08a62e();
}

34
test/intrinsics/gen/textureDimensions/08a62e.wgsl.expected.hlsl
View File

@ -1,34 +0,0 @@
Texture1D<int4> arg_0 : register(t0, space1);
void textureDimensions_08a62e() {
int tint_tmp;
arg_0.GetDimensions(tint_tmp);
int res = tint_tmp;
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
textureDimensions_08a62e();
return float4(0.0f, 0.0f, 0.0f, 0.0f);
}
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_08a62e();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
textureDimensions_08a62e();
return;
}

33
test/intrinsics/gen/textureDimensions/08a62e.wgsl.expected.msl
View File

@ -1,33 +0,0 @@
#include <metal_stdlib>
using namespace metal;
struct tint_symbol {
float4 value [[position]];
};
void textureDimensions_08a62e(texture1d<int, access::read> tint_symbol_1) {
int res = int(tint_symbol_1.get_width());
}
float4 vertex_main_inner(texture1d<int, access::read> tint_symbol_2) {
textureDimensions_08a62e(tint_symbol_2);
return float4();
}
vertex tint_symbol vertex_main(texture1d<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(texture1d<int, access::read> tint_symbol_4 [[texture(0)]]) {
textureDimensions_08a62e(tint_symbol_4);
return;
}
kernel void compute_main(texture1d<int, access::read> tint_symbol_5 [[texture(0)]]) {
textureDimensions_08a62e(tint_symbol_5);
return;
}

76
test/intrinsics/gen/textureDimensions/08a62e.wgsl.expected.spvasm
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_08a62e "textureDimensions_08a62e"
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 NonWritable
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 Rgba32i
%_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_08a62e = 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_08a62e
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_08a62e
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %13
%34 = OpLabel
%35 = OpFunctionCall %void %textureDimensions_08a62e
OpReturn
OpFunctionEnd

21
test/intrinsics/gen/textureDimensions/08a62e.wgsl.expected.wgsl
View File

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

46
test/intrinsics/gen/textureDimensions/0a5fcf.wgsl
View File

@ -1,46 +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/intrinsic-gen
// using the template:
// test/intrinsics/intrinsics.wgsl.tmpl
// and the intrinsic defintion file:
// src/intrinsics.def
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
[[group(1), binding(0)]] var arg_0: texture_storage_2d_array<rgba16float, read>;
// fn textureDimensions(texture: texture_storage_2d_array<rgba16float, read>) -> vec2<i32>
fn textureDimensions_0a5fcf() {
var res: vec2<i32> = textureDimensions(arg_0);
}
[[stage(vertex)]]
fn vertex_main() -> [[builtin(position)]] vec4<f32> {
textureDimensions_0a5fcf();
return vec4<f32>();
}
[[stage(fragment)]]
fn fragment_main() {
textureDimensions_0a5fcf();
}
[[stage(compute), workgroup_size(1)]]
fn compute_main() {
textureDimensions_0a5fcf();
}

34
test/intrinsics/gen/textureDimensions/0a5fcf.wgsl.expected.hlsl
View File

@ -1,34 +0,0 @@
Texture2DArray<float4> arg_0 : register(t0, space1);
void textureDimensions_0a5fcf() {
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_0a5fcf();
return float4(0.0f, 0.0f, 0.0f, 0.0f);
}
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_0a5fcf();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
textureDimensions_0a5fcf();
return;
}

33
test/intrinsics/gen/textureDimensions/0a5fcf.wgsl.expected.msl
View File

@ -1,33 +0,0 @@
#include <metal_stdlib>
using namespace metal;
struct tint_symbol {
float4 value [[position]];
};
void textureDimensions_0a5fcf(texture2d_array<float, access::read> tint_symbol_1) {
int2 res = int2(tint_symbol_1.get_width(), tint_symbol_1.get_height());
}
float4 vertex_main_inner(texture2d_array<float, access::read> tint_symbol_2) {
textureDimensions_0a5fcf(tint_symbol_2);
return float4();
}
vertex tint_symbol vertex_main(texture2d_array<float, 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_array<float, access::read> tint_symbol_4 [[texture(0)]]) {
textureDimensions_0a5fcf(tint_symbol_4);
return;
}
kernel void compute_main(texture2d_array<float, access::read> tint_symbol_5 [[texture(0)]]) {
textureDimensions_0a5fcf(tint_symbol_5);
return;
}

78
test/intrinsics/gen/textureDimensions/0a5fcf.wgsl.expected.spvasm
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_0a5fcf "textureDimensions_0a5fcf"
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 NonWritable
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 2 Rgba16f
%_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
%_ptr_Function_v2int = OpTypePointer Function %v2int
%24 = OpConstantNull %v2int
%25 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%textureDimensions_0a5fcf = OpFunction %void None %12
%15 = OpLabel
%res = OpVariable %_ptr_Function_v2int Function %24
%21 = OpLoad %11 %arg_0
%19 = OpImageQuerySize %v3int %21
%16 = OpVectorShuffle %v2int %19 %19 0 1
OpStore %res %16
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %25
%27 = OpLabel
%28 = OpFunctionCall %void %textureDimensions_0a5fcf
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %12
%30 = OpLabel
%31 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %31
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %12
%34 = OpLabel
%35 = OpFunctionCall %void %textureDimensions_0a5fcf
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %12
%37 = OpLabel
%38 = OpFunctionCall %void %textureDimensions_0a5fcf
OpReturn
OpFunctionEnd

21
test/intrinsics/gen/textureDimensions/0a5fcf.wgsl.expected.wgsl
View File

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

46
test/intrinsics/gen/textureDimensions/0bab57.wgsl
View File

@ -1,46 +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/intrinsic-gen
// using the template:
// test/intrinsics/intrinsics.wgsl.tmpl
// and the intrinsic defintion file:
// src/intrinsics.def
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
[[group(1), binding(0)]] var arg_0: texture_storage_3d<rgba16sint, read>;
// fn textureDimensions(texture: texture_storage_3d<rgba16sint, read>) -> vec3<i32>
fn textureDimensions_0bab57() {
var res: vec3<i32> = textureDimensions(arg_0);
}
[[stage(vertex)]]
fn vertex_main() -> [[builtin(position)]] vec4<f32> {
textureDimensions_0bab57();
return vec4<f32>();
}
[[stage(fragment)]]
fn fragment_main() {
textureDimensions_0bab57();
}
[[stage(compute), workgroup_size(1)]]
fn compute_main() {
textureDimensions_0bab57();
}

34
test/intrinsics/gen/textureDimensions/0bab57.wgsl.expected.hlsl
View File

@ -1,34 +0,0 @@
Texture3D<int4> arg_0 : register(t0, space1);
void textureDimensions_0bab57() {
int3 tint_tmp;
arg_0.GetDimensions(tint_tmp.x, tint_tmp.y, tint_tmp.z);
int3 res = tint_tmp;
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
textureDimensions_0bab57();
return float4(0.0f, 0.0f, 0.0f, 0.0f);
}
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_0bab57();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
textureDimensions_0bab57();
return;
}

33
test/intrinsics/gen/textureDimensions/0bab57.wgsl.expected.msl
View File

@ -1,33 +0,0 @@
#include <metal_stdlib>
using namespace metal;
struct tint_symbol {
float4 value [[position]];
};
void textureDimensions_0bab57(texture3d<int, access::read> tint_symbol_1) {
int3 res = int3(tint_symbol_1.get_width(), tint_symbol_1.get_height(), tint_symbol_1.get_depth());
}
float4 vertex_main_inner(texture3d<int, access::read> tint_symbol_2) {
textureDimensions_0bab57(tint_symbol_2);
return float4();
}
vertex tint_symbol vertex_main(texture3d<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(texture3d<int, access::read> tint_symbol_4 [[texture(0)]]) {
textureDimensions_0bab57(tint_symbol_4);
return;
}
kernel void compute_main(texture3d<int, access::read> tint_symbol_5 [[texture(0)]]) {
textureDimensions_0bab57(tint_symbol_5);
return;
}

76
test/intrinsics/gen/textureDimensions/0bab57.wgsl.expected.spvasm
View File

@ -1,76 +0,0 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 37
; 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_0bab57 "textureDimensions_0bab57"
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 NonWritable
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 3D 0 0 0 2 Rgba16i
%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
%arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant
%void = OpTypeVoid
%13 = OpTypeFunction %void
%v3int = OpTypeVector %int 3
%_ptr_Function_v3int = OpTypePointer Function %v3int
%22 = OpConstantNull %v3int
%23 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%textureDimensions_0bab57 = OpFunction %void None %13
%16 = OpLabel
%res = OpVariable %_ptr_Function_v3int Function %22
%19 = OpLoad %11 %arg_0
%17 = OpImageQuerySize %v3int %19
OpStore %res %17
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %23
%25 = OpLabel
%26 = OpFunctionCall %void %textureDimensions_0bab57
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %13
%28 = OpLabel
%29 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %29
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %13
%32 = OpLabel
%33 = OpFunctionCall %void %textureDimensions_0bab57
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %13
%35 = OpLabel
%36 = OpFunctionCall %void %textureDimensions_0bab57
OpReturn
OpFunctionEnd

21
test/intrinsics/gen/textureDimensions/0bab57.wgsl.expected.wgsl
View File

@ -1,21 +0,0 @@
[[group(1), binding(0)]] var arg_0 : texture_storage_3d<rgba16sint, read>;
fn textureDimensions_0bab57() {
var res : vec3<i32> = textureDimensions(arg_0);
}
[[stage(vertex)]]
fn vertex_main() -> [[builtin(position)]] vec4<f32> {
textureDimensions_0bab57();
return vec4<f32>();
}
[[stage(fragment)]]
fn fragment_main() {
textureDimensions_0bab57();
}
[[stage(compute), workgroup_size(1)]]
fn compute_main() {
textureDimensions_0bab57();
}

46
test/intrinsics/gen/textureDimensions/1147d6.wgsl
View File

@ -1,46 +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/intrinsic-gen
// using the template:
// test/intrinsics/intrinsics.wgsl.tmpl
// and the intrinsic defintion file:
// src/intrinsics.def
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
[[group(1), binding(0)]] var arg_0: texture_storage_2d_array<rgba16sint, read>;
// fn textureDimensions(texture: texture_storage_2d_array<rgba16sint, read>) -> vec2<i32>
fn textureDimensions_1147d6() {
var res: vec2<i32> = textureDimensions(arg_0);
}
[[stage(vertex)]]
fn vertex_main() -> [[builtin(position)]] vec4<f32> {
textureDimensions_1147d6();
return vec4<f32>();
}
[[stage(fragment)]]
fn fragment_main() {
textureDimensions_1147d6();
}
[[stage(compute), workgroup_size(1)]]
fn compute_main() {
textureDimensions_1147d6();
}

34
test/intrinsics/gen/textureDimensions/1147d6.wgsl.expected.hlsl
View File

@ -1,34 +0,0 @@
Texture2DArray<int4> arg_0 : register(t0, space1);
void textureDimensions_1147d6() {
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_1147d6();
return float4(0.0f, 0.0f, 0.0f, 0.0f);
}
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_1147d6();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
textureDimensions_1147d6();
return;
}

33
test/intrinsics/gen/textureDimensions/1147d6.wgsl.expected.msl
View File

@ -1,33 +0,0 @@
#include <metal_stdlib>
using namespace metal;
struct tint_symbol {
float4 value [[position]];
};
void textureDimensions_1147d6(texture2d_array<int, access::read> tint_symbol_1) {
int2 res = int2(tint_symbol_1.get_width(), tint_symbol_1.get_height());
}
float4 vertex_main_inner(texture2d_array<int, access::read> tint_symbol_2) {
textureDimensions_1147d6(tint_symbol_2);
return float4();
}
vertex tint_symbol vertex_main(texture2d_array<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_array<int, access::read> tint_symbol_4 [[texture(0)]]) {
textureDimensions_1147d6(tint_symbol_4);
return;
}
kernel void compute_main(texture2d_array<int, access::read> tint_symbol_5 [[texture(0)]]) {
textureDimensions_1147d6(tint_symbol_5);
return;
}

78
test/intrinsics/gen/textureDimensions/1147d6.wgsl.expected.spvasm
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_1147d6 "textureDimensions_1147d6"
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 NonWritable
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 Rgba16i
%_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_1147d6 = 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_1147d6
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_1147d6
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %13
%37 = OpLabel
%38 = OpFunctionCall %void %textureDimensions_1147d6
OpReturn
OpFunctionEnd

21
test/intrinsics/gen/textureDimensions/1147d6.wgsl.expected.wgsl
View File

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

46
test/intrinsics/gen/textureDimensions/168fcc.wgsl
View File

@ -1,46 +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/intrinsic-gen
// using the template:
// test/intrinsics/intrinsics.wgsl.tmpl
// and the intrinsic defintion file:
// src/intrinsics.def
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
[[group(1), binding(0)]] var arg_0: texture_storage_2d<rgba8unorm, read>;
// fn textureDimensions(texture: texture_storage_2d<rgba8unorm, read>) -> vec2<i32>
fn textureDimensions_168fcc() {
var res: vec2<i32> = textureDimensions(arg_0);
}
[[stage(vertex)]]
fn vertex_main() -> [[builtin(position)]] vec4<f32> {
textureDimensions_168fcc();
return vec4<f32>();
}
[[stage(fragment)]]
fn fragment_main() {
textureDimensions_168fcc();
}
[[stage(compute), workgroup_size(1)]]
fn compute_main() {
textureDimensions_168fcc();
}

34
test/intrinsics/gen/textureDimensions/168fcc.wgsl.expected.hlsl
View File

@ -1,34 +0,0 @@
Texture2D<float4> arg_0 : register(t0, space1);
void textureDimensions_168fcc() {
int2 tint_tmp;
arg_0.GetDimensions(tint_tmp.x, tint_tmp.y);
int2 res = tint_tmp;
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
textureDimensions_168fcc();
return float4(0.0f, 0.0f, 0.0f, 0.0f);
}
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_168fcc();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
textureDimensions_168fcc();
return;
}

33
test/intrinsics/gen/textureDimensions/168fcc.wgsl.expected.msl
View File

@ -1,33 +0,0 @@
#include <metal_stdlib>
using namespace metal;
struct tint_symbol {
float4 value [[position]];
};
void textureDimensions_168fcc(texture2d<float, access::read> tint_symbol_1) {
int2 res = int2(tint_symbol_1.get_width(), tint_symbol_1.get_height());
}
float4 vertex_main_inner(texture2d<float, access::read> tint_symbol_2) {
textureDimensions_168fcc(tint_symbol_2);
return float4();
}
vertex tint_symbol vertex_main(texture2d<float, 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<float, access::read> tint_symbol_4 [[texture(0)]]) {
textureDimensions_168fcc(tint_symbol_4);
return;
}
kernel void compute_main(texture2d<float, access::read> tint_symbol_5 [[texture(0)]]) {
textureDimensions_168fcc(tint_symbol_5);
return;
}

76
test/intrinsics/gen/textureDimensions/168fcc.wgsl.expected.spvasm
View File

@ -1,76 +0,0 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 37
; 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_168fcc "textureDimensions_168fcc"
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 NonWritable
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 0 0 2 Rgba8
%_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
%_ptr_Function_v2int = OpTypePointer Function %v2int
%22 = OpConstantNull %v2int
%23 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%textureDimensions_168fcc = OpFunction %void None %12
%15 = OpLabel
%res = OpVariable %_ptr_Function_v2int Function %22
%19 = OpLoad %11 %arg_0
%16 = OpImageQuerySize %v2int %19
OpStore %res %16
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %23
%25 = OpLabel
%26 = OpFunctionCall %void %textureDimensions_168fcc
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_168fcc
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %12
%35 = OpLabel
%36 = OpFunctionCall %void %textureDimensions_168fcc
OpReturn
OpFunctionEnd

21
test/intrinsics/gen/textureDimensions/168fcc.wgsl.expected.wgsl
View File

@ -1,21 +0,0 @@
[[group(1), binding(0)]] var arg_0 : texture_storage_2d<rgba8unorm, read>;
fn textureDimensions_168fcc() {
var res : vec2<i32> = textureDimensions(arg_0);
}
[[stage(vertex)]]
fn vertex_main() -> [[builtin(position)]] vec4<f32> {
textureDimensions_168fcc();
return vec4<f32>();
}
[[stage(fragment)]]
fn fragment_main() {
textureDimensions_168fcc();
}
[[stage(compute), workgroup_size(1)]]
fn compute_main() {
textureDimensions_168fcc();
}

46
test/intrinsics/gen/textureDimensions/18bd57.wgsl
View File

@ -1,46 +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/intrinsic-gen
// using the template:
// test/intrinsics/intrinsics.wgsl.tmpl
// and the intrinsic defintion file:
// src/intrinsics.def
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
[[group(1), binding(0)]] var arg_0: texture_storage_2d_array<rg32uint, read>;
// fn textureDimensions(texture: texture_storage_2d_array<rg32uint, read>) -> vec2<i32>
fn textureDimensions_18bd57() {
var res: vec2<i32> = textureDimensions(arg_0);
}
[[stage(vertex)]]
fn vertex_main() -> [[builtin(position)]] vec4<f32> {
textureDimensions_18bd57();
return vec4<f32>();
}
[[stage(fragment)]]
fn fragment_main() {
textureDimensions_18bd57();
}
[[stage(compute), workgroup_size(1)]]
fn compute_main() {
textureDimensions_18bd57();
}

34
test/intrinsics/gen/textureDimensions/18bd57.wgsl.expected.hlsl
View File

@ -1,34 +0,0 @@
Texture2DArray<uint4> arg_0 : register(t0, space1);
void textureDimensions_18bd57() {
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_18bd57();
return float4(0.0f, 0.0f, 0.0f, 0.0f);
}
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_18bd57();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
textureDimensions_18bd57();
return;
}

33
test/intrinsics/gen/textureDimensions/18bd57.wgsl.expected.msl
View File

@ -1,33 +0,0 @@
#include <metal_stdlib>
using namespace metal;
struct tint_symbol {
float4 value [[position]];
};
void textureDimensions_18bd57(texture2d_array<uint, access::read> tint_symbol_1) {
int2 res = int2(tint_symbol_1.get_width(), tint_symbol_1.get_height());
}
float4 vertex_main_inner(texture2d_array<uint, access::read> tint_symbol_2) {
textureDimensions_18bd57(tint_symbol_2);
return float4();
}
vertex tint_symbol vertex_main(texture2d_array<uint, 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_array<uint, access::read> tint_symbol_4 [[texture(0)]]) {
textureDimensions_18bd57(tint_symbol_4);
return;
}
kernel void compute_main(texture2d_array<uint, access::read> tint_symbol_5 [[texture(0)]]) {
textureDimensions_18bd57(tint_symbol_5);
return;
}

80
test/intrinsics/gen/textureDimensions/18bd57.wgsl.expected.spvasm
View File

@ -1,80 +0,0 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 40
; Schema: 0
OpCapability Shader
OpCapability StorageImageExtendedFormats
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_18bd57 "textureDimensions_18bd57"
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 NonWritable
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 2D 0 1 0 2 Rg32ui
%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
%arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant
%void = OpTypeVoid
%13 = OpTypeFunction %void
%int = OpTypeInt 32 1
%v2int = OpTypeVector %int 2
%v3int = OpTypeVector %int 3
%_ptr_Function_v2int = OpTypePointer Function %v2int
%25 = OpConstantNull %v2int
%26 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%textureDimensions_18bd57 = OpFunction %void None %13
%16 = OpLabel
%res = OpVariable %_ptr_Function_v2int Function %25
%22 = OpLoad %11 %arg_0
%20 = OpImageQuerySize %v3int %22
%17 = OpVectorShuffle %v2int %20 %20 0 1
OpStore %res %17
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %26
%28 = OpLabel
%29 = OpFunctionCall %void %textureDimensions_18bd57
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %13
%31 = OpLabel
%32 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %32
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %13
%35 = OpLabel
%36 = OpFunctionCall %void %textureDimensions_18bd57
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %13
%38 = OpLabel
%39 = OpFunctionCall %void %textureDimensions_18bd57
OpReturn
OpFunctionEnd

21
test/intrinsics/gen/textureDimensions/18bd57.wgsl.expected.wgsl
View File

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

46
test/intrinsics/gen/textureDimensions/19bffc.wgsl
View File

@ -1,46 +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/intrinsic-gen
// using the template:
// test/intrinsics/intrinsics.wgsl.tmpl
// and the intrinsic defintion file:
// src/intrinsics.def
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
[[group(1), binding(0)]] var arg_0: texture_storage_3d<rg32sint, read>;
// fn textureDimensions(texture: texture_storage_3d<rg32sint, read>) -> vec3<i32>
fn textureDimensions_19bffc() {
var res: vec3<i32> = textureDimensions(arg_0);
}
[[stage(vertex)]]
fn vertex_main() -> [[builtin(position)]] vec4<f32> {
textureDimensions_19bffc();
return vec4<f32>();
}
[[stage(fragment)]]
fn fragment_main() {
textureDimensions_19bffc();
}
[[stage(compute), workgroup_size(1)]]
fn compute_main() {
textureDimensions_19bffc();
}

34
test/intrinsics/gen/textureDimensions/19bffc.wgsl.expected.hlsl
View File

@ -1,34 +0,0 @@
Texture3D<int4> arg_0 : register(t0, space1);
void textureDimensions_19bffc() {
int3 tint_tmp;
arg_0.GetDimensions(tint_tmp.x, tint_tmp.y, tint_tmp.z);
int3 res = tint_tmp;
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
textureDimensions_19bffc();
return float4(0.0f, 0.0f, 0.0f, 0.0f);
}
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_19bffc();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
textureDimensions_19bffc();
return;
}

33
test/intrinsics/gen/textureDimensions/19bffc.wgsl.expected.msl
View File

@ -1,33 +0,0 @@
#include <metal_stdlib>
using namespace metal;
struct tint_symbol {
float4 value [[position]];
};
void textureDimensions_19bffc(texture3d<int, access::read> tint_symbol_1) {
int3 res = int3(tint_symbol_1.get_width(), tint_symbol_1.get_height(), tint_symbol_1.get_depth());
}
float4 vertex_main_inner(texture3d<int, access::read> tint_symbol_2) {
textureDimensions_19bffc(tint_symbol_2);
return float4();
}
vertex tint_symbol vertex_main(texture3d<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(texture3d<int, access::read> tint_symbol_4 [[texture(0)]]) {
textureDimensions_19bffc(tint_symbol_4);
return;
}
kernel void compute_main(texture3d<int, access::read> tint_symbol_5 [[texture(0)]]) {
textureDimensions_19bffc(tint_symbol_5);
return;
}

77
test/intrinsics/gen/textureDimensions/19bffc.wgsl.expected.spvasm
View File

@ -1,77 +0,0 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 37
; Schema: 0
OpCapability Shader
OpCapability StorageImageExtendedFormats
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_19bffc "textureDimensions_19bffc"
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 NonWritable
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 3D 0 0 0 2 Rg32i
%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
%arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant
%void = OpTypeVoid
%13 = OpTypeFunction %void
%v3int = OpTypeVector %int 3
%_ptr_Function_v3int = OpTypePointer Function %v3int
%22 = OpConstantNull %v3int
%23 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%textureDimensions_19bffc = OpFunction %void None %13
%16 = OpLabel
%res = OpVariable %_ptr_Function_v3int Function %22
%19 = OpLoad %11 %arg_0
%17 = OpImageQuerySize %v3int %19
OpStore %res %17
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %23
%25 = OpLabel
%26 = OpFunctionCall %void %textureDimensions_19bffc
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %13
%28 = OpLabel
%29 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %29
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %13
%32 = OpLabel
%33 = OpFunctionCall %void %textureDimensions_19bffc
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %13
%35 = OpLabel
%36 = OpFunctionCall %void %textureDimensions_19bffc
OpReturn
OpFunctionEnd

21
test/intrinsics/gen/textureDimensions/19bffc.wgsl.expected.wgsl
View File

@ -1,21 +0,0 @@
[[group(1), binding(0)]] var arg_0 : texture_storage_3d<rg32sint, read>;
fn textureDimensions_19bffc() {
var res : vec3<i32> = textureDimensions(arg_0);
}
[[stage(vertex)]]
fn vertex_main() -> [[builtin(position)]] vec4<f32> {
textureDimensions_19bffc();
return vec4<f32>();
}
[[stage(fragment)]]
fn fragment_main() {
textureDimensions_19bffc();
}
[[stage(compute), workgroup_size(1)]]
fn compute_main() {
textureDimensions_19bffc();
}

46
test/intrinsics/gen/textureDimensions/1a58e7.wgsl
View File

@ -1,46 +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/intrinsic-gen
// using the template:
// test/intrinsics/intrinsics.wgsl.tmpl
// and the intrinsic defintion file:
// src/intrinsics.def
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
[[group(1), binding(0)]] var arg_0: texture_storage_2d_array<rgba8unorm, read>;
// fn textureDimensions(texture: texture_storage_2d_array<rgba8unorm, read>) -> vec2<i32>
fn textureDimensions_1a58e7() {
var res: vec2<i32> = textureDimensions(arg_0);
}
[[stage(vertex)]]
fn vertex_main() -> [[builtin(position)]] vec4<f32> {
textureDimensions_1a58e7();
return vec4<f32>();
}
[[stage(fragment)]]
fn fragment_main() {
textureDimensions_1a58e7();
}
[[stage(compute), workgroup_size(1)]]
fn compute_main() {
textureDimensions_1a58e7();
}

34
test/intrinsics/gen/textureDimensions/1a58e7.wgsl.expected.hlsl
View File

@ -1,34 +0,0 @@
Texture2DArray<float4> arg_0 : register(t0, space1);
void textureDimensions_1a58e7() {
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_1a58e7();
return float4(0.0f, 0.0f, 0.0f, 0.0f);
}
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_1a58e7();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
textureDimensions_1a58e7();
return;
}

33
test/intrinsics/gen/textureDimensions/1a58e7.wgsl.expected.msl
View File

@ -1,33 +0,0 @@
#include <metal_stdlib>
using namespace metal;
struct tint_symbol {
float4 value [[position]];
};
void textureDimensions_1a58e7(texture2d_array<float, access::read> tint_symbol_1) {
int2 res = int2(tint_symbol_1.get_width(), tint_symbol_1.get_height());
}
float4 vertex_main_inner(texture2d_array<float, access::read> tint_symbol_2) {
textureDimensions_1a58e7(tint_symbol_2);
return float4();
}
vertex tint_symbol vertex_main(texture2d_array<float, 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_array<float, access::read> tint_symbol_4 [[texture(0)]]) {
textureDimensions_1a58e7(tint_symbol_4);
return;
}
kernel void compute_main(texture2d_array<float, access::read> tint_symbol_5 [[texture(0)]]) {
textureDimensions_1a58e7(tint_symbol_5);
return;
}

78
test/intrinsics/gen/textureDimensions/1a58e7.wgsl.expected.spvasm
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_1a58e7 "textureDimensions_1a58e7"
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 NonWritable
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 2 Rgba8
%_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
%_ptr_Function_v2int = OpTypePointer Function %v2int
%24 = OpConstantNull %v2int
%25 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%textureDimensions_1a58e7 = OpFunction %void None %12
%15 = OpLabel
%res = OpVariable %_ptr_Function_v2int Function %24
%21 = OpLoad %11 %arg_0
%19 = OpImageQuerySize %v3int %21
%16 = OpVectorShuffle %v2int %19 %19 0 1
OpStore %res %16
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %25
%27 = OpLabel
%28 = OpFunctionCall %void %textureDimensions_1a58e7
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %12
%30 = OpLabel
%31 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %31
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %12
%34 = OpLabel
%35 = OpFunctionCall %void %textureDimensions_1a58e7
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %12
%37 = OpLabel
%38 = OpFunctionCall %void %textureDimensions_1a58e7
OpReturn
OpFunctionEnd

21
test/intrinsics/gen/textureDimensions/1a58e7.wgsl.expected.wgsl
View File

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

46
test/intrinsics/gen/textureDimensions/1aa199.wgsl
View File

@ -1,46 +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/intrinsic-gen
// using the template:
// test/intrinsics/intrinsics.wgsl.tmpl
// and the intrinsic defintion file:
// src/intrinsics.def
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
[[group(1), binding(0)]] var arg_0: texture_storage_2d_array<rgba32sint, read>;
// fn textureDimensions(texture: texture_storage_2d_array<rgba32sint, read>) -> vec2<i32>
fn textureDimensions_1aa199() {
var res: vec2<i32> = textureDimensions(arg_0);
}
[[stage(vertex)]]
fn vertex_main() -> [[builtin(position)]] vec4<f32> {
textureDimensions_1aa199();
return vec4<f32>();
}
[[stage(fragment)]]
fn fragment_main() {
textureDimensions_1aa199();
}
[[stage(compute), workgroup_size(1)]]
fn compute_main() {
textureDimensions_1aa199();
}

34
test/intrinsics/gen/textureDimensions/1aa199.wgsl.expected.hlsl
View File

@ -1,34 +0,0 @@
Texture2DArray<int4> arg_0 : register(t0, space1);
void textureDimensions_1aa199() {
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_1aa199();
return float4(0.0f, 0.0f, 0.0f, 0.0f);
}
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_1aa199();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
textureDimensions_1aa199();
return;
}

33
test/intrinsics/gen/textureDimensions/1aa199.wgsl.expected.msl
View File

@ -1,33 +0,0 @@
#include <metal_stdlib>
using namespace metal;
struct tint_symbol {
float4 value [[position]];
};
void textureDimensions_1aa199(texture2d_array<int, access::read> tint_symbol_1) {
int2 res = int2(tint_symbol_1.get_width(), tint_symbol_1.get_height());
}
float4 vertex_main_inner(texture2d_array<int, access::read> tint_symbol_2) {
textureDimensions_1aa199(tint_symbol_2);
return float4();
}
vertex tint_symbol vertex_main(texture2d_array<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_array<int, access::read> tint_symbol_4 [[texture(0)]]) {
textureDimensions_1aa199(tint_symbol_4);
return;
}
kernel void compute_main(texture2d_array<int, access::read> tint_symbol_5 [[texture(0)]]) {
textureDimensions_1aa199(tint_symbol_5);
return;
}

78
test/intrinsics/gen/textureDimensions/1aa199.wgsl.expected.spvasm
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_1aa199 "textureDimensions_1aa199"
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 NonWritable
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 Rgba32i
%_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_1aa199 = 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_1aa199
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_1aa199
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %13
%37 = OpLabel
%38 = OpFunctionCall %void %textureDimensions_1aa199
OpReturn
OpFunctionEnd

21
test/intrinsics/gen/textureDimensions/1aa199.wgsl.expected.wgsl
View File

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

46
test/intrinsics/gen/textureDimensions/1e189c.wgsl
View File

@ -1,46 +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/intrinsic-gen
// using the template:
// test/intrinsics/intrinsics.wgsl.tmpl
// and the intrinsic defintion file:
// src/intrinsics.def
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
[[group(1), binding(0)]] var arg_0: texture_storage_3d<r32sint, read>;
// fn textureDimensions(texture: texture_storage_3d<r32sint, read>) -> vec3<i32>
fn textureDimensions_1e189c() {
var res: vec3<i32> = textureDimensions(arg_0);
}
[[stage(vertex)]]
fn vertex_main() -> [[builtin(position)]] vec4<f32> {
textureDimensions_1e189c();
return vec4<f32>();
}
[[stage(fragment)]]
fn fragment_main() {
textureDimensions_1e189c();
}
[[stage(compute), workgroup_size(1)]]
fn compute_main() {
textureDimensions_1e189c();
}

34
test/intrinsics/gen/textureDimensions/1e189c.wgsl.expected.hlsl
View File

@ -1,34 +0,0 @@
Texture3D<int4> arg_0 : register(t0, space1);
void textureDimensions_1e189c() {
int3 tint_tmp;
arg_0.GetDimensions(tint_tmp.x, tint_tmp.y, tint_tmp.z);
int3 res = tint_tmp;
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
textureDimensions_1e189c();
return float4(0.0f, 0.0f, 0.0f, 0.0f);
}
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_1e189c();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
textureDimensions_1e189c();
return;
}

33
test/intrinsics/gen/textureDimensions/1e189c.wgsl.expected.msl
View File

@ -1,33 +0,0 @@
#include <metal_stdlib>
using namespace metal;
struct tint_symbol {
float4 value [[position]];
};
void textureDimensions_1e189c(texture3d<int, access::read> tint_symbol_1) {
int3 res = int3(tint_symbol_1.get_width(), tint_symbol_1.get_height(), tint_symbol_1.get_depth());
}
float4 vertex_main_inner(texture3d<int, access::read> tint_symbol_2) {
textureDimensions_1e189c(tint_symbol_2);
return float4();
}
vertex tint_symbol vertex_main(texture3d<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(texture3d<int, access::read> tint_symbol_4 [[texture(0)]]) {
textureDimensions_1e189c(tint_symbol_4);
return;
}
kernel void compute_main(texture3d<int, access::read> tint_symbol_5 [[texture(0)]]) {
textureDimensions_1e189c(tint_symbol_5);
return;
}

76
test/intrinsics/gen/textureDimensions/1e189c.wgsl.expected.spvasm
View File

@ -1,76 +0,0 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 37
; 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_1e189c "textureDimensions_1e189c"
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 NonWritable
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 3D 0 0 0 2 R32i
%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
%arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant
%void = OpTypeVoid
%13 = OpTypeFunction %void
%v3int = OpTypeVector %int 3
%_ptr_Function_v3int = OpTypePointer Function %v3int
%22 = OpConstantNull %v3int
%23 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%textureDimensions_1e189c = OpFunction %void None %13
%16 = OpLabel
%res = OpVariable %_ptr_Function_v3int Function %22
%19 = OpLoad %11 %arg_0
%17 = OpImageQuerySize %v3int %19
OpStore %res %17
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %23
%25 = OpLabel
%26 = OpFunctionCall %void %textureDimensions_1e189c
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %13
%28 = OpLabel
%29 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %29
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %13
%32 = OpLabel
%33 = OpFunctionCall %void %textureDimensions_1e189c
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %13
%35 = OpLabel
%36 = OpFunctionCall %void %textureDimensions_1e189c
OpReturn
OpFunctionEnd

21
test/intrinsics/gen/textureDimensions/1e189c.wgsl.expected.wgsl
View File

@ -1,21 +0,0 @@
[[group(1), binding(0)]] var arg_0 : texture_storage_3d<r32sint, read>;
fn textureDimensions_1e189c() {
var res : vec3<i32> = textureDimensions(arg_0);
}
[[stage(vertex)]]
fn vertex_main() -> [[builtin(position)]] vec4<f32> {
textureDimensions_1e189c();
return vec4<f32>();
}
[[stage(fragment)]]
fn fragment_main() {
textureDimensions_1e189c();
}
[[stage(compute), workgroup_size(1)]]
fn compute_main() {
textureDimensions_1e189c();
}

46
test/intrinsics/gen/textureDimensions/214b7b.wgsl
View File

@ -1,46 +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/intrinsic-gen
// using the template:
// test/intrinsics/intrinsics.wgsl.tmpl
// and the intrinsic defintion file:
// src/intrinsics.def
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
[[group(1), binding(0)]] var arg_0: texture_storage_1d<rg32float, read>;
// fn textureDimensions(texture: texture_storage_1d<rg32float, read>) -> i32
fn textureDimensions_214b7b() {
var res: i32 = textureDimensions(arg_0);
}
[[stage(vertex)]]
fn vertex_main() -> [[builtin(position)]] vec4<f32> {
textureDimensions_214b7b();
return vec4<f32>();
}
[[stage(fragment)]]
fn fragment_main() {
textureDimensions_214b7b();
}
[[stage(compute), workgroup_size(1)]]
fn compute_main() {
textureDimensions_214b7b();
}

34
test/intrinsics/gen/textureDimensions/214b7b.wgsl.expected.hlsl
View File

@ -1,34 +0,0 @@
Texture1D<float4> arg_0 : register(t0, space1);
void textureDimensions_214b7b() {
int tint_tmp;
arg_0.GetDimensions(tint_tmp);
int res = tint_tmp;
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
textureDimensions_214b7b();
return float4(0.0f, 0.0f, 0.0f, 0.0f);
}
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_214b7b();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
textureDimensions_214b7b();
return;
}

33
test/intrinsics/gen/textureDimensions/214b7b.wgsl.expected.msl
View File

@ -1,33 +0,0 @@
#include <metal_stdlib>
using namespace metal;
struct tint_symbol {
float4 value [[position]];
};
void textureDimensions_214b7b(texture1d<float, access::read> tint_symbol_1) {
int res = int(tint_symbol_1.get_width());
}
float4 vertex_main_inner(texture1d<float, access::read> tint_symbol_2) {
textureDimensions_214b7b(tint_symbol_2);
return float4();
}
vertex tint_symbol vertex_main(texture1d<float, 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(texture1d<float, access::read> tint_symbol_4 [[texture(0)]]) {
textureDimensions_214b7b(tint_symbol_4);
return;
}
kernel void compute_main(texture1d<float, access::read> tint_symbol_5 [[texture(0)]]) {
textureDimensions_214b7b(tint_symbol_5);
return;
}

77
test/intrinsics/gen/textureDimensions/214b7b.wgsl.expected.spvasm
View File

@ -1,77 +0,0 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 36
; Schema: 0
OpCapability Shader
OpCapability Image1D
OpCapability StorageImageExtendedFormats
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_214b7b "textureDimensions_214b7b"
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 NonWritable
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 2 Rg32f
%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
%arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant
%void = OpTypeVoid
%12 = OpTypeFunction %void
%int = OpTypeInt 32 1
%_ptr_Function_int = OpTypePointer Function %int
%21 = OpConstantNull %int
%22 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%textureDimensions_214b7b = OpFunction %void None %12
%15 = OpLabel
%res = OpVariable %_ptr_Function_int Function %21
%18 = OpLoad %11 %arg_0
%16 = OpImageQuerySize %int %18
OpStore %res %16
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %22
%24 = OpLabel
%25 = OpFunctionCall %void %textureDimensions_214b7b
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %12
%27 = OpLabel
%28 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %28
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %12
%31 = OpLabel
%32 = OpFunctionCall %void %textureDimensions_214b7b
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %12
%34 = OpLabel
%35 = OpFunctionCall %void %textureDimensions_214b7b
OpReturn
OpFunctionEnd

21
test/intrinsics/gen/textureDimensions/214b7b.wgsl.expected.wgsl
View File

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

46
test/intrinsics/gen/textureDimensions/2d32ae.wgsl
View File

@ -1,46 +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/intrinsic-gen
// using the template:
// test/intrinsics/intrinsics.wgsl.tmpl
// and the intrinsic defintion file:
// src/intrinsics.def
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
[[group(1), binding(0)]] var arg_0: texture_storage_3d<rgba32uint, read>;
// fn textureDimensions(texture: texture_storage_3d<rgba32uint, read>) -> vec3<i32>
fn textureDimensions_2d32ae() {
var res: vec3<i32> = textureDimensions(arg_0);
}
[[stage(vertex)]]
fn vertex_main() -> [[builtin(position)]] vec4<f32> {
textureDimensions_2d32ae();
return vec4<f32>();
}
[[stage(fragment)]]
fn fragment_main() {
textureDimensions_2d32ae();
}
[[stage(compute), workgroup_size(1)]]
fn compute_main() {
textureDimensions_2d32ae();
}

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