diff options
-rw-r--r-- | docs/amber_script.md | 36 | ||||
-rw-r--r-- | src/CMakeLists.txt | 1 | ||||
-rw-r--r-- | src/amberscript/parser.cc | 174 | ||||
-rw-r--r-- | src/amberscript/parser.h | 1 | ||||
-rw-r--r-- | src/amberscript/parser_image_test.cc | 272 | ||||
-rw-r--r-- | src/amberscript/parser_sampler_test.cc | 31 | ||||
-rw-r--r-- | src/buffer.h | 16 | ||||
-rw-r--r-- | src/image.h | 24 | ||||
-rw-r--r-- | src/sampler.h | 8 | ||||
-rw-r--r-- | src/vulkan/frame_buffer.cc | 7 | ||||
-rw-r--r-- | src/vulkan/image_descriptor.cc | 41 | ||||
-rw-r--r-- | src/vulkan/sampler.cc | 4 | ||||
-rw-r--r-- | src/vulkan/transfer_image.cc | 26 | ||||
-rw-r--r-- | src/vulkan/transfer_image.h | 3 | ||||
-rw-r--r-- | tests/cases/glsl_read_and_write_image3d_rgba32i.amber | 76 | ||||
-rw-r--r-- | tests/cases/opencl_read_and_write_image3d_rgba32i.amber | 74 | ||||
-rw-r--r-- | tests/cases/opencl_read_image.amber | 2 | ||||
-rw-r--r-- | tests/cases/opencl_write_image.amber | 2 | ||||
-rwxr-xr-x | tests/run_tests.py | 3 |
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", |