[spirv][ir] Convert spirv image types.
Add initial conversions of the `spirv.image` and `spirv.sampled_image`
types into their texture equivalents. This will be expanded as the
texture methods are implemented.
Bug: 391482238, 391482434
Change-Id: I4cacb9bf413571245a17038b51a0794f91f5c8e2
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/240494
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: James Price <jrprice@google.com>
Auto-Submit: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/spirv/reader/lower/BUILD.bazel b/src/tint/lang/spirv/reader/lower/BUILD.bazel
index f5499f7..7b1610f 100644
--- a/src/tint/lang/spirv/reader/lower/BUILD.bazel
+++ b/src/tint/lang/spirv/reader/lower/BUILD.bazel
@@ -65,6 +65,7 @@
"//src/tint/lang/spirv",
"//src/tint/lang/spirv/intrinsic",
"//src/tint/lang/spirv/ir",
+ "//src/tint/lang/spirv/type",
"//src/tint/utils",
"//src/tint/utils/containers",
"//src/tint/utils/diagnostic",
diff --git a/src/tint/lang/spirv/reader/lower/BUILD.cmake b/src/tint/lang/spirv/reader/lower/BUILD.cmake
index 5c9c9a3..2eeb9bf 100644
--- a/src/tint/lang/spirv/reader/lower/BUILD.cmake
+++ b/src/tint/lang/spirv/reader/lower/BUILD.cmake
@@ -64,6 +64,7 @@
tint_lang_spirv
tint_lang_spirv_intrinsic
tint_lang_spirv_ir
+ tint_lang_spirv_type
tint_utils
tint_utils_containers
tint_utils_diagnostic
diff --git a/src/tint/lang/spirv/reader/lower/BUILD.gn b/src/tint/lang/spirv/reader/lower/BUILD.gn
index 87c9217..31224fb 100644
--- a/src/tint/lang/spirv/reader/lower/BUILD.gn
+++ b/src/tint/lang/spirv/reader/lower/BUILD.gn
@@ -70,6 +70,7 @@
"${tint_src_dir}/lang/spirv",
"${tint_src_dir}/lang/spirv/intrinsic",
"${tint_src_dir}/lang/spirv/ir",
+ "${tint_src_dir}/lang/spirv/type",
"${tint_src_dir}/utils",
"${tint_src_dir}/utils/containers",
"${tint_src_dir}/utils/diagnostic",
diff --git a/src/tint/lang/spirv/reader/lower/texture.cc b/src/tint/lang/spirv/reader/lower/texture.cc
index 98afede..f8ba558 100644
--- a/src/tint/lang/spirv/reader/lower/texture.cc
+++ b/src/tint/lang/spirv/reader/lower/texture.cc
@@ -32,6 +32,9 @@
#include "src/tint/lang/core/ir/builder.h"
#include "src/tint/lang/core/ir/module.h"
#include "src/tint/lang/core/ir/validator.h"
+#include "src/tint/lang/core/type/sampled_texture.h"
+#include "src/tint/lang/core/type/storage_texture.h"
+#include "src/tint/lang/spirv/type/image.h"
namespace tint::spirv::reader::lower {
namespace {
@@ -51,7 +54,74 @@
core::type::Manager& ty{ir.Types()};
/// Process the module.
- void Process() {}
+ void Process() {
+ for (auto* inst : *ir.root_block) {
+ auto* var = inst->As<core::ir::Var>();
+ if (!var) {
+ continue;
+ }
+
+ auto* ptr = var->Result()->Type()->As<core::type::Pointer>();
+ TINT_ASSERT(ptr);
+
+ auto* type = ptr->UnwrapPtr();
+ if (!type->Is<spirv::type::Image>()) {
+ continue;
+ }
+
+ auto* new_ty = TypeFor(type);
+ var->Result()->SetType(ty.ptr(ptr->AddressSpace(), new_ty, ptr->Access()));
+
+ // TOOD(dsinclair): Replace var usages
+ }
+ }
+
+ const core::type::Type* TypeFor(const core::type::Type* src_ty) {
+ TINT_ASSERT(src_ty->Is<spirv::type::Image>());
+
+ if (auto* img = src_ty->As<spirv::type::Image>()) {
+ return TypeForImage(img);
+ }
+
+ TINT_UNREACHABLE();
+ }
+
+ core::type::TextureDimension ConvertDim(spirv::type::Dim dim, spirv::type::Arrayed arrayed) {
+ switch (dim) {
+ case spirv::type::Dim::kD1:
+ return core::type::TextureDimension::k1d;
+ case spirv::type::Dim::kD2:
+ return arrayed == spirv::type::Arrayed::kArrayed
+ ? core::type::TextureDimension::k2dArray
+ : core::type::TextureDimension::k2d;
+ case spirv::type::Dim::kD3:
+ return core::type::TextureDimension::k3d;
+ case spirv::type::Dim::kCube:
+ return arrayed == spirv::type::Arrayed::kArrayed
+ ? core::type::TextureDimension::kCubeArray
+ : core::type::TextureDimension::kCube;
+ default:
+ TINT_UNREACHABLE();
+ }
+ }
+
+ const core::type::Type* TypeForImage(const spirv::type::Image* img) {
+ if (img->GetDim() == spirv::type::Dim::kSubpassData) {
+ return ty.input_attachment(img->GetSampledType());
+ }
+
+ if (img->GetSampled() == spirv::type::Sampled::kReadWriteOpCompatible) {
+ return ty.storage_texture(ConvertDim(img->GetDim(), img->GetArrayed()),
+ img->GetTexelFormat(), img->GetAccess());
+ }
+
+ // TODO(dsinclair): Handle determining depth texture by usage
+ if (img->GetDepth() == spirv::type::Depth::kDepth) {
+ return ty.depth_texture(ConvertDim(img->GetDim(), img->GetArrayed()));
+ }
+ return ty.sampled_texture(ConvertDim(img->GetDim(), img->GetArrayed()),
+ img->GetSampledType());
+ }
};
} // namespace
diff --git a/src/tint/lang/spirv/reader/lower/texture_test.cc b/src/tint/lang/spirv/reader/lower/texture_test.cc
index 33081ea..0843602 100644
--- a/src/tint/lang/spirv/reader/lower/texture_test.cc
+++ b/src/tint/lang/spirv/reader/lower/texture_test.cc
@@ -46,7 +46,7 @@
using Multisampled = spirv::type::Multisampled;
using Sampled = spirv::type::Sampled;
-TEST_F(SpirvReader_TextureTest, DISABLED_Type_Image_1d) {
+TEST_F(SpirvReader_TextureTest, Type_Image_1d) {
b.Append(mod.root_block, [&] {
auto* v = b.Var("wg", ty.ptr(handle,
ty.Get<spirv::type::Image>(
@@ -75,7 +75,7 @@
ASSERT_EQ(expect, str());
}
-TEST_F(SpirvReader_TextureTest, DISABLED_Type_Image_2d) {
+TEST_F(SpirvReader_TextureTest, Type_Image_2d) {
b.Append(mod.root_block, [&] {
auto* v = b.Var("wg", ty.ptr(handle,
ty.Get<spirv::type::Image>(
@@ -104,7 +104,7 @@
ASSERT_EQ(expect, str());
}
-TEST_F(SpirvReader_TextureTest, DISABLED_Type_Image_3d) {
+TEST_F(SpirvReader_TextureTest, Type_Image_3d) {
b.Append(mod.root_block, [&] {
auto* v = b.Var("wg", ty.ptr(handle,
ty.Get<spirv::type::Image>(
@@ -133,7 +133,7 @@
ASSERT_EQ(expect, str());
}
-TEST_F(SpirvReader_TextureTest, DISABLED_Type_Image_Cube) {
+TEST_F(SpirvReader_TextureTest, Type_Image_Cube) {
b.Append(mod.root_block, [&] {
auto* v =
b.Var("wg", ty.ptr(handle,
@@ -163,7 +163,7 @@
ASSERT_EQ(expect, str());
}
-TEST_F(SpirvReader_TextureTest, DISABLED_Type_Image_SubpassData) {
+TEST_F(SpirvReader_TextureTest, Type_Image_SubpassData) {
b.Append(mod.root_block, [&] {
auto* v = b.Var(
"wg", ty.ptr(handle,
@@ -193,11 +193,11 @@
ASSERT_EQ(expect, str());
}
-TEST_F(SpirvReader_TextureTest, DISABLED_Type_Image_Depth) {
+TEST_F(SpirvReader_TextureTest, Type_Image_Depth) {
b.Append(mod.root_block, [&] {
auto* v = b.Var("wg", ty.ptr(handle,
ty.Get<spirv::type::Image>(
- ty.f32(), Dim::kD1, Depth::kDepth, Arrayed::kNonArrayed,
+ ty.f32(), Dim::kD2, Depth::kDepth, Arrayed::kNonArrayed,
Multisampled::kSingleSampled, Sampled::kSamplingCompatible,
core::TexelFormat::kUndefined, core::Access::kRead),
read));
@@ -206,7 +206,7 @@
auto* src = R"(
$B1: { # root
- %wg:ptr<handle, spirv.image<f32, 1d, depth, non_arrayed, single_sampled, sampling_compatible, undefined, read>, read> = var undef @binding_point(1, 2)
+ %wg:ptr<handle, spirv.image<f32, 2d, depth, non_arrayed, single_sampled, sampling_compatible, undefined, read>, read> = var undef @binding_point(1, 2)
}
)";
@@ -215,7 +215,7 @@
auto* expect = R"(
$B1: { # root
- %wg:ptr<handle, texture_depth_1d<f32>, read> = var undef @binding_point(1, 2)
+ %wg:ptr<handle, texture_depth_2d, read> = var undef @binding_point(1, 2)
}
)";
@@ -251,7 +251,7 @@
ASSERT_EQ(expect, str());
}
-TEST_F(SpirvReader_TextureTest, DISABLED_Type_Image_Arrayed) {
+TEST_F(SpirvReader_TextureTest, Type_Image_Arrayed) {
b.Append(mod.root_block, [&] {
auto* v = b.Var("wg", ty.ptr(handle,
ty.Get<spirv::type::Image>(
@@ -280,21 +280,21 @@
ASSERT_EQ(expect, str());
}
-TEST_F(SpirvReader_TextureTest, DISABLED_Type_Image_RW_Ops) {
+TEST_F(SpirvReader_TextureTest, Type_Image_RW_Ops) {
b.Append(mod.root_block, [&] {
auto* v =
b.Var("wg", ty.ptr(handle,
ty.Get<spirv::type::Image>(
ty.f32(), Dim::kD1, Depth::kNotDepth, Arrayed::kNonArrayed,
Multisampled::kSingleSampled, Sampled::kReadWriteOpCompatible,
- core::TexelFormat::kUndefined, core::Access::kRead),
+ core::TexelFormat::kRg32Float, core::Access::kRead),
read));
v->SetBindingPoint(1, 2);
});
auto* src = R"(
$B1: { # root
- %wg:ptr<handle, spirv.image<f32, 1d, not_depth, non_arrayed, single_sampled, rw_op_compatbile, undefined, read>, read> = var undef @binding_point(1, 2)
+ %wg:ptr<handle, spirv.image<f32, 1d, not_depth, non_arrayed, single_sampled, rw_op_compatible, rg32float, read>, read> = var undef @binding_point(1, 2)
}
)";
@@ -303,14 +303,14 @@
auto* expect = R"(
$B1: { # root
- %wg:ptr<handle, texture_1d<f32>, read> = var undef @binding_point(1, 2)
+ %wg:ptr<handle, texture_storage_1d<rg32float, read>, read> = var undef @binding_point(1, 2)
}
)";
ASSERT_EQ(expect, str());
}
-TEST_F(SpirvReader_TextureTest, DISABLED_Type_Image_TexelFormat) {
+TEST_F(SpirvReader_TextureTest, Type_Image_TexelFormat) {
b.Append(mod.root_block, [&] {
auto* v = b.Var("wg", ty.ptr(handle,
ty.Get<spirv::type::Image>(
@@ -323,7 +323,7 @@
auto* src = R"(
$B1: { # root
- %wg:ptr<handle, spirv.image<f32, 2d, not_depth, non_arrayed, single_sampled, sampling_compatible, undefined, read>, read> = var undef @binding_point(1, 2)
+ %wg:ptr<handle, spirv.image<f32, 1d, not_depth, non_arrayed, single_sampled, sampling_compatible, rg32float, read>, read> = var undef @binding_point(1, 2)
}
)";
@@ -332,37 +332,7 @@
auto* expect = R"(
$B1: { # root
- %wg:ptr<handle, texture_storage_1d<rg32float, read>, read_write> = var undef @binding_point(1, 2)
-}
-
-)";
- ASSERT_EQ(expect, str());
-}
-
-TEST_F(SpirvReader_TextureTest, DISABLED_Type_SampledImage) {
- b.Append(mod.root_block, [&] {
- auto* v = b.Var("wg", ty.ptr(handle,
- ty.Get<spirv::type::SampledImage>(ty.Get<spirv::type::Image>(
- ty.f32(), Dim::kD1, Depth::kNotDepth, Arrayed::kNonArrayed,
- Multisampled::kSingleSampled, Sampled::kSamplingCompatible,
- core::TexelFormat::kRg32Float, core::Access::kRead)),
- read));
- v->SetBindingPoint(1, 2);
- });
-
- auto* src = R"(
-$B1: { # root
- %wg:ptr<handle, spirv.sampled_image<spirv.image<f32, 2d, not_depth, non_arrayed, single_sampled, sampling_compatible, undefined, read>>, read> = var undef @binding_point(1, 2)
-}
-
-)";
- ASSERT_EQ(src, str());
- Run(Texture);
-
- auto* expect = R"(
-$B1: { # root
- %wg_1:ptr<handle, sampler, read> = var undef @binding_point(1, 2)
- %wg:ptr<handle, texture_storage_1d<rg32float, read>, read> = var undef @binding_point(1, 2)
+ %wg:ptr<handle, texture_1d<f32>, read> = var undef @binding_point(1, 2)
}
)";
diff --git a/src/tint/lang/spirv/reader/parser/image_sampler_test.cc b/src/tint/lang/spirv/reader/parser/image_sampler_test.cc
index 905007d..0d3aaae 100644
--- a/src/tint/lang/spirv/reader/parser/image_sampler_test.cc
+++ b/src/tint/lang/spirv/reader/parser/image_sampler_test.cc
@@ -416,7 +416,8 @@
)",
R"(
$B1: { # root
- %1:ptr<handle, spirv.sampled_image<spirv.image<f32, 1d, not_depth, non_arrayed, single_sampled, sampling_compatible, rg32float, read>>, read> = var undef @binding_point(1, 2)
+ %1:ptr<handle, sampler, read> = var undef @binding_point(1, 2)
+ %2:ptr<handle, spirv.image<f32, 1d, not_depth, non_arrayed, single_sampled, sampling_compatible, rg32float, read>, read> = var undef @binding_point(1, 2)
}
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index 8b3cde0..31e7f70 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -43,6 +43,7 @@
TINT_BEGIN_DISABLE_WARNING(WEAK_VTABLES);
TINT_BEGIN_DISABLE_WARNING(UNSAFE_BUFFER_USAGE);
#include "source/opt/build_module.h"
+#include "source/opt/split_combined_image_sampler_pass.h"
TINT_END_DISABLE_WARNING(UNSAFE_BUFFER_USAGE);
TINT_END_DISABLE_WARNING(WEAK_VTABLES);
TINT_END_DISABLE_WARNING(SIGN_CONVERSION);
@@ -89,6 +90,14 @@
return Failure("failed to build the internal representation of the module");
}
+ {
+ spvtools::opt::SplitCombinedImageSamplerPass pass;
+ auto status = pass.Run(spirv_context_.get());
+ if (status == spvtools::opt::Pass::Status::Failure) {
+ return Failure("failed to run SplitCombinedImageSamplerPass in SPIR-V opt");
+ }
+ }
+
// Check for unsupported extensions.
for (const auto& ext : spirv_context_->extensions()) {
auto name = ext.GetOperand(0).AsString();
@@ -419,8 +428,8 @@
texel_format, access);
}
case spvtools::opt::analysis::Type::kSampledImage: {
- auto* sampled = type->AsSampledImage();
- return ty_.Get<spirv::type::SampledImage>(Type(sampled->image_type()));
+ TINT_UNREACHABLE() << "OpTypeSampledImage should have been removed by the "
+ "SplitCombinedImageSamplerPass";
}
default: {
TINT_UNIMPLEMENTED() << "unhandled SPIR-V type: " << type->str();
