From 26f80d727aa1aa54c265fe5e3cc51fb7dab09509 Mon Sep 17 00:00:00 2001 From: Alan Baker Date: Thu, 18 Jul 2019 16:37:31 -0400 Subject: [PATCH 1/9] Add OpenCL BIND commands Fixes #426 * New BIND subcommands to specify kernel args by name or ordinal * clspv compiles parse descriptor map (partially) into Pipeline::ShaderInfo * Vulkan engine binds OpenCL buffers --- src/amberscript/parser.cc | 51 +++++++--- src/amberscript/parser_bind_test.cc | 142 ++++++++++++++++++++++++++++ src/clspv_helper.cc | 50 +++++++++- src/clspv_helper.h | 3 +- src/executor.cc | 2 +- src/pipeline.cc | 30 ++++++ src/pipeline.h | 35 +++++++ src/shader_compiler.cc | 11 ++- src/shader_compiler.h | 5 +- src/shader_compiler_test.cc | 21 ++-- src/vulkan/engine_vulkan.cc | 24 ++++- 11 files changed, 341 insertions(+), 33 deletions(-) diff --git a/src/amberscript/parser.cc b/src/amberscript/parser.cc index 412f33fab..04d2a5595 100644 --- a/src/amberscript/parser.cc +++ b/src/amberscript/parser.cc @@ -687,22 +687,47 @@ Result Parser::ParsePipelineBind(Pipeline* pipeline) { return Result("buffer type does not match intended usage"); token = tokenizer_->NextToken(); - if (!token->IsString() || token->AsString() != "DESCRIPTOR_SET") - return Result("missing DESCRIPTOR_SET for BIND command"); + if (token->IsString() && token->AsString() == "DESCRIPTOR_SET") { + token = tokenizer_->NextToken(); + if (!token->IsInteger()) + return Result("invalid value for DESCRIPTOR_SET in BIND command"); + uint32_t descriptor_set = token->AsUint32(); - token = tokenizer_->NextToken(); - if (!token->IsInteger()) - return Result("invalid value for DESCRIPTOR_SET in BIND command"); - uint32_t descriptor_set = token->AsUint32(); + token = tokenizer_->NextToken(); + if (!token->IsString() || token->AsString() != "BINDING") + return Result("missing BINDING for BIND command"); - token = tokenizer_->NextToken(); - if (!token->IsString() || token->AsString() != "BINDING") - return Result("missing BINDING for BIND command"); + token = tokenizer_->NextToken(); + if (!token->IsInteger()) + return Result("invalid value for BINDING in BIND command"); + pipeline->AddBuffer(buffer, descriptor_set, token->AsUint32()); + } else { + if (!token->IsString()) + return Result("missing DESCRIPTOR_SET for BIND command"); - token = tokenizer_->NextToken(); - if (!token->IsInteger()) - return Result("invalid value for BINDING in BIND command"); - pipeline->AddBuffer(buffer, descriptor_set, token->AsUint32()); + if (token->AsString() != "KERNEL") + return Result("missing DESCRIPTOR_SET for BIND command"); + + token = tokenizer_->NextToken(); + if (!token->IsString()) + return Result("missing kernel arg identifier"); + + if (token->AsString() == "ARG") { + token = tokenizer_->NextToken(); + if (!token->IsString()) + return Result("expected argument identifier"); + + pipeline->AddBuffer(buffer, token->AsString()); + } else if (token->AsString() == "ARGNO") { + token = tokenizer_->NextToken(); + if (!token->IsInteger()) + return Result("expected argument identifier number"); + + pipeline->AddBuffer(buffer, token->AsUint32()); + } else { + return Result("missing ARG or ARGNO keyword"); + } + } } return ValidateEndOfStatement("BIND command"); diff --git a/src/amberscript/parser_bind_test.cc b/src/amberscript/parser_bind_test.cc index 767229417..ee8d9ac0b 100644 --- a/src/amberscript/parser_bind_test.cc +++ b/src/amberscript/parser_bind_test.cc @@ -1149,5 +1149,147 @@ END)"; EXPECT_EQ("12: extra parameters after BIND command", r.Error()); } +TEST_F(AmberScriptParserTest, BindBufferOpenCLArgName) { + std::string in = R"( +SHADER compute my_shader OPENCL-C +#shader +END +BUFFER my_buf DATA_TYPE uint32 DATA 1 END + +PIPELINE compute my_pipeline + ATTACH my_shader + BIND BUFFER my_buf AS storage KERNEL ARG arg +END)"; + + Parser parser; + Result r = parser.Parse(in); + ASSERT_TRUE(r.IsSuccess()); +} + +TEST_F(AmberScriptParserTest, BindBufferOpenCLArgNo) { + std::string in = R"( +SHADER compute my_shader OPENCL-C +#shader +END +BUFFER my_buf DATA_TYPE uint32 DATA 1 END + +PIPELINE compute my_pipeline + ATTACH my_shader + BIND BUFFER my_buf AS storage KERNEL ARGNO 0 +END)"; + + Parser parser; + Result r = parser.Parse(in); + ASSERT_TRUE(r.IsSuccess()); +} + +TEST_F(AmberScriptParserTest, BindBufferOpenCLMissingKernel) { + std::string in = R"( +SHADER compute my_shader OPENCL-C +#shader +END +BUFFER my_buf DATA_TYPE uint32 DATA 1 END + +PIPELINE compute my_pipeline + ATTACH my_shader + BIND BUFFER my_buf AS storage ARG arg +END)"; + + Parser parser; + Result r = parser.Parse(in); + ASSERT_FALSE(r.IsSuccess()); + EXPECT_EQ("9: missing DESCRIPTOR_SET for BIND command", r.Error()); +} + +TEST_F(AmberScriptParserTest, BindBufferOpenCLMissingArg) { + std::string in = R"( +SHADER compute my_shader OPENCL-C +#shader +END +BUFFER my_buf DATA_TYPE uint32 DATA 1 END + +PIPELINE compute my_pipeline + ATTACH my_shader + BIND BUFFER my_buf AS storage KERNEL arg +END)"; + + Parser parser; + Result r = parser.Parse(in); + ASSERT_FALSE(r.IsSuccess()); + EXPECT_EQ("9: missing ARG or ARGNO keyword", r.Error()); +} + +TEST_F(AmberScriptParserTest, BindBufferOpenCLMissingArgName) { + std::string in = R"( +SHADER compute my_shader OPENCL-C +#shader +END +BUFFER my_buf DATA_TYPE uint32 DATA 1 END + +PIPELINE compute my_pipeline + ATTACH my_shader + BIND BUFFER my_buf AS storage KERNEL ARG +END)"; + + Parser parser; + Result r = parser.Parse(in); + ASSERT_FALSE(r.IsSuccess()); + EXPECT_EQ("10: expected argument identifier", r.Error()); +} + +TEST_F(AmberScriptParserTest, BindBufferOpenCLMissingArgNo) { + std::string in = R"( +SHADER compute my_shader OPENCL-C +#shader +END +BUFFER my_buf DATA_TYPE uint32 DATA 1 END + +PIPELINE compute my_pipeline + ATTACH my_shader + BIND BUFFER my_buf AS storage KERNEL ARGNO +END)"; + + Parser parser; + Result r = parser.Parse(in); + ASSERT_FALSE(r.IsSuccess()); + EXPECT_EQ("10: expected argument identifier number", r.Error()); +} + +TEST_F(AmberScriptParserTest, BindBufferOpenCLArgNameNotString) { + std::string in = R"( +SHADER compute my_shader OPENCL-C +#shader +END +BUFFER my_buf DATA_TYPE uint32 DATA 1 END + +PIPELINE compute my_pipeline + ATTACH my_shader + BIND BUFFER my_buf AS storage KERNEL ARG 0 +END)"; + + Parser parser; + Result r = parser.Parse(in); + ASSERT_FALSE(r.IsSuccess()); + EXPECT_EQ("9: expected argument identifier", r.Error()); +} + +TEST_F(AmberScriptParserTest, BindBufferOpenCLArgNoNotInteger) { + std::string in = R"( +SHADER compute my_shader OPENCL-C +#shader +END +BUFFER my_buf DATA_TYPE uint32 DATA 1 END + +PIPELINE compute my_pipeline + ATTACH my_shader + BIND BUFFER my_buf AS storage KERNEL ARGNO in +END)"; + + Parser parser; + Result r = parser.Parse(in); + ASSERT_FALSE(r.IsSuccess()); + EXPECT_EQ("9: expected argument identifier number", r.Error()); +} + } // namespace amberscript } // namespace amber diff --git a/src/clspv_helper.cc b/src/clspv_helper.cc index 056b5d511..058ab80a7 100644 --- a/src/clspv_helper.cc +++ b/src/clspv_helper.cc @@ -14,20 +14,66 @@ #include "src/clspv_helper.h" +#include "clspv/ArgKind.h" #include "clspv/Compiler.h" namespace amber { namespace clspvhelper { -Result Compile(const std::string& src_str, +Result Compile(Pipeline::ShaderInfo* shader_info, std::vector* generated_binary) { - // TODO(alan-baker): Parse the descriptor map. std::vector entries; + const auto& src_str = shader_info->GetShader()->GetData(); if (clspv::CompileFromSourceString(src_str, "", "", generated_binary, &entries)) { return Result("Clspv compile failed"); } + for (auto& entry : entries) { + if (entry.kind != clspv::version0::DescriptorMapEntry::KernelArg) { + return Result( + "Only kernel argument descriptor entries are currently supported"); + } + + Pipeline::ShaderInfo::DescriptorMapEntry descriptor_entry; + descriptor_entry.descriptor_set = entry.descriptor_set; + descriptor_entry.binding = entry.binding; + descriptor_entry.pod_offset = 0; + descriptor_entry.pod_arg_size = 0; + switch (entry.kernel_arg_data.arg_kind) { + case clspv::ArgKind::Buffer: + descriptor_entry.kind = + Pipeline::ShaderInfo::DescriptorMapEntry::Kind::SSBO; + break; + case clspv::ArgKind::BufferUBO: + descriptor_entry.kind = + Pipeline::ShaderInfo::DescriptorMapEntry::Kind::UBO; + break; + case clspv::ArgKind::Pod: + descriptor_entry.kind = + Pipeline::ShaderInfo::DescriptorMapEntry::Kind::Pod; + break; + case clspv::ArgKind::PodUBO: + descriptor_entry.kind = + Pipeline::ShaderInfo::DescriptorMapEntry::Kind::PodUBO; + break; + default: + return Result("Unsupported kernel argument descriptor entry"); + } + + if (entry.kernel_arg_data.arg_kind == clspv::ArgKind::Pod || + entry.kernel_arg_data.arg_kind == clspv::ArgKind::PodUBO) { + descriptor_entry.pod_offset = entry.kernel_arg_data.pod_offset; + descriptor_entry.pod_arg_size = entry.kernel_arg_data.pod_arg_size; + } + + descriptor_entry.arg_name = entry.kernel_arg_data.arg_name; + descriptor_entry.arg_ordinal = entry.kernel_arg_data.arg_ordinal; + + shader_info->AddDescriptorEntry(entry.kernel_arg_data.kernel_name, + std::move(descriptor_entry)); + } + return Result(); } diff --git a/src/clspv_helper.h b/src/clspv_helper.h index 367253213..0df2d0594 100644 --- a/src/clspv_helper.h +++ b/src/clspv_helper.h @@ -19,13 +19,14 @@ #include #include "amber/result.h" +#include "src/pipeline.h" namespace amber { namespace clspvhelper { // Passes the OpenCL C source code to Clspv. // Returns the generated SPIR-V binary via |generated_binary| argument. -Result Compile(const std::string& src_str, +Result Compile(Pipeline::ShaderInfo* shader_info, std::vector* generated_binary); } // namespace clspvhelper diff --git a/src/executor.cc b/src/executor.cc index f74a93b0f..490109274 100644 --- a/src/executor.cc +++ b/src/executor.cc @@ -37,7 +37,7 @@ Result Executor::CompileShaders(const amber::Script* script, Result r; std::vector data; - std::tie(r, data) = sc.Compile(shader_info.GetShader(), shader_map); + std::tie(r, data) = sc.Compile(&shader_info, shader_map); if (!r.IsSuccess()) return r; diff --git a/src/pipeline.cc b/src/pipeline.cc index 1a32a42f2..4a6e837ef 100644 --- a/src/pipeline.cc +++ b/src/pipeline.cc @@ -330,4 +330,34 @@ void Pipeline::AddBuffer(Buffer* buf, info.binding = binding; } +void Pipeline::AddBuffer(Buffer* buf, const std::string& arg_name) { + // If this buffer binding already exists, overwrite with the new buffer. + for (auto& info : buffers_) { + if (info.arg_name == arg_name) { + info.buffer = buf; + return; + } + } + + buffers_.push_back(BufferInfo{buf}); + + auto& info = buffers_.back(); + info.arg_name = arg_name; +} + +void Pipeline::AddBuffer(Buffer* buf, uint32_t arg_no) { + // If this buffer binding already exists, overwrite with the new buffer. + for (auto& info : buffers_) { + if (info.arg_no == arg_no) { + info.buffer = buf; + return; + } + } + + buffers_.push_back(BufferInfo{buf}); + + auto& info = buffers_.back(); + info.arg_no = arg_no; +} + } // namespace amber diff --git a/src/pipeline.h b/src/pipeline.h index c17a67eb1..239defa97 100644 --- a/src/pipeline.h +++ b/src/pipeline.h @@ -18,6 +18,7 @@ #include #include #include +#include #include #include @@ -67,6 +68,33 @@ class Pipeline { specialization_[spec_id] = value; } + /// Descriptor information for an OpenCL-C shader. + struct DescriptorMapEntry { + std::string arg_name; + + enum class Kind : int { + SSBO, + UBO, + Pod, + PodUBO, + } kind; + + uint32_t descriptor_set; + uint32_t binding; + uint32_t arg_ordinal; + uint32_t pod_offset; + uint32_t pod_arg_size; + }; + + void AddDescriptorEntry(const std::string& kernel, + DescriptorMapEntry&& entry) { + descriptor_map_[kernel].emplace_back(std::move(entry)); + } + const std::unordered_map>& + GetDescriptorMap() const { + return descriptor_map_; + } + private: Shader* shader_ = nullptr; ShaderType shader_type_; @@ -74,6 +102,7 @@ class Pipeline { std::string entry_point_; std::vector data_; std::map specialization_; + std::unordered_map> descriptor_map_; }; /// Information on a buffer attached to the pipeline. @@ -88,6 +117,8 @@ class Pipeline { uint32_t descriptor_set = 0; uint32_t binding = 0; uint32_t location = 0; + std::string arg_name = ""; + uint32_t arg_no = 0; }; static const char* kGeneratedColorBuffer; @@ -165,6 +196,10 @@ class Pipeline { /// Adds |buf| to the pipeline at the given |descriptor_set| and |binding|. void AddBuffer(Buffer* buf, uint32_t descriptor_set, uint32_t binding); + /// Adds |buf| to the pipeline at the given |arg_name|. + void AddBuffer(Buffer* buf, const std::string& arg_name); + /// Adds |buf| to the pipeline at the given |arg_no|. + void AddBuffer(Buffer* buf, uint32_t arg_no); /// Returns information on all buffers in this pipeline. const std::vector& GetBuffers() const { return buffers_; } diff --git a/src/shader_compiler.cc b/src/shader_compiler.cc index 8bf40ae24..916ad9aa7 100644 --- a/src/shader_compiler.cc +++ b/src/shader_compiler.cc @@ -51,8 +51,9 @@ ShaderCompiler::ShaderCompiler(const std::string& env) : spv_env_(env) {} ShaderCompiler::~ShaderCompiler() = default; std::pair> ShaderCompiler::Compile( - const Shader* shader, + Pipeline::ShaderInfo* shader_info, const ShaderMap& shader_map) const { + const auto shader = shader_info->GetShader(); auto it = shader_map.find(shader->GetName()); if (it != shader_map.end()) return {{}, it->second}; @@ -122,7 +123,7 @@ std::pair> ShaderCompiler::Compile( #if AMBER_ENABLE_CLSPV } else if (shader->GetFormat() == kShaderFormatOpenCLC) { - Result r = CompileOpenCLC(shader, &results); + Result r = CompileOpenCLC(shader_info, &results); if (!r.IsSuccess()) return {r, {}}; #endif // AMBER_ENABLE_CLSPV @@ -241,12 +242,12 @@ Result ShaderCompiler::CompileHlsl(const Shader*, #endif // AMBER_ENABLE_DXC #if AMBER_ENABLE_CLSPV -Result ShaderCompiler::CompileOpenCLC(const Shader* shader, +Result ShaderCompiler::CompileOpenCLC(Pipeline::ShaderInfo* shader_info, std::vector* result) const { - return clspvhelper::Compile(shader->GetData(), result); + return clspvhelper::Compile(shader_info, result); } #else -Result ShaderCompiler::CompileOpenCLC(const Shader*, +Result ShaderCompiler::CompileOpenCLC(Pipeline::ShaderInfo*, std::vector*) const { return {}; } diff --git a/src/shader_compiler.h b/src/shader_compiler.h index 3c4b56fbb..02769b46c 100644 --- a/src/shader_compiler.h +++ b/src/shader_compiler.h @@ -22,6 +22,7 @@ #include "amber/amber.h" #include "amber/result.h" #include "src/shader.h" +#include "src/pipeline.h" namespace amber { @@ -37,14 +38,14 @@ class ShaderCompiler { /// compilation result is copied from that entry. Otherwise a compiler is /// invoked to produce the compilation result. std::pair> Compile( - const Shader* shader, + Pipeline::ShaderInfo* shader_info, const ShaderMap& shader_map) const; private: Result ParseHex(const std::string& data, std::vector* result) const; Result CompileGlsl(const Shader* shader, std::vector* result) const; Result CompileHlsl(const Shader* shader, std::vector* result) const; - Result CompileOpenCLC(const Shader* shader, + Result CompileOpenCLC(Pipeline::ShaderInfo* shader, std::vector* result) const; std::string spv_env_; diff --git a/src/shader_compiler_test.cc b/src/shader_compiler_test.cc index 188a74109..06dae021b 100644 --- a/src/shader_compiler_test.cc +++ b/src/shader_compiler_test.cc @@ -105,7 +105,8 @@ void main() { ShaderCompiler sc; Result r; std::vector binary; - std::tie(r, binary) = sc.Compile(&shader, ShaderMap()); + Pipeline::ShaderInfo shader_info(&shader, kShaderTypeCompute); + std::tie(r, binary) = sc.Compile(&shader_info, ShaderMap()); ASSERT_TRUE(r.IsSuccess()) << r.Error(); EXPECT_FALSE(binary.empty()); EXPECT_EQ(0x07230203, binary[0]); // Verify SPIR-V header present. @@ -122,7 +123,8 @@ TEST_F(ShaderCompilerTest, CompilesSpirvAsm) { ShaderCompiler sc; Result r; std::vector binary; - std::tie(r, binary) = sc.Compile(&shader, ShaderMap()); + Pipeline::ShaderInfo shader_info(&shader, kShaderTypeCompute); + std::tie(r, binary) = sc.Compile(&shader_info, ShaderMap()); ASSERT_TRUE(r.IsSuccess()); EXPECT_FALSE(binary.empty()); EXPECT_EQ(0x07230203, binary[0]); // Verify SPIR-V header present. @@ -140,7 +142,8 @@ TEST_F(ShaderCompilerTest, InvalidSpirvHex) { ShaderCompiler sc; Result r; std::vector binary; - std::tie(r, binary) = sc.Compile(&shader, ShaderMap()); + Pipeline::ShaderInfo shader_info(&shader, kShaderTypeCompute); + std::tie(r, binary) = sc.Compile(&shader_info, ShaderMap()); ASSERT_FALSE(r.IsSuccess()); EXPECT_EQ("Invalid shader: error: line 0: Invalid SPIR-V magic number.\n", r.Error()); @@ -155,7 +158,8 @@ TEST_F(ShaderCompilerTest, InvalidHex) { ShaderCompiler sc; Result r; std::vector binary; - std::tie(r, binary) = sc.Compile(&shader, ShaderMap()); + Pipeline::ShaderInfo shader_info(&shader, kShaderTypeCompute); + std::tie(r, binary) = sc.Compile(&shader_info, ShaderMap()); ASSERT_FALSE(r.IsSuccess()); EXPECT_EQ("Invalid shader: error: line 0: Invalid SPIR-V magic number.\n", r.Error()); @@ -171,7 +175,8 @@ TEST_F(ShaderCompilerTest, CompilesSpirvHex) { ShaderCompiler sc; Result r; std::vector binary; - std::tie(r, binary) = sc.Compile(&shader, ShaderMap()); + Pipeline::ShaderInfo shader_info(&shader, kShaderTypeCompute); + std::tie(r, binary) = sc.Compile(&shader_info, ShaderMap()); ASSERT_TRUE(r.IsSuccess()); EXPECT_FALSE(binary.empty()); EXPECT_EQ(0x07230203, binary[0]); // Verify SPIR-V header present. @@ -188,7 +193,8 @@ TEST_F(ShaderCompilerTest, FailsOnInvalidShader) { ShaderCompiler sc; Result r; std::vector binary; - std::tie(r, binary) = sc.Compile(&shader, ShaderMap()); + Pipeline::ShaderInfo shader_info(&shader, kShaderTypeCompute); + std::tie(r, binary) = sc.Compile(&shader_info, ShaderMap()); ASSERT_FALSE(r.IsSuccess()); } @@ -211,7 +217,8 @@ TEST_F(ShaderCompilerTest, ReturnsCachedShader) { ShaderCompiler sc; Result r; std::vector binary; - std::tie(r, binary) = sc.Compile(&shader, map); + Pipeline::ShaderInfo shader_info(&shader, kShaderTypeCompute); + std::tie(r, binary) = sc.Compile(&shader_info, map); ASSERT_TRUE(r.IsSuccess()) << r.Error(); ASSERT_EQ(binary.size(), src_bytes.size()); diff --git a/src/vulkan/engine_vulkan.cc b/src/vulkan/engine_vulkan.cc index f8cfb396d..b8d0b903d 100644 --- a/src/vulkan/engine_vulkan.cc +++ b/src/vulkan/engine_vulkan.cc @@ -237,9 +237,29 @@ Result EngineVulkan::CreatePipeline(amber::Pipeline* pipeline) { buf_info.buffer->GetBufferType()))); } + // Bind OpenCL arguments if they are specified. + uint32_t descriptor_set = buf_info.descriptor_set; + uint32_t binding = buf_info.binding; + const auto& descriptor_map = pipeline->GetShaders()[0].GetDescriptorMap(); + if (!descriptor_map.empty()) { + const std::string& entry_point = + pipeline->GetShaders()[0].GetEntryPoint(); + auto iter = descriptor_map.find(entry_point); + if (iter != descriptor_map.end()) { + for (const auto& entry : iter->second) { + if (entry.arg_name == buf_info.arg_name || + entry.arg_ordinal == buf_info.arg_no) { + descriptor_set = entry.descriptor_set; + binding = entry.binding; + break; + } + } + } + } + auto cmd = MakeUnique(type, pipeline); - cmd->SetDescriptorSet(buf_info.descriptor_set); - cmd->SetBinding(buf_info.binding); + cmd->SetDescriptorSet(descriptor_set); + cmd->SetBinding(binding); cmd->SetBuffer(buf_info.buffer); r = info.vk_pipeline->AddDescriptor(cmd.get()); From 4101314fb6437183aa1845c70ae36db12a4cce9b Mon Sep 17 00:00:00 2001 From: Alan Baker Date: Fri, 19 Jul 2019 10:53:11 -0400 Subject: [PATCH 2/9] Improvements and fixes * Refactored binding updates to Pipeline function * Fixed how unused buffer elements are set to prevent mismatches * added some extra error handling * OPENCL-C kernels disallow caching * local memory args are allowed, but only supported via specialization constants * added compiler and pipeline tests for new functionality * added new end-to-end test --- src/clspv_helper.cc | 5 +++ src/executor.cc | 1 + src/pipeline.cc | 35 ++++++++++++++++ src/pipeline.h | 16 +++++--- src/pipeline_test.cc | 47 ++++++++++++++++++++++ src/shader_compiler.cc | 9 ++++- src/shader_compiler_test.cc | 47 ++++++++++++++++++++++ src/vulkan/engine_vulkan.cc | 24 +---------- tests/cases/opencl_bind_buffer.amber | 60 ++++++++++++++++++++++++++++ 9 files changed, 215 insertions(+), 29 deletions(-) create mode 100644 tests/cases/opencl_bind_buffer.amber diff --git a/src/clspv_helper.cc b/src/clspv_helper.cc index 058ab80a7..38343f92f 100644 --- a/src/clspv_helper.cc +++ b/src/clspv_helper.cc @@ -57,12 +57,17 @@ Result Compile(Pipeline::ShaderInfo* shader_info, descriptor_entry.kind = Pipeline::ShaderInfo::DescriptorMapEntry::Kind::PodUBO; break; + case clspv::ArgKind::Local: + // Local arguments are handled via specialization constants. + break; default: return Result("Unsupported kernel argument descriptor entry"); } if (entry.kernel_arg_data.arg_kind == clspv::ArgKind::Pod || entry.kernel_arg_data.arg_kind == clspv::ArgKind::PodUBO) { + if (entry.kernel_arg_data.pod_offset != 0) + return Result("Clustered PoD arguments are not currently supported"); descriptor_entry.pod_offset = entry.kernel_arg_data.pod_offset; descriptor_entry.pod_arg_size = entry.kernel_arg_data.pod_arg_size; } diff --git a/src/executor.cc b/src/executor.cc index 490109274..3bc0bdb20 100644 --- a/src/executor.cc +++ b/src/executor.cc @@ -60,6 +60,7 @@ Result Executor::Execute(Engine* engine, return r; for (auto& pipeline : script->GetPipelines()) { + pipeline->UpdateOpenCLBufferBindings(); r = engine->CreatePipeline(pipeline.get()); if (!r.IsSuccess()) return r; diff --git a/src/pipeline.cc b/src/pipeline.cc index 4a6e837ef..c8a9d322f 100644 --- a/src/pipeline.cc +++ b/src/pipeline.cc @@ -15,6 +15,7 @@ #include "src/pipeline.h" #include +#include #include #include "src/format_parser.h" @@ -343,6 +344,9 @@ void Pipeline::AddBuffer(Buffer* buf, const std::string& arg_name) { auto& info = buffers_.back(); info.arg_name = arg_name; + info.descriptor_set = std::numeric_limits::max(); + info.binding = std::numeric_limits::max(); + info.arg_no = std::numeric_limits::max(); } void Pipeline::AddBuffer(Buffer* buf, uint32_t arg_no) { @@ -358,6 +362,37 @@ void Pipeline::AddBuffer(Buffer* buf, uint32_t arg_no) { auto& info = buffers_.back(); info.arg_no = arg_no; + info.descriptor_set = std::numeric_limits::max(); + info.binding = std::numeric_limits::max(); +} + +void Pipeline::UpdateOpenCLBufferBindings() { + if (!IsCompute() || + GetShaders().empty() || + GetShaders()[0].GetShader()->GetFormat() != kShaderFormatOpenCLC) + return; + + const auto& shader_info = GetShaders()[0]; + const auto& descriptor_map = shader_info.GetDescriptorMap(); + if (descriptor_map.empty()) + return; + + const auto iter = descriptor_map.find(shader_info.GetEntryPoint()); + if (iter == descriptor_map.end()) + return; + + for (auto& info : buffers_) { + if (info.descriptor_set == std::numeric_limits::max() && + info.binding == std::numeric_limits::max()) { + for (const auto& entry : iter->second) { + if (entry.arg_name == info.arg_name || + entry.arg_ordinal == info.arg_no) { + info.descriptor_set = entry.descriptor_set; + info.binding = entry.binding; + } + } + } + } } } // namespace amber diff --git a/src/pipeline.h b/src/pipeline.h index 239defa97..fcd151398 100644 --- a/src/pipeline.h +++ b/src/pipeline.h @@ -70,7 +70,7 @@ class Pipeline { /// Descriptor information for an OpenCL-C shader. struct DescriptorMapEntry { - std::string arg_name; + std::string arg_name = ""; enum class Kind : int { SSBO, @@ -79,11 +79,11 @@ class Pipeline { PodUBO, } kind; - uint32_t descriptor_set; - uint32_t binding; - uint32_t arg_ordinal; - uint32_t pod_offset; - uint32_t pod_arg_size; + uint32_t descriptor_set = 0; + uint32_t binding = 0; + uint32_t arg_ordinal = 0; + uint32_t pod_offset = 0; + uint32_t pod_arg_size = 0; }; void AddDescriptorEntry(const std::string& kernel, @@ -203,6 +203,10 @@ class Pipeline { /// Returns information on all buffers in this pipeline. const std::vector& GetBuffers() const { return buffers_; } + /// Updates the descriptor set and binding info for the OpenCL-C kernel bound + /// to the pipeline. No effect for other shader formats. + void UpdateOpenCLBufferBindings(); + /// Returns the buffer which is currently bound to this pipeline at /// |descriptor_set| and |binding|. Buffer* GetBufferForBinding(uint32_t descriptor_set, uint32_t binding) const; diff --git a/src/pipeline_test.cc b/src/pipeline_test.cc index 8c61a0b53..a5cc56bfb 100644 --- a/src/pipeline_test.cc +++ b/src/pipeline_test.cc @@ -396,4 +396,51 @@ TEST_F(PipelineTest, Clone) { EXPECT_EQ(2U, bufs[1].binding); } +#if AMBER_ENABLE_CLSPV +TEST_F(PipelineTest, ClspvUpdateBindings) { + Pipeline p(PipelineType::kCompute); + p.SetName("my_pipeline"); + + Shader cs(kShaderTypeCompute); + cs.SetFormat(kShaderFormatOpenCLC); + p.AddShader(&cs, kShaderTypeCompute); + p.SetShaderEntryPoint(&cs, "my_main"); + + Pipeline::ShaderInfo::DescriptorMapEntry entry1; + entry1.kind = Pipeline::ShaderInfo::DescriptorMapEntry::Kind::SSBO; + entry1.descriptor_set = 4; + entry1.binding = 5; + entry1.arg_name = "arg_a"; + entry1.arg_ordinal = 0; + p.GetShaders()[0].AddDescriptorEntry("my_main", std::move(entry1)); + + Pipeline::ShaderInfo::DescriptorMapEntry entry2; + entry2.kind = Pipeline::ShaderInfo::DescriptorMapEntry::Kind::SSBO; + entry2.descriptor_set = 3; + entry2.binding = 1; + entry2.arg_name = "arg_b"; + entry2.arg_ordinal = 1; + p.GetShaders()[0].AddDescriptorEntry("my_main", std::move(entry2)); + + auto a_buf = MakeUnique(BufferType::kStorage); + a_buf->SetName("buf1"); + p.AddBuffer(a_buf.get(), "arg_a"); + + auto b_buf = MakeUnique(BufferType::kStorage); + b_buf->SetName("buf2"); + p.AddBuffer(b_buf.get(), 1); + + p.UpdateOpenCLBufferBindings(); + + auto& bufs = p.GetBuffers(); + ASSERT_EQ(2U, bufs.size()); + EXPECT_EQ("buf1", bufs[0].buffer->GetName()); + EXPECT_EQ(4U, bufs[0].descriptor_set); + EXPECT_EQ(5U, bufs[0].binding); + EXPECT_EQ("buf2", bufs[1].buffer->GetName()); + EXPECT_EQ(3U, bufs[1].descriptor_set); + EXPECT_EQ(1U, bufs[1].binding); +} +#endif // AMBER_ENABLE_CLSPV + } // namespace amber diff --git a/src/shader_compiler.cc b/src/shader_compiler.cc index 916ad9aa7..f2dc41bae 100644 --- a/src/shader_compiler.cc +++ b/src/shader_compiler.cc @@ -55,8 +55,15 @@ std::pair> ShaderCompiler::Compile( const ShaderMap& shader_map) const { const auto shader = shader_info->GetShader(); auto it = shader_map.find(shader->GetName()); - if (it != shader_map.end()) + if (it != shader_map.end()) { +#if AMBER_ENABLE_CLSPV + if (shader->GetFormat() == kShaderFormatOpenCLC) { + return {Result("OPENCL-C shaders do not support pre-compiled shaders"), + {}}; + } +#endif // AMBER_ENABLE_CLSPV return {{}, it->second}; + } #if AMBER_ENABLE_SPIRV_TOOLS std::string spv_errors; diff --git a/src/shader_compiler_test.cc b/src/shader_compiler_test.cc index 06dae021b..e33d6ddf3 100644 --- a/src/shader_compiler_test.cc +++ b/src/shader_compiler_test.cc @@ -227,6 +227,53 @@ TEST_F(ShaderCompilerTest, ReturnsCachedShader) { } } +#if AMBER_ENABLE_CLSPV +TEST_F(ShaderCompilerTest, ClspvCompile) { + Shader shader(kShaderTypeCompute); + shader.SetName("TestShader"); + shader.SetFormat(kShaderFormatOpenCLC); + shader.SetData(R"( +kernel void TestShader(global int* in, global int* out) { + *out = *in; +} + )"); + + ShaderCompiler sc; + Result r; + std::vector binary; + Pipeline::ShaderInfo shader_info(&shader, kShaderTypeCompute); + std::tie(r, binary) = sc.Compile(&shader_info, ShaderMap()); + ASSERT_TRUE(r.IsSuccess()); + EXPECT_FALSE(binary.empty()); + EXPECT_EQ(0x07230203, binary[0]); // Verify SPIR-V header present. +} + +TEST_F(ShaderCompilerTest, ClspvDisallowCaching) { + Shader shader(kShaderTypeCompute); + std::string name = "TestShader"; + shader.SetName(name); + shader.SetFormat(kShaderFormatOpenCLC); + shader.SetData(R"( +kernel void TestShader(global int* in, global int* out) { + *out = *in; +} + )"); + + std::vector src_bytes = {1, 2, 3, 4, 5}; + + ShaderMap map; + map[name] = src_bytes; + + ShaderCompiler sc; + Result r; + std::vector binary; + Pipeline::ShaderInfo shader_info(&shader, kShaderTypeCompute); + std::tie(r, binary) = sc.Compile(&shader_info, map); + ASSERT_FALSE(r.IsSuccess()); + EXPECT_TRUE(binary.empty()); +} +#endif // AMBER_ENABLE_CLSPV + struct ParseSpvEnvCase { std::string env_str; bool ok; diff --git a/src/vulkan/engine_vulkan.cc b/src/vulkan/engine_vulkan.cc index b8d0b903d..f8cfb396d 100644 --- a/src/vulkan/engine_vulkan.cc +++ b/src/vulkan/engine_vulkan.cc @@ -237,29 +237,9 @@ Result EngineVulkan::CreatePipeline(amber::Pipeline* pipeline) { buf_info.buffer->GetBufferType()))); } - // Bind OpenCL arguments if they are specified. - uint32_t descriptor_set = buf_info.descriptor_set; - uint32_t binding = buf_info.binding; - const auto& descriptor_map = pipeline->GetShaders()[0].GetDescriptorMap(); - if (!descriptor_map.empty()) { - const std::string& entry_point = - pipeline->GetShaders()[0].GetEntryPoint(); - auto iter = descriptor_map.find(entry_point); - if (iter != descriptor_map.end()) { - for (const auto& entry : iter->second) { - if (entry.arg_name == buf_info.arg_name || - entry.arg_ordinal == buf_info.arg_no) { - descriptor_set = entry.descriptor_set; - binding = entry.binding; - break; - } - } - } - } - auto cmd = MakeUnique(type, pipeline); - cmd->SetDescriptorSet(descriptor_set); - cmd->SetBinding(binding); + cmd->SetDescriptorSet(buf_info.descriptor_set); + cmd->SetBinding(buf_info.binding); cmd->SetBuffer(buf_info.buffer); r = info.vk_pipeline->AddDescriptor(cmd.get()); diff --git a/tests/cases/opencl_bind_buffer.amber b/tests/cases/opencl_bind_buffer.amber new file mode 100644 index 000000000..a710a6b75 --- /dev/null +++ b/tests/cases/opencl_bind_buffer.amber @@ -0,0 +1,60 @@ +#!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 my_shader OPENCL-C +kernel void foo(global int* in, global int* out) { + unsigned int local_size_x = get_local_size(0); + unsigned int local_size_y = get_local_size(1); + unsigned int local_size_z = get_local_size(2); + unsigned int local_id_x = get_local_id(0); + unsigned int local_id_y = get_local_id(1); + unsigned int local_id_z = get_local_id(2); + unsigned int group_id_x = get_group_id(0); + unsigned int group_id_y = get_group_id(1); + unsigned int group_id_z = get_group_id(2); + unsigned int global_id_x = get_global_id(0); + unsigned int global_id_y = get_global_id(1); + unsigned int global_id_z = get_global_id(2); + unsigned int wgs_x = get_num_groups(0); + unsigned int wgs_y = get_num_groups(1); + unsigned int wgs_z = get_num_groups(2); + + unsigned int in_wg_id = (local_id_z * local_size_x * local_size_y) + + (local_id_y * local_size_x) + + local_id_x; + unsigned int prev_ids = (local_size_x * local_size_y * local_size_z) * + (group_id_z * wgs_y * wgs_x + group_id_y * wgs_x + group_id_x); + unsigned int linear_id = in_wg_id + prev_ids; + out[linear_id] = in[linear_id]; +} +END + +BUFFER in_buf DATA_TYPE uint32 SIZE 64 SERIES_FROM 1 INC_BY 1 +BUFFER out_buf DATA_TYPE uint32 SIZE 64 FILL 0 + +PIPELINE compute my_pipeline + ATTACH my_shader ENTRY_POINT foo \ + SPECIALIZE 0 AS uint32 2 \ + SPECIALIZE 1 AS uint32 2 \ + SPECIALIZE 2 AS uint32 2 + + BIND BUFFER in_buf AS storage KERNEL ARG in + BIND BUFFER out_buf AS storage KERNEL ARGNO 1 +END + +RUN my_pipeline 2 2 2 + +EXPECT out_buf EQ_BUFFER in_buf + From 96e7292b3cbdb164c80852705fc1e061a49ec098 Mon Sep 17 00:00:00 2001 From: Alan Baker Date: Fri, 19 Jul 2019 10:58:06 -0400 Subject: [PATCH 3/9] Add documentation --- docs/amber_script.md | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/docs/amber_script.md b/docs/amber_script.md index cbe5c737f..df07e9bb4 100644 --- a/docs/amber_script.md +++ b/docs/amber_script.md @@ -254,6 +254,12 @@ attachment content, depth/stencil content, uniform buffers, etc. # Bind the sampler at the given descriptor set and binding. BIND SAMPLER {sampler_name} DESCRIPTOR_SET _id_ BINDING _id_ + + # Bind OpenCL argument buffer by name. + BIND BUFFER {buffer_name} AS {buffer_type} KERNEL ARG _name_ + + # Bind OpenCL argument buffer by argument ordinal. + BIND BUFFER {buffer_name} AS {buffer_type} KERNEL ARGNO _number_ ``` ```groovy From a1fa85bccd8a2bfcf35ee26dad2e97b66ff67c97 Mon Sep 17 00:00:00 2001 From: Alan Baker Date: Fri, 19 Jul 2019 11:12:12 -0400 Subject: [PATCH 4/9] formatting --- src/clspv_helper.cc | 2 ++ src/pipeline.h | 3 ++- src/pipeline_test.cc | 2 +- 3 files changed, 5 insertions(+), 2 deletions(-) diff --git a/src/clspv_helper.cc b/src/clspv_helper.cc index 38343f92f..ce627e685 100644 --- a/src/clspv_helper.cc +++ b/src/clspv_helper.cc @@ -14,6 +14,8 @@ #include "src/clspv_helper.h" +#include + #include "clspv/ArgKind.h" #include "clspv/Compiler.h" diff --git a/src/pipeline.h b/src/pipeline.h index fcd151398..9dfe743d9 100644 --- a/src/pipeline.h +++ b/src/pipeline.h @@ -102,7 +102,8 @@ class Pipeline { std::string entry_point_; std::vector data_; std::map specialization_; - std::unordered_map> descriptor_map_; + std::unordered_map> + descriptor_map_; }; /// Information on a buffer attached to the pipeline. diff --git a/src/pipeline_test.cc b/src/pipeline_test.cc index a5cc56bfb..3c1180bf1 100644 --- a/src/pipeline_test.cc +++ b/src/pipeline_test.cc @@ -441,6 +441,6 @@ TEST_F(PipelineTest, ClspvUpdateBindings) { EXPECT_EQ(3U, bufs[1].descriptor_set); EXPECT_EQ(1U, bufs[1].binding); } -#endif // AMBER_ENABLE_CLSPV +#endif // AMBER_ENABLE_CLSPV } // namespace amber From 82ecc451182af585c1499688e8e41785ab4d4789 Mon Sep 17 00:00:00 2001 From: Alan Baker Date: Fri, 19 Jul 2019 11:18:45 -0400 Subject: [PATCH 5/9] formatting --- src/pipeline.cc | 3 +-- src/shader_compiler.h | 2 +- 2 files changed, 2 insertions(+), 3 deletions(-) diff --git a/src/pipeline.cc b/src/pipeline.cc index c8a9d322f..30d7a4696 100644 --- a/src/pipeline.cc +++ b/src/pipeline.cc @@ -367,8 +367,7 @@ void Pipeline::AddBuffer(Buffer* buf, uint32_t arg_no) { } void Pipeline::UpdateOpenCLBufferBindings() { - if (!IsCompute() || - GetShaders().empty() || + if (!IsCompute() || GetShaders().empty() || GetShaders()[0].GetShader()->GetFormat() != kShaderFormatOpenCLC) return; diff --git a/src/shader_compiler.h b/src/shader_compiler.h index 02769b46c..a6dc8f876 100644 --- a/src/shader_compiler.h +++ b/src/shader_compiler.h @@ -21,8 +21,8 @@ #include "amber/amber.h" #include "amber/result.h" -#include "src/shader.h" #include "src/pipeline.h" +#include "src/shader.h" namespace amber { From effbf7a39fbe00873eca0f12670b61a4ff98efb2 Mon Sep 17 00:00:00 2001 From: Alan Baker Date: Fri, 19 Jul 2019 11:54:22 -0400 Subject: [PATCH 6/9] More explicit opencl binding keywords * changes to ARG_NAME and ARG_NUMBER * update collateral --- docs/amber_script.md | 4 ++-- src/amberscript/parser.cc | 6 +++--- src/amberscript/parser_bind_test.cc | 16 ++++++++-------- tests/cases/opencl_bind_buffer.amber | 4 ++-- 4 files changed, 15 insertions(+), 15 deletions(-) diff --git a/docs/amber_script.md b/docs/amber_script.md index df07e9bb4..e10ebba66 100644 --- a/docs/amber_script.md +++ b/docs/amber_script.md @@ -256,10 +256,10 @@ attachment content, depth/stencil content, uniform buffers, etc. BIND SAMPLER {sampler_name} DESCRIPTOR_SET _id_ BINDING _id_ # Bind OpenCL argument buffer by name. - BIND BUFFER {buffer_name} AS {buffer_type} KERNEL ARG _name_ + BIND BUFFER {buffer_name} AS {buffer_type} KERNEL ARG_NAME _name_ # Bind OpenCL argument buffer by argument ordinal. - BIND BUFFER {buffer_name} AS {buffer_type} KERNEL ARGNO _number_ + BIND BUFFER {buffer_name} AS {buffer_type} KERNEL ARG_NUMBER _number_ ``` ```groovy diff --git a/src/amberscript/parser.cc b/src/amberscript/parser.cc index 04d2a5595..385d6032c 100644 --- a/src/amberscript/parser.cc +++ b/src/amberscript/parser.cc @@ -712,20 +712,20 @@ Result Parser::ParsePipelineBind(Pipeline* pipeline) { if (!token->IsString()) return Result("missing kernel arg identifier"); - if (token->AsString() == "ARG") { + if (token->AsString() == "ARG_NAME") { token = tokenizer_->NextToken(); if (!token->IsString()) return Result("expected argument identifier"); pipeline->AddBuffer(buffer, token->AsString()); - } else if (token->AsString() == "ARGNO") { + } else if (token->AsString() == "ARG_NUMBER") { token = tokenizer_->NextToken(); if (!token->IsInteger()) return Result("expected argument identifier number"); pipeline->AddBuffer(buffer, token->AsUint32()); } else { - return Result("missing ARG or ARGNO keyword"); + return Result("missing ARG_NAME or ARG_NUMBER keyword"); } } } diff --git a/src/amberscript/parser_bind_test.cc b/src/amberscript/parser_bind_test.cc index ee8d9ac0b..d533dfff0 100644 --- a/src/amberscript/parser_bind_test.cc +++ b/src/amberscript/parser_bind_test.cc @@ -1158,7 +1158,7 @@ BUFFER my_buf DATA_TYPE uint32 DATA 1 END PIPELINE compute my_pipeline ATTACH my_shader - BIND BUFFER my_buf AS storage KERNEL ARG arg + BIND BUFFER my_buf AS storage KERNEL ARG_NAME arg END)"; Parser parser; @@ -1175,7 +1175,7 @@ BUFFER my_buf DATA_TYPE uint32 DATA 1 END PIPELINE compute my_pipeline ATTACH my_shader - BIND BUFFER my_buf AS storage KERNEL ARGNO 0 + BIND BUFFER my_buf AS storage KERNEL ARG_NUMBER 0 END)"; Parser parser; @@ -1192,7 +1192,7 @@ BUFFER my_buf DATA_TYPE uint32 DATA 1 END PIPELINE compute my_pipeline ATTACH my_shader - BIND BUFFER my_buf AS storage ARG arg + BIND BUFFER my_buf AS storage ARG_NAME arg END)"; Parser parser; @@ -1216,7 +1216,7 @@ END)"; Parser parser; Result r = parser.Parse(in); ASSERT_FALSE(r.IsSuccess()); - EXPECT_EQ("9: missing ARG or ARGNO keyword", r.Error()); + EXPECT_EQ("9: missing ARG_NAME or ARG_NUMBER keyword", r.Error()); } TEST_F(AmberScriptParserTest, BindBufferOpenCLMissingArgName) { @@ -1228,7 +1228,7 @@ BUFFER my_buf DATA_TYPE uint32 DATA 1 END PIPELINE compute my_pipeline ATTACH my_shader - BIND BUFFER my_buf AS storage KERNEL ARG + BIND BUFFER my_buf AS storage KERNEL ARG_NAME END)"; Parser parser; @@ -1246,7 +1246,7 @@ BUFFER my_buf DATA_TYPE uint32 DATA 1 END PIPELINE compute my_pipeline ATTACH my_shader - BIND BUFFER my_buf AS storage KERNEL ARGNO + BIND BUFFER my_buf AS storage KERNEL ARG_NUMBER END)"; Parser parser; @@ -1264,7 +1264,7 @@ BUFFER my_buf DATA_TYPE uint32 DATA 1 END PIPELINE compute my_pipeline ATTACH my_shader - BIND BUFFER my_buf AS storage KERNEL ARG 0 + BIND BUFFER my_buf AS storage KERNEL ARG_NAME 0 END)"; Parser parser; @@ -1282,7 +1282,7 @@ BUFFER my_buf DATA_TYPE uint32 DATA 1 END PIPELINE compute my_pipeline ATTACH my_shader - BIND BUFFER my_buf AS storage KERNEL ARGNO in + BIND BUFFER my_buf AS storage KERNEL ARG_NUMBER in END)"; Parser parser; diff --git a/tests/cases/opencl_bind_buffer.amber b/tests/cases/opencl_bind_buffer.amber index a710a6b75..d25837ea7 100644 --- a/tests/cases/opencl_bind_buffer.amber +++ b/tests/cases/opencl_bind_buffer.amber @@ -50,8 +50,8 @@ PIPELINE compute my_pipeline SPECIALIZE 1 AS uint32 2 \ SPECIALIZE 2 AS uint32 2 - BIND BUFFER in_buf AS storage KERNEL ARG in - BIND BUFFER out_buf AS storage KERNEL ARGNO 1 + BIND BUFFER in_buf AS storage KERNEL ARG_NAME in + BIND BUFFER out_buf AS storage KERNEL ARG_NUMBER 1 END RUN my_pipeline 2 2 2 From 851b825066d928f0ea932e0db3f155f26c2b0982 Mon Sep 17 00:00:00 2001 From: Alan Baker Date: Fri, 19 Jul 2019 13:07:16 -0400 Subject: [PATCH 7/9] Improved parsing and error messages --- src/amberscript/parser.cc | 12 ++++-------- src/amberscript/parser_bind_test.cc | 8 ++++---- 2 files changed, 8 insertions(+), 12 deletions(-) diff --git a/src/amberscript/parser.cc b/src/amberscript/parser.cc index 385d6032c..d73de1f46 100644 --- a/src/amberscript/parser.cc +++ b/src/amberscript/parser.cc @@ -701,13 +701,7 @@ Result Parser::ParsePipelineBind(Pipeline* pipeline) { if (!token->IsInteger()) return Result("invalid value for BINDING in BIND command"); pipeline->AddBuffer(buffer, descriptor_set, token->AsUint32()); - } else { - if (!token->IsString()) - return Result("missing DESCRIPTOR_SET for BIND command"); - - if (token->AsString() != "KERNEL") - return Result("missing DESCRIPTOR_SET for BIND command"); - + } else if (token->IsString() && token->AsString() == "KERNEL") { token = tokenizer_->NextToken(); if (!token->IsString()) return Result("missing kernel arg identifier"); @@ -721,12 +715,14 @@ Result Parser::ParsePipelineBind(Pipeline* pipeline) { } else if (token->AsString() == "ARG_NUMBER") { token = tokenizer_->NextToken(); if (!token->IsInteger()) - return Result("expected argument identifier number"); + return Result("expected argument number"); pipeline->AddBuffer(buffer, token->AsUint32()); } else { return Result("missing ARG_NAME or ARG_NUMBER keyword"); } + } else { + return Result("missing DESCRIPTOR_SET or KERNEL for BIND command"); } } diff --git a/src/amberscript/parser_bind_test.cc b/src/amberscript/parser_bind_test.cc index d533dfff0..bf4800b1c 100644 --- a/src/amberscript/parser_bind_test.cc +++ b/src/amberscript/parser_bind_test.cc @@ -965,7 +965,7 @@ END)"; Parser parser; Result r = parser.Parse(in); ASSERT_FALSE(r.IsSuccess()); - EXPECT_EQ("12: missing DESCRIPTOR_SET for BIND command", r.Error()); + EXPECT_EQ("12: missing DESCRIPTOR_SET or KERNEL for BIND command", r.Error()); } TEST_F(AmberScriptParserTest, BindingBufferExtraParams) { @@ -1198,7 +1198,7 @@ END)"; Parser parser; Result r = parser.Parse(in); ASSERT_FALSE(r.IsSuccess()); - EXPECT_EQ("9: missing DESCRIPTOR_SET for BIND command", r.Error()); + EXPECT_EQ("9: missing DESCRIPTOR_SET or KERNEL for BIND command", r.Error()); } TEST_F(AmberScriptParserTest, BindBufferOpenCLMissingArg) { @@ -1252,7 +1252,7 @@ END)"; Parser parser; Result r = parser.Parse(in); ASSERT_FALSE(r.IsSuccess()); - EXPECT_EQ("10: expected argument identifier number", r.Error()); + EXPECT_EQ("10: expected argument number", r.Error()); } TEST_F(AmberScriptParserTest, BindBufferOpenCLArgNameNotString) { @@ -1288,7 +1288,7 @@ END)"; Parser parser; Result r = parser.Parse(in); ASSERT_FALSE(r.IsSuccess()); - EXPECT_EQ("9: expected argument identifier number", r.Error()); + EXPECT_EQ("9: expected argument number", r.Error()); } } // namespace amberscript From bbba07027af4baf1f1abf388efbeccea9c0a2188 Mon Sep 17 00:00:00 2001 From: Alan Baker Date: Fri, 19 Jul 2019 14:21:56 -0400 Subject: [PATCH 8/9] Revamp syntax and parsing * Buffer storage class specification is now optional for OpenCL-C shaders * updated tests and docs * Added a storage class consistency check when binding are updated * added a test for bad consistency * Made enum follow google style guide --- docs/amber_script.md | 14 +++-- src/amberscript/parser.cc | 83 +++++++++++++++------------- src/amberscript/parser_bind_test.cc | 44 +++++++++++++-- src/clspv_helper.cc | 4 +- src/executor.cc | 4 +- src/pipeline.cc | 44 +++++++++++++-- src/pipeline.h | 7 ++- src/pipeline_test.cc | 39 +++++++++++++ tests/cases/opencl_bind_buffer.amber | 4 +- 9 files changed, 183 insertions(+), 60 deletions(-) diff --git a/docs/amber_script.md b/docs/amber_script.md index e10ebba66..e42e7dad6 100644 --- a/docs/amber_script.md +++ b/docs/amber_script.md @@ -255,11 +255,15 @@ attachment content, depth/stencil content, uniform buffers, etc. # Bind the sampler at the given descriptor set and binding. BIND SAMPLER {sampler_name} DESCRIPTOR_SET _id_ BINDING _id_ - # Bind OpenCL argument buffer by name. - BIND BUFFER {buffer_name} AS {buffer_type} KERNEL ARG_NAME _name_ - - # Bind OpenCL argument buffer by argument ordinal. - BIND BUFFER {buffer_name} AS {buffer_type} KERNEL ARG_NUMBER _number_ + # Bind OpenCL argument buffer by name. Specifying the buffer type is optional. + # Amber will set the type as appropriate for the argument buffer. All uses + # of the buffer must be across all pipelines. + BIND BUFFER {buffer_name} [AS {buffer_type}] KERNEL ARG_NAME _name_ + + # Bind OpenCL argument buffer by argument ordinal. Specifying the buffer type + # is optional. Amber will set the type as appropriate for the argument + # buffer. All uses of the buffer must be across all pipelines. + BIND BUFFER {buffer_name} [AS {buffer_type}] KERNEL ARG_NUMBER _number_ ``` ```groovy diff --git a/src/amberscript/parser.cc b/src/amberscript/parser.cc index d73de1f46..1abe3a4f1 100644 --- a/src/amberscript/parser.cc +++ b/src/amberscript/parser.cc @@ -644,50 +644,57 @@ Result Parser::ParsePipelineBind(Pipeline* pipeline) { return Result("unknown buffer: " + token->AsString()); token = tokenizer_->NextToken(); - if (!token->IsString() || token->AsString() != "AS") - return Result("BUFFER command missing AS keyword"); - - token = tokenizer_->NextToken(); - if (!token->IsString()) - return Result("invalid token for BUFFER type"); - - if (token->AsString() == "color") { + if (token->IsString() && token->AsString() == "AS") { token = tokenizer_->NextToken(); - if (!token->IsString() || token->AsString() != "LOCATION") - return Result("BIND missing LOCATION"); + if (!token->IsString()) + return Result("invalid token for BUFFER type"); - token = tokenizer_->NextToken(); - if (!token->IsInteger()) - return Result("invalid value for BIND LOCATION"); + if (token->AsString() == "color") { + token = tokenizer_->NextToken(); + if (!token->IsString() || token->AsString() != "LOCATION") + return Result("BIND missing LOCATION"); - buffer->SetBufferType(BufferType::kColor); + token = tokenizer_->NextToken(); + if (!token->IsInteger()) + return Result("invalid value for BIND LOCATION"); - Result r = pipeline->AddColorAttachment(buffer, token->AsUint32()); - if (!r.IsSuccess()) - return r; - } else if (token->AsString() == "depth_stencil") { - buffer->SetBufferType(BufferType::kDepth); - Result r = pipeline->SetDepthBuffer(buffer); - if (!r.IsSuccess()) - return r; - } else if (token->AsString() == "push_constant") { - buffer->SetBufferType(BufferType::kPushConstant); - Result r = pipeline->SetPushConstantBuffer(buffer); - if (!r.IsSuccess()) - return r; - } else { - BufferType type = BufferType::kColor; - Result r = ToBufferType(token->AsString(), &type); - if (!r.IsSuccess()) - return r; + buffer->SetBufferType(BufferType::kColor); - if (buffer->GetBufferType() == BufferType::kUnknown) - buffer->SetBufferType(type); - else if (buffer->GetBufferType() != type) - return Result("buffer type does not match intended usage"); + Result r = pipeline->AddColorAttachment(buffer, token->AsUint32()); + if (!r.IsSuccess()) + return r; + } else if (token->AsString() == "depth_stencil") { + buffer->SetBufferType(BufferType::kDepth); + Result r = pipeline->SetDepthBuffer(buffer); + if (!r.IsSuccess()) + return r; + } else if (token->AsString() == "push_constant") { + buffer->SetBufferType(BufferType::kPushConstant); + Result r = pipeline->SetPushConstantBuffer(buffer); + if (!r.IsSuccess()) + return r; + } else { + BufferType type = BufferType::kColor; + Result r = ToBufferType(token->AsString(), &type); + if (!r.IsSuccess()) + return r; - token = tokenizer_->NextToken(); - if (token->IsString() && token->AsString() == "DESCRIPTOR_SET") { + if (buffer->GetBufferType() == BufferType::kUnknown) + buffer->SetBufferType(type); + else if (buffer->GetBufferType() != type) + return Result("buffer type does not match intended usage"); + } + } + + if (buffer->GetBufferType() == BufferType::kUnknown || + buffer->GetBufferType() == BufferType::kStorage || + buffer->GetBufferType() == BufferType::kUniform) { + // If AS was parsed above consume the next token. + if (buffer->GetBufferType() != BufferType::kUnknown) + token = tokenizer_->NextToken(); + // DESCRIPTOR_SET requires a buffer type to have been specified. + if (buffer->GetBufferType() != BufferType::kUnknown && token->IsString() && + token->AsString() == "DESCRIPTOR_SET") { token = tokenizer_->NextToken(); if (!token->IsInteger()) return Result("invalid value for DESCRIPTOR_SET in BIND command"); diff --git a/src/amberscript/parser_bind_test.cc b/src/amberscript/parser_bind_test.cc index bf4800b1c..5e6d87c29 100644 --- a/src/amberscript/parser_bind_test.cc +++ b/src/amberscript/parser_bind_test.cc @@ -1166,7 +1166,7 @@ END)"; ASSERT_TRUE(r.IsSuccess()); } -TEST_F(AmberScriptParserTest, BindBufferOpenCLArgNo) { +TEST_F(AmberScriptParserTest, BindBufferOpenCLArgNumber) { std::string in = R"( SHADER compute my_shader OPENCL-C #shader @@ -1183,6 +1183,40 @@ END)"; ASSERT_TRUE(r.IsSuccess()); } +TEST_F(AmberScriptParserTest, BindBufferOpenCLArgNameTypeless) { + std::string in = R"( +SHADER compute my_shader OPENCL-C +#shader +END +BUFFER my_buf DATA_TYPE uint32 DATA 1 END + +PIPELINE compute my_pipeline + ATTACH my_shader + BIND BUFFER my_buf KERNEL ARG_NAME arg +END)"; + + Parser parser; + Result r = parser.Parse(in); + ASSERT_TRUE(r.IsSuccess()); +} + +TEST_F(AmberScriptParserTest, BindBufferOpenCLArgNumberTypeless) { + std::string in = R"( +SHADER compute my_shader OPENCL-C +#shader +END +BUFFER my_buf DATA_TYPE uint32 DATA 1 END + +PIPELINE compute my_pipeline + ATTACH my_shader + BIND BUFFER my_buf KERNEL ARG_NUMBER 0 +END)"; + + Parser parser; + Result r = parser.Parse(in); + ASSERT_TRUE(r.IsSuccess()); +} + TEST_F(AmberScriptParserTest, BindBufferOpenCLMissingKernel) { std::string in = R"( SHADER compute my_shader OPENCL-C @@ -1228,7 +1262,7 @@ BUFFER my_buf DATA_TYPE uint32 DATA 1 END PIPELINE compute my_pipeline ATTACH my_shader - BIND BUFFER my_buf AS storage KERNEL ARG_NAME + BIND BUFFER my_buf KERNEL ARG_NAME END)"; Parser parser; @@ -1237,7 +1271,7 @@ END)"; EXPECT_EQ("10: expected argument identifier", r.Error()); } -TEST_F(AmberScriptParserTest, BindBufferOpenCLMissingArgNo) { +TEST_F(AmberScriptParserTest, BindBufferOpenCLMissingArgNumber) { std::string in = R"( SHADER compute my_shader OPENCL-C #shader @@ -1273,7 +1307,7 @@ END)"; EXPECT_EQ("9: expected argument identifier", r.Error()); } -TEST_F(AmberScriptParserTest, BindBufferOpenCLArgNoNotInteger) { +TEST_F(AmberScriptParserTest, BindBufferOpenCLArgNumberNotInteger) { std::string in = R"( SHADER compute my_shader OPENCL-C #shader @@ -1282,7 +1316,7 @@ BUFFER my_buf DATA_TYPE uint32 DATA 1 END PIPELINE compute my_pipeline ATTACH my_shader - BIND BUFFER my_buf AS storage KERNEL ARG_NUMBER in + BIND BUFFER my_buf KERNEL ARG_NUMBER in END)"; Parser parser; diff --git a/src/clspv_helper.cc b/src/clspv_helper.cc index ce627e685..e0d853ea6 100644 --- a/src/clspv_helper.cc +++ b/src/clspv_helper.cc @@ -53,11 +53,11 @@ Result Compile(Pipeline::ShaderInfo* shader_info, break; case clspv::ArgKind::Pod: descriptor_entry.kind = - Pipeline::ShaderInfo::DescriptorMapEntry::Kind::Pod; + Pipeline::ShaderInfo::DescriptorMapEntry::Kind::POD; break; case clspv::ArgKind::PodUBO: descriptor_entry.kind = - Pipeline::ShaderInfo::DescriptorMapEntry::Kind::PodUBO; + Pipeline::ShaderInfo::DescriptorMapEntry::Kind::POD_UBO; break; case clspv::ArgKind::Local: // Local arguments are handled via specialization constants. diff --git a/src/executor.cc b/src/executor.cc index 3bc0bdb20..3bbec7a8d 100644 --- a/src/executor.cc +++ b/src/executor.cc @@ -60,7 +60,9 @@ Result Executor::Execute(Engine* engine, return r; for (auto& pipeline : script->GetPipelines()) { - pipeline->UpdateOpenCLBufferBindings(); + r = pipeline->UpdateOpenCLBufferBindings(); + if (!r.IsSuccess()) + return r; r = engine->CreatePipeline(pipeline.get()); if (!r.IsSuccess()) return r; diff --git a/src/pipeline.cc b/src/pipeline.cc index 30d7a4696..8f214f0fa 100644 --- a/src/pipeline.cc +++ b/src/pipeline.cc @@ -366,19 +366,19 @@ void Pipeline::AddBuffer(Buffer* buf, uint32_t arg_no) { info.binding = std::numeric_limits::max(); } -void Pipeline::UpdateOpenCLBufferBindings() { +Result Pipeline::UpdateOpenCLBufferBindings() { if (!IsCompute() || GetShaders().empty() || GetShaders()[0].GetShader()->GetFormat() != kShaderFormatOpenCLC) - return; + return {}; const auto& shader_info = GetShaders()[0]; const auto& descriptor_map = shader_info.GetDescriptorMap(); if (descriptor_map.empty()) - return; + return {}; const auto iter = descriptor_map.find(shader_info.GetEntryPoint()); if (iter == descriptor_map.end()) - return; + return {}; for (auto& info : buffers_) { if (info.descriptor_set == std::numeric_limits::max() && @@ -386,12 +386,48 @@ void Pipeline::UpdateOpenCLBufferBindings() { for (const auto& entry : iter->second) { if (entry.arg_name == info.arg_name || entry.arg_ordinal == info.arg_no) { + // Buffer storage class consistency checks. + if (info.buffer->GetBufferType() == BufferType::kUnknown) { + // Set the appropriate buffer type. + switch (entry.kind) { + case Pipeline::ShaderInfo::DescriptorMapEntry::Kind::UBO: + case Pipeline::ShaderInfo::DescriptorMapEntry::Kind::POD_UBO: + info.buffer->SetBufferType(BufferType::kUniform); + break; + case Pipeline::ShaderInfo::DescriptorMapEntry::Kind::SSBO: + case Pipeline::ShaderInfo::DescriptorMapEntry::Kind::POD: + info.buffer->SetBufferType(BufferType::kStorage); + break; + default: + return Result("Unhandled buffer type for OPENCL-C shader"); + } + } else if (info.buffer->GetBufferType() == BufferType::kUniform) { + if (entry.kind != + Pipeline::ShaderInfo::DescriptorMapEntry::Kind::UBO && + entry.kind != + Pipeline::ShaderInfo::DescriptorMapEntry::Kind::POD_UBO) { + return Result("Buffer " + info.buffer->GetName() + + " must be an uniform binding"); + } + } else if (info.buffer->GetBufferType() == BufferType::kStorage) { + if (entry.kind != + Pipeline::ShaderInfo::DescriptorMapEntry::Kind::SSBO && + entry.kind != + Pipeline::ShaderInfo::DescriptorMapEntry::Kind::POD) { + return Result("Buffer " + info.buffer->GetName() + + " must be a storage binding"); + } + } else { + return Result("Unhandled buffer type for OPENCL-C shader"); + } info.descriptor_set = entry.descriptor_set; info.binding = entry.binding; } } } } + + return {}; } } // namespace amber diff --git a/src/pipeline.h b/src/pipeline.h index 9dfe743d9..8859888c7 100644 --- a/src/pipeline.h +++ b/src/pipeline.h @@ -73,10 +73,11 @@ class Pipeline { std::string arg_name = ""; enum class Kind : int { + UNKNOWN, SSBO, UBO, - Pod, - PodUBO, + POD, + POD_UBO, } kind; uint32_t descriptor_set = 0; @@ -206,7 +207,7 @@ class Pipeline { /// Updates the descriptor set and binding info for the OpenCL-C kernel bound /// to the pipeline. No effect for other shader formats. - void UpdateOpenCLBufferBindings(); + Result UpdateOpenCLBufferBindings(); /// Returns the buffer which is currently bound to this pipeline at /// |descriptor_set| and |binding|. diff --git a/src/pipeline_test.cc b/src/pipeline_test.cc index 3c1180bf1..8e0ddccc2 100644 --- a/src/pipeline_test.cc +++ b/src/pipeline_test.cc @@ -441,6 +441,45 @@ TEST_F(PipelineTest, ClspvUpdateBindings) { EXPECT_EQ(3U, bufs[1].descriptor_set); EXPECT_EQ(1U, bufs[1].binding); } + +TEST_F(PipelineTest, ClspvUpdateBindingTypeMismatch) { + Pipeline p(PipelineType::kCompute); + p.SetName("my_pipeline"); + + Shader cs(kShaderTypeCompute); + cs.SetFormat(kShaderFormatOpenCLC); + p.AddShader(&cs, kShaderTypeCompute); + p.SetShaderEntryPoint(&cs, "my_main"); + + Pipeline::ShaderInfo::DescriptorMapEntry entry1; + entry1.kind = Pipeline::ShaderInfo::DescriptorMapEntry::Kind::SSBO; + entry1.descriptor_set = 4; + entry1.binding = 5; + entry1.arg_name = "arg_a"; + entry1.arg_ordinal = 0; + p.GetShaders()[0].AddDescriptorEntry("my_main", std::move(entry1)); + + Pipeline::ShaderInfo::DescriptorMapEntry entry2; + entry2.kind = Pipeline::ShaderInfo::DescriptorMapEntry::Kind::SSBO; + entry2.descriptor_set = 3; + entry2.binding = 1; + entry2.arg_name = "arg_b"; + entry2.arg_ordinal = 1; + p.GetShaders()[0].AddDescriptorEntry("my_main", std::move(entry2)); + + auto a_buf = MakeUnique(BufferType::kStorage); + a_buf->SetName("buf1"); + p.AddBuffer(a_buf.get(), "arg_a"); + + auto b_buf = MakeUnique(BufferType::kUniform); + b_buf->SetName("buf2"); + p.AddBuffer(b_buf.get(), 1); + + auto r = p.UpdateOpenCLBufferBindings(); + + ASSERT_FALSE(r.IsSuccess()); + EXPECT_EQ("Buffer buf2 must be an uniform binding", r.Error()); +} #endif // AMBER_ENABLE_CLSPV } // namespace amber diff --git a/tests/cases/opencl_bind_buffer.amber b/tests/cases/opencl_bind_buffer.amber index d25837ea7..b47739a0b 100644 --- a/tests/cases/opencl_bind_buffer.amber +++ b/tests/cases/opencl_bind_buffer.amber @@ -50,8 +50,8 @@ PIPELINE compute my_pipeline SPECIALIZE 1 AS uint32 2 \ SPECIALIZE 2 AS uint32 2 - BIND BUFFER in_buf AS storage KERNEL ARG_NAME in - BIND BUFFER out_buf AS storage KERNEL ARG_NUMBER 1 + BIND BUFFER in_buf KERNEL ARG_NAME in + BIND BUFFER out_buf KERNEL ARG_NUMBER 1 END RUN my_pipeline 2 2 2 From 18e3b1bd386702b2265fdaff860ba4aba6ddeddb Mon Sep 17 00:00:00 2001 From: Alan Baker Date: Fri, 19 Jul 2019 14:57:37 -0400 Subject: [PATCH 9/9] Improve docs --- docs/amber_script.md | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/docs/amber_script.md b/docs/amber_script.md index e42e7dad6..a7a552c0b 100644 --- a/docs/amber_script.md +++ b/docs/amber_script.md @@ -257,12 +257,13 @@ attachment content, depth/stencil content, uniform buffers, etc. # Bind OpenCL argument buffer by name. Specifying the buffer type is optional. # Amber will set the type as appropriate for the argument buffer. All uses - # of the buffer must be across all pipelines. + # of the buffer must have a consistent |buffer_type| across all pipelines. BIND BUFFER {buffer_name} [AS {buffer_type}] KERNEL ARG_NAME _name_ - # Bind OpenCL argument buffer by argument ordinal. Specifying the buffer type - # is optional. Amber will set the type as appropriate for the argument - # buffer. All uses of the buffer must be across all pipelines. + # Bind OpenCL argument buffer by argument ordinal. Arguments use 0-based + # numbering. Specifying the buffer type is optional. Amber will set the + # type as appropriate for the argument buffer. All uses of the buffer + # must have a consistent |buffer_type| across all pipelines. BIND BUFFER {buffer_name} [AS {buffer_type}] KERNEL ARG_NUMBER _number_ ```