aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--docs/amber_script.md36
-rw-r--r--src/CMakeLists.txt1
-rw-r--r--src/amberscript/parser.cc174
-rw-r--r--src/amberscript/parser.h1
-rw-r--r--src/amberscript/parser_image_test.cc272
-rw-r--r--src/amberscript/parser_sampler_test.cc31
-rw-r--r--src/buffer.h16
-rw-r--r--src/image.h24
-rw-r--r--src/sampler.h8
-rw-r--r--src/vulkan/frame_buffer.cc7
-rw-r--r--src/vulkan/image_descriptor.cc41
-rw-r--r--src/vulkan/sampler.cc4
-rw-r--r--src/vulkan/transfer_image.cc26
-rw-r--r--src/vulkan/transfer_image.h3
-rw-r--r--tests/cases/glsl_read_and_write_image3d_rgba32i.amber76
-rw-r--r--tests/cases/opencl_read_and_write_image3d_rgba32i.amber74
-rw-r--r--tests/cases/opencl_read_image.amber2
-rw-r--r--tests/cases/opencl_write_image.amber2
-rwxr-xr-xtests/run_tests.py3
19 files changed, 768 insertions, 33 deletions
diff --git a/docs/amber_script.md b/docs/amber_script.md
index 3a403a4..03a23b0 100644
--- a/docs/amber_script.md
+++ b/docs/amber_script.md
@@ -137,10 +137,11 @@ END
BUFFER {name} DATA_TYPE {type} {STD140 | STD430} SIZE _size_in_items_ \
{initializer}
+# Deprecated
# Defines a buffer with width and height and filled by data as specified by the
# `initializer`.
BUFFER {name} DATA_TYPE {type} {STD140 | STD430} WIDTH {w} HEIGHT {h} \
- {initializer}
+ {initializer}
# Creates a buffer which will store the given `FORMAT` of data. These
# buffers are used as image and depth buffers in the `PIPELINE` commands.
@@ -149,6 +150,31 @@ BUFFER {name} FORMAT {format_string} \
[ MIP_LEVELS _mip_levels_ (default 1) ]
```
+#### Images
+
+An AmberScript image is a specialized buffer that specifies image-specific
+attributes.
+
+##### Dimensionality
+ * `DIM_1D` -- A 1-dimensional image
+ * `DIM_2D` -- A 2-dimensional image
+ * `DIM_3D` -- A 3-dimensional image
+
+```groovy
+# Specify an image buffer with a format. HEIGHT is necessary for DIM_2D and
+# DIM_3D. DEPTH is necessary for DIM_3D.
+IMAGE {name} FORMAT {format_string} [ MIP_LEVELS _mip_levels_ (default 1) ] \
+ {dimensionality} \
+ WIDTH {w} [ HEIGHT {h} [ DEPTH {d} ] ] \
+ {initializer}
+
+# Specify an image buffer with a data type. HEIGHT is necessary for DIM_2D and
+# DIM_3D. DEPTH is necessary for DIM_3D.
+IMAGE {name} DATA_TYPE {type} {dimensionality} \
+ WIDTH {w} [ HEIGHT {h} [ DEPTH {d} ] ] \
+ {intializer}
+```
+
#### Buffer Initializers
```groovy
@@ -177,8 +203,6 @@ COPY {buffer_from} TO {buffer_to}
Samplers are used for sampling buffers that are bound to a pipeline as
sampled image or combined image sampler.
-The samplers use normalized coordinates in the range of [0..1].
-
#### Filter types
* `nearest`
* `linear`
@@ -206,11 +230,15 @@ SAMPLER {name} \
[ MIN_FILTER {filter_type} (default nearest) ] \
[ ADDRESS_MODE_U {address_mode} (default repeat) ] \
[ ADDRESS_MODE_V {address_mode} (default repeat) ] \
+ [ ADDRESS_MODE_W {address_mode} (default repeat) ] \
[ BORDER_COLOR {border_color} (default float_transparent_black) ] \
[ MIN_LOD _val_ (default 0.0) ] \
- [ MAX_LOD _val_ (default 1.0) ]
+ [ MAX_LOD _val_ (default 1.0) ] \
+ [ NORMALIZED_COORDS | UNNORMALIZED_COORDS (default NORMALIZED_COORDS) ]
```
+Note: unnormalized coordinates will override MIN\_LOD and MAX\_LOD to 0.0.
+
### Pipelines
#### Pipeline type
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index cd05481..ae937f6 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -125,6 +125,7 @@ if (${AMBER_ENABLE_TESTS})
amberscript/parser_expect_test.cc
amberscript/parser_extension_test.cc
amberscript/parser_framebuffer_test.cc
+ amberscript/parser_image_test.cc
amberscript/parser_pipeline_test.cc
amberscript/parser_pipeline_set_test.cc
amberscript/parser_repeat_test.cc
diff --git a/src/amberscript/parser.cc b/src/amberscript/parser.cc
index 12edfbc..9a207d4 100644
--- a/src/amberscript/parser.cc
+++ b/src/amberscript/parser.cc
@@ -21,6 +21,7 @@
#include <utility>
#include <vector>
+#include "src/image.h"
#include "src/make_unique.h"
#include "src/sampler.h"
#include "src/shader_data.h"
@@ -140,6 +141,17 @@ AddressMode StrToAddressMode(std::string str) {
return AddressMode::kUnknown;
}
+ImageDimension StrToImageDimension(const std::string& str) {
+ if (str == "DIM_1D")
+ return ImageDimension::k1D;
+ if (str == "DIM_2D")
+ return ImageDimension::k2D;
+ if (str == "DIM_3D")
+ return ImageDimension::k3D;
+
+ return ImageDimension::kUnknown;
+}
+
} // namespace
Parser::Parser() : amber::Parser() {}
@@ -172,6 +184,8 @@ Result Parser::Parse(const std::string& data) {
r = ParseDeviceFeature();
} else if (tok == "DEVICE_EXTENSION") {
r = ParseDeviceExtension();
+ } else if (tok == "IMAGE") {
+ r = ParseImage();
} else if (tok == "INSTANCE_EXTENSION") {
r = ParseInstanceExtension();
} else if (tok == "PIPELINE") {
@@ -1205,6 +1219,146 @@ Result Parser::ParseBuffer() {
return {};
}
+Result Parser::ParseImage() {
+ auto token = tokenizer_->NextToken();
+ if (!token->IsString())
+ return Result("invalid IMAGE name provided");
+
+ auto name = token->AsString();
+ if (name == "DATA_TYPE" || name == "FORMAT")
+ return Result("missing IMAGE name");
+
+ token = tokenizer_->NextToken();
+ if (!token->IsString())
+ return Result("invalid IMAGE command provided");
+
+ std::unique_ptr<Buffer> buffer = MakeUnique<Buffer>();
+ buffer->SetName(name);
+ auto& cmd = token->AsString();
+ if (cmd == "DATA_TYPE") {
+ token = tokenizer_->NextToken();
+ if (!token->IsString())
+ return Result("IMAGE invalid data type");
+
+ auto type = script_->ParseType(token->AsString());
+ std::unique_ptr<Format> fmt;
+ if (type != nullptr) {
+ fmt = MakeUnique<Format>(type);
+ buffer->SetFormat(fmt.get());
+ } else {
+ auto new_type = ToType(token->AsString());
+ if (!new_type)
+ return Result("invalid data type '" + token->AsString() + "' provided");
+
+ fmt = MakeUnique<Format>(new_type.get());
+ buffer->SetFormat(fmt.get());
+ script_->RegisterType(std::move(new_type));
+ }
+ script_->RegisterFormat(std::move(fmt));
+ } else if (cmd == "FORMAT") {
+ token = tokenizer_->NextToken();
+ if (!token->IsString())
+ return Result("IMAGE FORMAT must be a string");
+
+ auto type = script_->ParseType(token->AsString());
+ if (!type)
+ return Result("invalid IMAGE FORMAT");
+
+ auto fmt = MakeUnique<Format>(type);
+ buffer->SetFormat(fmt.get());
+ script_->RegisterFormat(std::move(fmt));
+
+ token = tokenizer_->PeekNextToken();
+ if (token->IsString() && token->AsString() == "MIP_LEVELS") {
+ tokenizer_->NextToken();
+ token = tokenizer_->NextToken();
+
+ if (!token->IsInteger())
+ return Result("invalid value for MIP_LEVELS");
+
+ buffer->SetMipLevels(token->AsUint32());
+ }
+ } else {
+ return Result("unknown IMAGE command provided: " + cmd);
+ }
+
+ token = tokenizer_->NextToken();
+ if (!token->IsString()) {
+ return Result("IMAGE dimensionality must be a string: " +
+ token->ToOriginalString());
+ }
+
+ auto dim = StrToImageDimension(token->AsString());
+ if (dim == ImageDimension::kUnknown)
+ return Result("unknown IMAGE dimensionality");
+ buffer->SetImageDimension(dim);
+
+ token = tokenizer_->NextToken();
+ if (!token->IsString() || token->AsString() != "WIDTH")
+ return Result("expected IMAGE WIDTH");
+
+ // Parse image dimensions.
+ uint32_t width = 1;
+ uint32_t height = 1;
+ uint32_t depth = 1;
+ token = tokenizer_->NextToken();
+ if (!token->IsInteger() || token->AsUint32() == 0)
+ return Result("expected positive IMAGE WIDTH");
+ width = token->AsUint32();
+ buffer->SetWidth(width);
+
+ if (dim == ImageDimension::k2D || dim == ImageDimension::k3D) {
+ token = tokenizer_->NextToken();
+ if (!token->IsString() || token->AsString() != "HEIGHT")
+ return Result("expected IMAGE HEIGHT");
+
+ token = tokenizer_->NextToken();
+ if (!token->IsInteger() || token->AsUint32() == 0)
+ return Result("expected positive IMAGE HEIGHT");
+ height = token->AsUint32();
+ buffer->SetHeight(height);
+ }
+
+ if (dim == ImageDimension::k3D) {
+ token = tokenizer_->NextToken();
+ if (!token->IsString() || token->AsString() != "DEPTH")
+ return Result("expected IMAGE DEPTH");
+
+ token = tokenizer_->NextToken();
+ if (!token->IsInteger() || token->AsUint32() == 0)
+ return Result("expected positive IMAGE DEPTH");
+ depth = token->AsUint32();
+ buffer->SetDepth(depth);
+ }
+
+ const uint32_t size_in_items = width * height * depth;
+ buffer->SetElementCount(size_in_items);
+
+ // Parse initializers.
+ token = tokenizer_->NextToken();
+ if (token->IsString()) {
+ if (token->AsString() == "FILL") {
+ Result r = ParseBufferInitializerFill(buffer.get(), size_in_items);
+ if (!r.IsSuccess())
+ return r;
+ } else if (token->AsString() == "SERIES_FROM") {
+ Result r = ParseBufferInitializerSeries(buffer.get(), size_in_items);
+ if (!r.IsSuccess())
+ return r;
+ } else {
+ return Result("unexpected IMAGE token: " + token->AsString());
+ }
+ } else if (!token->IsEOL() && !token->IsEOS()) {
+ return Result("unexpected IMAGE token: " + token->ToOriginalString());
+ }
+
+ Result r = script_->AddBuffer(std::move(buffer));
+ if (!r.IsSuccess())
+ return r;
+
+ return {};
+}
+
Result Parser::ParseBufferInitializer(Buffer* buffer) {
auto token = tokenizer_->NextToken();
if (!token->IsString())
@@ -1252,6 +1406,7 @@ Result Parser::ParseBufferInitializer(Buffer* buffer) {
if (width == 0)
return Result("expected WIDTH to be positive");
buffer->SetWidth(width);
+ buffer->SetImageDimension(ImageDimension::k2D);
token = tokenizer_->NextToken();
if (token->AsString() != "HEIGHT")
@@ -2321,6 +2476,19 @@ Result Parser::ParseSampler() {
return Result("invalid ADDRESS_MODE_V value " + mode_str);
sampler->SetAddressModeV(mode);
+ } else if (param == "ADDRESS_MODE_W") {
+ token = tokenizer_->NextToken();
+
+ if (!token->IsString())
+ return Result("invalid token when looking for ADDRESS_MODE_W value");
+
+ auto mode_str = token->AsString();
+ auto mode = StrToAddressMode(mode_str);
+
+ if (mode == AddressMode::kUnknown)
+ return Result("invalid ADDRESS_MODE_W value " + mode_str);
+
+ sampler->SetAddressModeW(mode);
} else if (param == "BORDER_COLOR") {
token = tokenizer_->NextToken();
@@ -2357,6 +2525,12 @@ Result Parser::ParseSampler() {
return Result("invalid token when looking for MAX_LOD value");
sampler->SetMaxLOD(token->AsFloat());
+ } else if (param == "NORMALIZED_COORDS") {
+ sampler->SetNormalizedCoords(true);
+ } else if (param == "UNNORMALIZED_COORDS") {
+ sampler->SetNormalizedCoords(false);
+ sampler->SetMinLOD(0.0f);
+ sampler->SetMaxLOD(0.0f);
} else {
return Result("unexpected sampler parameter " + param);
}
diff --git a/src/amberscript/parser.h b/src/amberscript/parser.h
index 18f6ad9..ef0eab3 100644
--- a/src/amberscript/parser.h
+++ b/src/amberscript/parser.h
@@ -50,6 +50,7 @@ class Parser : public amber::Parser {
Result ParseStruct();
Result ParseBuffer();
+ Result ParseImage();
Result ParseBufferInitializer(Buffer*);
Result ParseBufferInitializerSize(Buffer*);
Result ParseBufferInitializerFill(Buffer*, uint32_t);
diff --git a/src/amberscript/parser_image_test.cc b/src/amberscript/parser_image_test.cc
new file mode 100644
index 0000000..4acba8b
--- /dev/null
+++ b/src/amberscript/parser_image_test.cc
@@ -0,0 +1,272 @@
+// Copyright 2019 The Amber 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 parseried.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "gtest/gtest.h"
+#include "src/amberscript/parser.h"
+
+namespace amber {
+namespace amberscript {
+
+using AmberScriptParserTest = testing::Test;
+
+TEST_F(AmberScriptParserTest, ImageNameMissing1) {
+ std::string in = R"(
+IMAGE
+)";
+
+ Parser parser;
+ Result r = parser.Parse(in);
+ ASSERT_FALSE(r.IsSuccess());
+ EXPECT_EQ("3: invalid IMAGE name provided", r.Error());
+}
+
+TEST_F(AmberScriptParserTest, ImageNameMissing2) {
+ std::string in = R"(
+IMAGE DATA_TYPE
+)";
+
+ Parser parser;
+ Result r = parser.Parse(in);
+ ASSERT_FALSE(r.IsSuccess());
+ EXPECT_EQ("2: missing IMAGE name", r.Error());
+}
+
+TEST_F(AmberScriptParserTest, ImageNameMissing3) {
+ std::string in = R"(
+IMAGE FORMAT
+)";
+
+ Parser parser;
+ Result r = parser.Parse(in);
+ ASSERT_FALSE(r.IsSuccess());
+ EXPECT_EQ("2: missing IMAGE name", r.Error());
+}
+
+TEST_F(AmberScriptParserTest, ImageNameInvalid) {
+ std::string in = R"(
+IMAGE 1
+)";
+
+ Parser parser;
+ Result r = parser.Parse(in);
+ ASSERT_FALSE(r.IsSuccess());
+ EXPECT_EQ("2: invalid IMAGE name provided", r.Error());
+}
+
+TEST_F(AmberScriptParserTest, ImageDataTypeInvalid) {
+ std::string in = R"(
+IMAGE image DATA_TYPE blah
+)";
+
+ Parser parser;
+ Result r = parser.Parse(in);
+ ASSERT_FALSE(r.IsSuccess());
+ EXPECT_EQ("2: invalid data type 'blah' provided", r.Error());
+}
+
+TEST_F(AmberScriptParserTest, ImageFormatInvalid) {
+ std::string in = R"(
+IMAGE image FORMAT blah
+)";
+
+ Parser parser;
+ Result r = parser.Parse(in);
+ ASSERT_FALSE(r.IsSuccess());
+ EXPECT_EQ("2: invalid IMAGE FORMAT", r.Error());
+}
+
+TEST_F(AmberScriptParserTest, ImageMipLevelsInvalid) {
+ std::string in = R"(
+IMAGE image FORMAT R32G32B32A32_SFLOAT MIP_LEVELS mips
+)";
+
+ Parser parser;
+ Result r = parser.Parse(in);
+ ASSERT_FALSE(r.IsSuccess());
+ EXPECT_EQ("2: invalid value for MIP_LEVELS", r.Error());
+}
+
+TEST_F(AmberScriptParserTest, ImageMissingDataTypeCommand) {
+ std::string in = R"(
+IMAGE image OTHER
+)";
+
+ Parser parser;
+ Result r = parser.Parse(in);
+ ASSERT_FALSE(r.IsSuccess());
+ EXPECT_EQ("2: unknown IMAGE command provided: OTHER", r.Error());
+}
+
+TEST_F(AmberScriptParserTest, ImageDimensionalityInvalid) {
+ std::string in = R"(
+IMAGE image DATA_TYPE uint32 DIM_WRONG
+)";
+
+ Parser parser;
+ Result r = parser.Parse(in);
+ ASSERT_FALSE(r.IsSuccess());
+ EXPECT_EQ("2: unknown IMAGE dimensionality", r.Error());
+}
+
+TEST_F(AmberScriptParserTest, ImageDimensionalityInvalid2) {
+ std::string in = R"(
+IMAGE image DATA_TYPE uint32 4
+)";
+
+ Parser parser;
+ Result r = parser.Parse(in);
+ ASSERT_FALSE(r.IsSuccess());
+ EXPECT_EQ("2: IMAGE dimensionality must be a string: 4", r.Error());
+}
+
+TEST_F(AmberScriptParserTest, ImageWidthMissing) {
+ std::string in = R"(
+IMAGE image DATA_TYPE uint32 DIM_3D HEIGHT 2 DEPTH 2 FILL 0
+)";
+
+ Parser parser;
+ Result r = parser.Parse(in);
+ ASSERT_FALSE(r.IsSuccess());
+ EXPECT_EQ("2: expected IMAGE WIDTH", r.Error());
+}
+
+TEST_F(AmberScriptParserTest, ImageHeightMissing) {
+ std::string in = R"(
+IMAGE image DATA_TYPE uint32 DIM_3D WIDTH 2 DEPTH 2 FILL 0
+)";
+
+ Parser parser;
+ Result r = parser.Parse(in);
+ ASSERT_FALSE(r.IsSuccess());
+ EXPECT_EQ("2: expected IMAGE HEIGHT", r.Error());
+}
+
+TEST_F(AmberScriptParserTest, ImageDepthMissing) {
+ std::string in = R"(
+IMAGE image DATA_TYPE uint32 DIM_3D WIDTH 2 HEIGHT 2 FILL 0
+)";
+
+ Parser parser;
+ Result r = parser.Parse(in);
+ ASSERT_FALSE(r.IsSuccess());
+ EXPECT_EQ("2: expected IMAGE DEPTH", r.Error());
+}
+
+TEST_F(AmberScriptParserTest, ImageWidthMissingNumber) {
+ std::string in = R"(
+IMAGE image DATA_TYPE uint32 DIM_3D WIDTH HEIGHT 2 DEPTH 2 FILL 0
+)";
+
+ Parser parser;
+ Result r = parser.Parse(in);
+ ASSERT_FALSE(r.IsSuccess());
+ EXPECT_EQ("2: expected positive IMAGE WIDTH", r.Error());
+}
+
+TEST_F(AmberScriptParserTest, ImageHeightMissingNumber) {
+ std::string in = R"(
+IMAGE image DATA_TYPE uint32 DIM_3D WIDTH 2 HEIGHT DEPTH 2 FILL 0
+)";
+
+ Parser parser;
+ Result r = parser.Parse(in);
+ ASSERT_FALSE(r.IsSuccess());
+ EXPECT_EQ("2: expected positive IMAGE HEIGHT", r.Error());
+}
+
+TEST_F(AmberScriptParserTest, ImageDepthMissingNumber) {
+ std::string in = R"(
+IMAGE image DATA_TYPE uint32 DIM_3D WIDTH 2 HEIGHT 2 DEPTH FILL 0
+)";
+
+ Parser parser;
+ Result r = parser.Parse(in);
+ ASSERT_FALSE(r.IsSuccess());
+ EXPECT_EQ("2: expected positive IMAGE DEPTH", r.Error());
+}
+
+TEST_F(AmberScriptParserTest, Image1D) {
+ std::string in = R"(
+IMAGE image DATA_TYPE uint32 DIM_1D WIDTH 4
+)";
+
+ Parser parser;
+ Result r = parser.Parse(in);
+ ASSERT_TRUE(r.IsSuccess());
+ auto script = parser.GetScript();
+ const auto& buffers = script->GetBuffers();
+ ASSERT_EQ(1U, buffers.size());
+
+ ASSERT_TRUE(buffers[0] != nullptr);
+ EXPECT_EQ("image", buffers[0]->GetName());
+
+ auto* buffer = buffers[0].get();
+ EXPECT_TRUE(buffer->GetFormat()->IsUint32());
+ EXPECT_EQ(ImageDimension::k1D, buffer->GetImageDimension());
+ EXPECT_EQ(4, buffer->GetWidth());
+ EXPECT_EQ(1, buffer->GetHeight());
+ EXPECT_EQ(1, buffer->GetDepth());
+ EXPECT_EQ(4, buffer->ElementCount());
+}
+
+TEST_F(AmberScriptParserTest, Image2D) {
+ std::string in = R"(
+IMAGE image DATA_TYPE uint32 DIM_2D WIDTH 3 HEIGHT 4
+)";
+
+ Parser parser;
+ Result r = parser.Parse(in);
+ ASSERT_TRUE(r.IsSuccess());
+ auto script = parser.GetScript();
+ const auto& buffers = script->GetBuffers();
+ ASSERT_EQ(1U, buffers.size());
+
+ ASSERT_TRUE(buffers[0] != nullptr);
+ EXPECT_EQ("image", buffers[0]->GetName());
+
+ auto* buffer = buffers[0].get();
+ EXPECT_TRUE(buffer->GetFormat()->IsUint32());
+ EXPECT_EQ(ImageDimension::k2D, buffer->GetImageDimension());
+ EXPECT_EQ(3, buffer->GetWidth());
+ EXPECT_EQ(4, buffer->GetHeight());
+ EXPECT_EQ(1, buffer->GetDepth());
+ EXPECT_EQ(12, buffer->ElementCount());
+}
+
+TEST_F(AmberScriptParserTest, Image3D) {
+ std::string in = R"(
+IMAGE image DATA_TYPE uint32 DIM_3D WIDTH 3 HEIGHT 4 DEPTH 5
+)";
+
+ Parser parser;
+ Result r = parser.Parse(in);
+ ASSERT_TRUE(r.IsSuccess());
+ auto script = parser.GetScript();
+ const auto& buffers = script->GetBuffers();
+ ASSERT_EQ(1U, buffers.size());
+
+ ASSERT_TRUE(buffers[0] != nullptr);
+ EXPECT_EQ("image", buffers[0]->GetName());
+
+ auto* buffer = buffers[0].get();
+ EXPECT_TRUE(buffer->GetFormat()->IsUint32());
+ EXPECT_EQ(ImageDimension::k3D, buffer->GetImageDimension());
+ EXPECT_EQ(3, buffer->GetWidth());
+ EXPECT_EQ(4, buffer->GetHeight());
+ EXPECT_EQ(5, buffer->GetDepth());
+ EXPECT_EQ(60, buffer->ElementCount());
+}
+
+} // namespace amberscript
+} // namespace amber
diff --git a/src/amberscript/parser_sampler_test.cc b/src/amberscript/parser_sampler_test.cc
index 65dc05a..6cba69d 100644
--- a/src/amberscript/parser_sampler_test.cc
+++ b/src/amberscript/parser_sampler_test.cc
@@ -40,9 +40,11 @@ TEST_F(AmberScriptParserTest, SamplerDefaultValues) {
EXPECT_EQ(FilterType::kNearest, sampler->GetMipmapMode());
EXPECT_EQ(AddressMode::kRepeat, sampler->GetAddressModeU());
EXPECT_EQ(AddressMode::kRepeat, sampler->GetAddressModeV());
+ EXPECT_EQ(AddressMode::kRepeat, sampler->GetAddressModeW());
EXPECT_EQ(BorderColor::kFloatTransparentBlack, sampler->GetBorderColor());
EXPECT_EQ(0.0, sampler->GetMinLOD());
EXPECT_EQ(1.0, sampler->GetMaxLOD());
+ EXPECT_EQ(true, sampler->GetNormalizedCoords());
}
TEST_F(AmberScriptParserTest, SamplerCustomValues) {
@@ -51,9 +53,11 @@ SAMPLER sampler MAG_FILTER linear \
MIN_FILTER linear \
ADDRESS_MODE_U clamp_to_edge \
ADDRESS_MODE_V clamp_to_border \
+ ADDRESS_MODE_W mirrored_repeat \
BORDER_COLOR float_opaque_white \
MIN_LOD 2.5 \
- MAX_LOD 5.0)";
+ MAX_LOD 5.0 \
+ NORMALIZED_COORDS)";
Parser parser;
Result r = parser.Parse(in);
@@ -72,9 +76,11 @@ SAMPLER sampler MAG_FILTER linear \
EXPECT_EQ(FilterType::kNearest, sampler->GetMipmapMode());
EXPECT_EQ(AddressMode::kClampToEdge, sampler->GetAddressModeU());
EXPECT_EQ(AddressMode::kClampToBorder, sampler->GetAddressModeV());
+ EXPECT_EQ(AddressMode::kMirroredRepeat, sampler->GetAddressModeW());
EXPECT_EQ(BorderColor::kFloatOpaqueWhite, sampler->GetBorderColor());
EXPECT_EQ(2.5, sampler->GetMinLOD());
EXPECT_EQ(5.0, sampler->GetMaxLOD());
+ EXPECT_EQ(true, sampler->GetNormalizedCoords());
}
TEST_F(AmberScriptParserTest, SamplerUnexpectedParameter) {
@@ -162,5 +168,28 @@ TEST_F(AmberScriptParserTest, SamplerMaxLodSmallerThanMinLod) {
r.Error());
}
+TEST_F(AmberScriptParserTest, SamplerUnnormalizedCoordsSetsLod) {
+ std::string in = R"(
+SAMPLER sampler \
+ MIN_LOD 2.0 \
+ MAX_LOD 3.0 \
+ UNNORMALIZED_COORDS
+)";
+
+ Parser parser;
+ Result r = parser.Parse(in);
+ ASSERT_TRUE(r.IsSuccess());
+ auto script = parser.GetScript();
+ const auto& samplers = script->GetSamplers();
+ ASSERT_EQ(1U, samplers.size());
+
+ ASSERT_TRUE(samplers[0] != nullptr);
+ EXPECT_EQ("sampler", samplers[0]->GetName());
+
+ auto* sampler = samplers[0].get();
+ EXPECT_EQ(0.0f, sampler->GetMinLOD());
+ EXPECT_EQ(0.0f, sampler->GetMaxLOD());
+}
+
} // namespace amberscript
} // namespace amber
diff --git a/src/buffer.h b/src/buffer.h
index 6a0eca3..d867a5a 100644
--- a/src/buffer.h
+++ b/src/buffer.h
@@ -24,6 +24,7 @@
#include "amber/result.h"
#include "amber/value.h"
#include "src/format.h"
+#include "src/image.h"
namespace amber {
@@ -93,6 +94,15 @@ class Buffer {
uint32_t GetHeight() const { return height_; }
/// Set the number of elements high for the buffer.
void SetHeight(uint32_t height) { height_ = height; }
+ /// Get the number of elements this buffer is deep.
+ uint32_t GetDepth() const { return depth_; }
+ /// Set the number of elements this buffer is deep.
+ void SetDepth(uint32_t depth) { depth_ = depth; }
+
+ /// Get the image dimensionality.
+ ImageDimension GetImageDimension() const { return image_dim_; }
+ /// Set the image dimensionality.
+ void SetImageDimension(ImageDimension dim) { image_dim_ = dim; }
// | ---------- Element ---------- | ElementCount == 1
// | Value | Value | Value | Value | ValueCount == 4
@@ -233,13 +243,15 @@ class Buffer {
/// over all ubo, ssbo size and ssbo subdata size calls.
uint32_t max_size_in_bytes_ = 0;
uint32_t element_count_ = 0;
- uint32_t width_ = 0;
- uint32_t height_ = 0;
+ uint32_t width_ = 1;
+ uint32_t height_ = 1;
+ uint32_t depth_ = 1;
uint32_t mip_levels_ = 1;
bool format_is_default_ = false;
std::vector<uint8_t> bytes_;
Format* format_ = nullptr;
Sampler* sampler_ = nullptr;
+ ImageDimension image_dim_ = ImageDimension::kUnknown;
};
} // namespace amber
diff --git a/src/image.h b/src/image.h
new file mode 100644
index 0000000..a74c3ac
--- /dev/null
+++ b/src/image.h
@@ -0,0 +1,24 @@
+// Copyright 2019 The Amber 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.
+
+#ifndef SRC_IMAGE_H_
+#define SRC_IMAGE_H_
+
+namespace amber {
+
+enum class ImageDimension : int8_t { kUnknown = -1, k1D = 0, k2D = 1, k3D = 2 };
+
+} // namespace amber
+
+#endif // SRC_IMAGE_H_
diff --git a/src/sampler.h b/src/sampler.h
index f7c1aad..c44316a 100644
--- a/src/sampler.h
+++ b/src/sampler.h
@@ -71,6 +71,9 @@ class Sampler {
void SetAddressModeV(AddressMode mode) { address_mode_v_ = mode; }
AddressMode GetAddressModeV() const { return address_mode_v_; }
+ void SetAddressModeW(AddressMode mode) { address_mode_w_ = mode; }
+ AddressMode GetAddressModeW() const { return address_mode_w_; }
+
void SetBorderColor(BorderColor color) { border_color_ = color; }
BorderColor GetBorderColor() const { return border_color_; }
@@ -80,6 +83,9 @@ class Sampler {
void SetMaxLOD(float max_lod) { max_lod_ = max_lod; }
float GetMaxLOD() const { return max_lod_; }
+ void SetNormalizedCoords(bool norm) { normalized_coords_ = norm; }
+ bool GetNormalizedCoords() const { return normalized_coords_; }
+
private:
std::string name_;
FilterType min_filter_ = FilterType::kNearest;
@@ -87,9 +93,11 @@ class Sampler {
FilterType mipmap_mode_ = FilterType::kNearest;
AddressMode address_mode_u_ = AddressMode::kRepeat;
AddressMode address_mode_v_ = AddressMode::kRepeat;
+ AddressMode address_mode_w_ = AddressMode::kRepeat;
BorderColor border_color_ = BorderColor::kFloatTransparentBlack;
float min_lod_ = 0.0f;
float max_lod_ = 1.0f;
+ bool normalized_coords_ = true;
};
} // namespace amber
diff --git a/src/vulkan/frame_buffer.cc b/src/vulkan/frame_buffer.cc
index c73184f..d042bd4 100644
--- a/src/vulkan/frame_buffer.cc
+++ b/src/vulkan/frame_buffer.cc
@@ -63,8 +63,9 @@ Result FrameBuffer::Initialize(VkRenderPass render_pass,
for (auto* info : color_attachments_) {
color_images_.push_back(MakeUnique<TransferImage>(
device_, *info->buffer->GetFormat(), VK_IMAGE_ASPECT_COLOR_BIT,
- width_ << info->base_mip_level, height_ << info->base_mip_level,
- depth_, info->buffer->GetMipLevels(), info->base_mip_level, 1u));
+ VK_IMAGE_TYPE_2D, width_ << info->base_mip_level,
+ height_ << info->base_mip_level, depth_, info->buffer->GetMipLevels(),
+ info->base_mip_level, 1u));
Result r = color_images_.back()->Initialize(
VK_IMAGE_USAGE_TRANSFER_SRC_BIT | VK_IMAGE_USAGE_TRANSFER_DST_BIT |
@@ -83,7 +84,7 @@ Result FrameBuffer::Initialize(VkRenderPass render_pass,
? VK_IMAGE_ASPECT_DEPTH_BIT |
VK_IMAGE_ASPECT_STENCIL_BIT
: VK_IMAGE_ASPECT_DEPTH_BIT),
- width_, height_, depth_, 1u, 0u, 1u);
+ VK_IMAGE_TYPE_2D, width_, height_, depth_, 1u, 0u, 1u);
Result r = depth_image_->Initialize(
VK_IMAGE_USAGE_TRANSFER_SRC_BIT | VK_IMAGE_USAGE_TRANSFER_DST_BIT |
diff --git a/src/vulkan/image_descriptor.cc b/src/vulkan/image_descriptor.cc
index 3f689e5..ee52787 100644
--- a/src/vulkan/image_descriptor.cc
+++ b/src/vulkan/image_descriptor.cc
@@ -37,17 +37,9 @@ void ImageDescriptor::RecordCopyDataToResourceIfNeeded(CommandBuffer* command) {
BufferBackedDescriptor::RecordCopyDataToResourceIfNeeded(command);
- if (type_ == DescriptorType::kStorageImage) {
- // Change to general layout as it's required for storage images.
- transfer_image_->ImageBarrier(command, VK_IMAGE_LAYOUT_GENERAL,
- VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT);
- } else {
- // Use the earliest shader stage as we don't know which stage the image is
- // used in.
- transfer_image_->ImageBarrier(command,
- VK_IMAGE_LAYOUT_SHADER_READ_ONLY_OPTIMAL,
- VK_PIPELINE_STAGE_VERTEX_SHADER_BIT);
- }
+ // Just do this as early as possible.
+ transfer_image_->ImageBarrier(command, VK_IMAGE_LAYOUT_GENERAL,
+ VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT);
}
Result ImageDescriptor::CreateResourceIfNeeded() {
@@ -62,10 +54,27 @@ Result ImageDescriptor::CreateResourceIfNeeded() {
if (amber_buffer && amber_buffer->ValuePtr()->empty())
return {};
+ // Default to 2D image.
+ VkImageType image_type = VK_IMAGE_TYPE_2D;
+ switch (amber_buffer->GetImageDimension()) {
+ case ImageDimension::k1D:
+ image_type = VK_IMAGE_TYPE_1D;
+ break;
+ case ImageDimension::k2D:
+ image_type = VK_IMAGE_TYPE_2D;
+ break;
+ case ImageDimension::k3D:
+ image_type = VK_IMAGE_TYPE_3D;
+ break;
+ default:
+ break;
+ }
+
transfer_image_ = MakeUnique<TransferImage>(
device_, *amber_buffer->GetFormat(), VK_IMAGE_ASPECT_COLOR_BIT,
- amber_buffer->GetWidth(), amber_buffer->GetHeight(), 1u,
- amber_buffer->GetMipLevels(), base_mip_level_, VK_REMAINING_MIP_LEVELS);
+ image_type, amber_buffer->GetWidth(), amber_buffer->GetHeight(),
+ amber_buffer->GetDepth(), amber_buffer->GetMipLevels(), base_mip_level_,
+ VK_REMAINING_MIP_LEVELS);
VkImageUsageFlags usage =
VK_IMAGE_USAGE_TRANSFER_SRC_BIT | VK_IMAGE_USAGE_TRANSFER_DST_BIT;
@@ -113,10 +122,8 @@ void ImageDescriptor::UpdateDescriptorSetIfNeeded(
if (!is_descriptor_set_update_needed_)
return;
- VkImageLayout layout = VK_IMAGE_LAYOUT_SHADER_READ_ONLY_OPTIMAL;
-
- if (type_ == DescriptorType::kStorageImage)
- layout = VK_IMAGE_LAYOUT_GENERAL;
+ // Always use general layout.
+ VkImageLayout layout = VK_IMAGE_LAYOUT_GENERAL;
VkDescriptorImageInfo image_info = {vulkan_sampler_.GetVkSampler(),
transfer_image_->GetVkImageView(),
diff --git a/src/vulkan/sampler.cc b/src/vulkan/sampler.cc
index 0d514c7..979c098 100644
--- a/src/vulkan/sampler.cc
+++ b/src/vulkan/sampler.cc
@@ -70,10 +70,12 @@ Result Sampler::CreateSampler(amber::Sampler* sampler) {
: VK_SAMPLER_MIPMAP_MODE_NEAREST;
sampler_info.addressModeU = GetVkAddressMode(sampler->GetAddressModeU());
sampler_info.addressModeV = GetVkAddressMode(sampler->GetAddressModeV());
- sampler_info.addressModeW = VK_SAMPLER_ADDRESS_MODE_REPEAT;
+ sampler_info.addressModeW = GetVkAddressMode(sampler->GetAddressModeW());
sampler_info.borderColor = GetVkBorderColor(sampler->GetBorderColor());
sampler_info.minLod = sampler->GetMinLOD();
sampler_info.maxLod = sampler->GetMaxLOD();
+ sampler_info.unnormalizedCoordinates =
+ (sampler->GetNormalizedCoords() ? VK_FALSE : VK_TRUE);
if (device_->GetPtrs()->vkCreateSampler(device_->GetVkDevice(), &sampler_info,
nullptr, &sampler_) != VK_SUCCESS) {
diff --git a/src/vulkan/transfer_image.cc b/src/vulkan/transfer_image.cc
index ca79514..c7d0058 100644
--- a/src/vulkan/transfer_image.cc
+++ b/src/vulkan/transfer_image.cc
@@ -48,6 +48,7 @@ const VkImageCreateInfo kDefaultImageInfo = {
TransferImage::TransferImage(Device* device,
const Format& format,
VkImageAspectFlags aspect,
+ VkImageType image_type,
uint32_t x,
uint32_t y,
uint32_t z,
@@ -61,6 +62,7 @@ TransferImage::TransferImage(Device* device,
base_mip_level_(base_mip_level),
used_mip_levels_(used_mip_levels) {
image_info_.format = device_->GetVkFormat(format);
+ image_info_.imageType = image_type;
image_info_.extent = {x, y, z};
image_info_.mipLevels = mip_levels;
}
@@ -133,12 +135,29 @@ Result TransferImage::Initialize(VkImageUsageFlags usage) {
return MapMemory(host_accessible_memory_);
}
+VkImageViewType TransferImage::GetImageViewType() const {
+ // TODO(alan-baker): handle other view types.
+ // 1D-array, 2D-array, Cube, Cube-array.
+ switch (image_info_.imageType) {
+ case VK_IMAGE_TYPE_1D:
+ return VK_IMAGE_VIEW_TYPE_1D;
+ case VK_IMAGE_TYPE_2D:
+ return VK_IMAGE_VIEW_TYPE_2D;
+ case VK_IMAGE_TYPE_3D:
+ return VK_IMAGE_VIEW_TYPE_3D;
+ default:
+ break;
+ }
+
+ // Default to 2D image view.
+ return VK_IMAGE_VIEW_TYPE_2D;
+}
+
Result TransferImage::CreateVkImageView() {
VkImageViewCreateInfo image_view_info = VkImageViewCreateInfo();
image_view_info.sType = VK_STRUCTURE_TYPE_IMAGE_VIEW_CREATE_INFO;
image_view_info.image = image_;
- // TODO(jaebaek): Set .viewType correctly
- image_view_info.viewType = VK_IMAGE_VIEW_TYPE_2D;
+ image_view_info.viewType = GetImageViewType();
image_view_info.format = image_info_.format;
image_view_info.components = {
VK_COMPONENT_SWIZZLE_R,
@@ -178,7 +197,8 @@ VkBufferImageCopy TransferImage::CreateBufferImageCopy(uint32_t mip_level) {
};
copy_region.imageOffset = {0, 0, 0};
copy_region.imageExtent = {image_info_.extent.width >> mip_level,
- image_info_.extent.height >> mip_level, 1};
+ image_info_.extent.height >> mip_level,
+ image_info_.extent.depth};
return copy_region;
}
diff --git a/src/vulkan/transfer_image.h b/src/vulkan/transfer_image.h
index 2beb9d8..e68b88e 100644
--- a/src/vulkan/transfer_image.h
+++ b/src/vulkan/transfer_image.h
@@ -32,6 +32,7 @@ class TransferImage : public Resource {
TransferImage(Device* device,
const Format& format,
VkImageAspectFlags aspect,
+ VkImageType image_type,
uint32_t x,
uint32_t y,
uint32_t z,
@@ -63,6 +64,8 @@ class TransferImage : public Resource {
uint32_t* memory_type_index);
VkBufferImageCopy CreateBufferImageCopy(uint32_t mip_level);
+ VkImageViewType GetImageViewType() const;
+
/// An extra `VkBuffer` is used to facilitate the transfer of data from the
/// host into the `VkImage` on the device.
VkBuffer host_accessible_buffer_ = VK_NULL_HANDLE;
diff --git a/tests/cases/glsl_read_and_write_image3d_rgba32i.amber b/tests/cases/glsl_read_and_write_image3d_rgba32i.amber
new file mode 100644
index 0000000..75e8d8b
--- /dev/null
+++ b/tests/cases/glsl_read_and_write_image3d_rgba32i.amber
@@ -0,0 +1,76 @@
+#!amber
+# Copyright 2019 The Amber 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
+#
+# https://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.
+
+SHADER compute write GLSL
+#version 450
+layout(set=0, binding=0, rgba32i) uniform iimage3D im3d;
+void main() {
+ uvec3 gid = gl_GlobalInvocationID;
+ ivec3 coord = ivec3(gid.x, gid.y, gid.z);
+ ivec4 data = ivec4(gid.x + 1, gid.y + 1, gid.z + 1, 0);
+ imageStore(im3d, coord, data);
+}
+END
+
+SHADER compute read GLSL
+#version 450
+
+layout(set=0, binding=0) uniform itexture3D im3d;
+layout(set=0, binding=1) uniform sampler s;
+layout(set=0, binding=2) buffer A { ivec4 x[]; } data;
+
+void main() {
+ uvec3 gid = gl_GlobalInvocationID;
+ ivec3 coord = ivec3(gid.x, gid.y, gid.z);
+ uint linear = 4 * gid.z + 2 * gid.y + gid.x;
+ data.x[linear] = texture(isampler3D(im3d, s), coord);
+}
+END
+
+BUFFER out_buf DATA_TYPE vec4<int32> SIZE 8 FILL 15
+IMAGE im3d DATA_TYPE vec4<int32> DIM_3D \
+ WIDTH 2 HEIGHT 2 DEPTH 2 FILL 0
+SAMPLER sampler \
+ ADDRESS_MODE_U clamp_to_edge \
+ ADDRESS_MODE_V clamp_to_edge \
+ ADDRESS_MODE_W clamp_to_edge \
+ MIN_FILTER nearest \
+ MAG_FILTER nearest \
+ MIN_LOD 0.0 \
+ MAX_LOD 0.0
+
+PIPELINE compute write_pipe
+ ATTACH write
+ BIND BUFFER im3d AS storage_image DESCRIPTOR_SET 0 BINDING 0
+END
+
+PIPELINE compute read_pipe
+ ATTACH read
+ BIND BUFFER im3d AS sampled_image DESCRIPTOR_SET 0 BINDING 0
+ BIND SAMPLER sampler DESCRIPTOR_SET 0 BINDING 1
+ BIND BUFFER out_buf AS storage DESCRIPTOR_SET 0 BINDING 2
+END
+
+RUN write_pipe 2 2 2
+RUN read_pipe 2 2 2
+
+EXPECT out_buf IDX 0 EQ 1 1 1 0
+EXPECT out_buf IDX 16 EQ 2 1 1 0
+EXPECT out_buf IDX 32 EQ 1 2 1 0
+EXPECT out_buf IDX 48 EQ 2 2 1 0
+EXPECT out_buf IDX 64 EQ 1 1 2 0
+EXPECT out_buf IDX 80 EQ 2 1 2 0
+EXPECT out_buf IDX 96 EQ 1 2 2 0
+EXPECT out_buf IDX 112 EQ 2 2 2 0
diff --git a/tests/cases/opencl_read_and_write_image3d_rgba32i.amber b/tests/cases/opencl_read_and_write_image3d_rgba32i.amber
new file mode 100644
index 0000000..9c84a58
--- /dev/null
+++ b/tests/cases/opencl_read_and_write_image3d_rgba32i.amber
@@ -0,0 +1,74 @@
+#!amber
+# Copyright 2019 The Amber 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
+#
+# https://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.
+
+SHADER compute write OPENCL-C
+#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable
+kernel void write_foo(write_only image3d_t im) {
+ int4 dim = get_image_dim(im);
+ int gid_x = get_global_id(0);
+ int gid_y = get_global_id(1);
+ int gid_z = get_global_id(2);
+ int4 coord = (int4)(gid_x, gid_y, gid_z, 0);
+ int4 data = (int4)(gid_x + 1, gid_y + 1, gid_z + 1, 0);
+ write_imagei(im, coord, data);
+}
+END
+
+SHADER compute read OPENCL-C
+kernel void read_foo(read_only image3d_t im, sampler_t sampler, global int4* out) {
+ int gid_x = get_global_id(0);
+ int gid_y = get_global_id(1);
+ int gid_z = get_global_id(2);
+ float4 coord = (float4)(gid_x, gid_y, gid_z, 0);
+ int linear = 4 * gid_z + 2 * gid_y + gid_x;
+ out[linear] = read_imagei(im, sampler, coord);
+}
+END
+
+BUFFER out_buf DATA_TYPE vec4<int32> SIZE 8 FILL 15
+IMAGE im3d DATA_TYPE vec4<int32> DIM_3D \
+ WIDTH 2 HEIGHT 2 DEPTH 2 FILL 0
+SAMPLER sampler \
+ ADDRESS_MODE_U clamp_to_edge \
+ ADDRESS_MODE_V clamp_to_edge \
+ ADDRESS_MODE_W clamp_to_edge \
+ MIN_FILTER nearest \
+ MAG_FILTER nearest \
+ MIN_LOD 0.0 \
+ MAX_LOD 0.0
+
+PIPELINE compute write_pipe
+ ATTACH write ENTRY_POINT write_foo
+ BIND BUFFER im3d KERNEL ARG_NAME im
+END
+
+PIPELINE compute read_pipe
+ ATTACH read ENTRY_POINT read_foo
+ BIND BUFFER im3d KERNEL ARG_NAME im
+ BIND SAMPLER sampler KERNEL ARG_NAME sampler
+ BIND BUFFER out_buf KERNEL ARG_NAME out
+END
+
+RUN write_pipe 2 2 2
+RUN read_pipe 2 2 2
+
+EXPECT out_buf IDX 0 EQ 1 1 1 0
+EXPECT out_buf IDX 16 EQ 2 1 1 0
+EXPECT out_buf IDX 32 EQ 1 2 1 0
+EXPECT out_buf IDX 48 EQ 2 2 1 0
+EXPECT out_buf IDX 64 EQ 1 1 2 0
+EXPECT out_buf IDX 80 EQ 2 1 2 0
+EXPECT out_buf IDX 96 EQ 1 2 2 0
+EXPECT out_buf IDX 112 EQ 2 2 2 0
diff --git a/tests/cases/opencl_read_image.amber b/tests/cases/opencl_read_image.amber
index dfb7f16..a25a298 100644
--- a/tests/cases/opencl_read_image.amber
+++ b/tests/cases/opencl_read_image.amber
@@ -33,7 +33,7 @@ kernel void foo(read_only image2d_t image, sampler_t sampler, global float4* out
}
END
-BUFFER texture DATA_TYPE vec4<float> SIZE 4 FILL 0.0
+IMAGE texture DATA_TYPE vec4<float> DIM_2D WIDTH 2 HEIGHT 2 FILL 0.0
BUFFER out_buf DATA_TYPE vec4<float> DATA
2.0 2.0 2.0 2.0
2.0 2.0 2.0 2.0
diff --git a/tests/cases/opencl_write_image.amber b/tests/cases/opencl_write_image.amber
index 62a3e9c..28ada5a 100644
--- a/tests/cases/opencl_write_image.amber
+++ b/tests/cases/opencl_write_image.amber
@@ -23,7 +23,7 @@ kernel void foo(write_only image2d_t image, global float4* data) {
}
END
-BUFFER texture DATA_TYPE vec4<float> WIDTH 2 HEIGHT 2 FILL 0.0
+IMAGE texture DATA_TYPE vec4<float> DIM_2D WIDTH 2 HEIGHT 2 FILL 0.0
BUFFER data DATA_TYPE vec4<float> DATA
1.0 2.0 3.0 4.0
2.0 3.0 4.0 1.0
diff --git a/tests/run_tests.py b/tests/run_tests.py
index f461bff..e514fd7 100755
--- a/tests/run_tests.py
+++ b/tests/run_tests.py
@@ -80,12 +80,15 @@ SUPPRESSIONS_SWIFTSHADER = [
# Exceeded maxBoundDescriptorSets limit of physical device
"multiple_ssbo_with_sparse_descriptor_set_in_compute_pipeline.vkscript",
# shaderStorageImageWriteWithoutFormat but is not enabled on the device
+ "opencl_read_and_write_image3d_rgba32i.amber",
"opencl_write_image.amber",
+ "glsl_read_and_write_image3d_rgba32i.amber",
]
OPENCL_CASES = [
"opencl_bind_buffer.amber",
"opencl_c_copy.amber",
+ "opencl_read_and_write_image3d_rgba32i.amber",
"opencl_read_image.amber",
"opencl_set_arg.amber",
"opencl_write_image.amber",
generated by cgit v1.2.3 (git 2.25.1) at 2024-07-09 09:21:27 +0000