Bind opencl buffers (#584)

Fixes #426

* New BIND subcommands to specify kernel args by name or ordinal
  * storage class specification is optional, but checked for consistency
* clspv compiles parse descriptor map (partially) into
Pipeline::ShaderInfo
* Bindings updated before pipeline creation
* added compiler and pipeline tests for new functionality
* added new end-to-end test
* Add documentation
* Buffer storage class specification is now optional for OpenCL-C

13 files changed, 691 insertions, 69 deletions

diff --git a/docs/amber_script.md b/docs/amber_script.md
index cbe5c73..a7a552c 100644
--- a/docs/amber_script.md
+++ b/docs/amber_script.md
@@ -254,6 +254,17 @@ 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. 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_NAME _name_
+
+  # 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_
 ```
 
 ```groovy
diff --git a/src/amberscript/parser.cc b/src/amberscript/parser.cc
index 412f33f..1abe3a4 100644
--- a/src/amberscript/parser.cc
+++ b/src/amberscript/parser.cc
@@ -644,65 +644,93 @@ 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");
+  if (token->IsString() && token->AsString() == "AS") {
+    token = tokenizer_->NextToken();
+    if (!token->IsString())
+      return Result("invalid token for BUFFER type");
 
-  token = tokenizer_->NextToken();
-  if (!token->IsString())
-    return Result("invalid token for BUFFER type");
+    if (token->AsString() == "color") {
+      token = tokenizer_->NextToken();
+      if (!token->IsString() || token->AsString() != "LOCATION")
+        return Result("BIND missing LOCATION");
 
-  if (token->AsString() == "color") {
-    token = tokenizer_->NextToken();
-    if (!token->IsString() || token->AsString() != "LOCATION")
-      return Result("BIND missing LOCATION");
+      token = tokenizer_->NextToken();
+      if (!token->IsInteger())
+        return Result("invalid value for BIND LOCATION");
 
-    token = tokenizer_->NextToken();
-    if (!token->IsInteger())
-      return Result("invalid value for BIND LOCATION");
+      buffer->SetBufferType(BufferType::kColor);
 
-    buffer->SetBufferType(BufferType::kColor);
+      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;
 
-    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;
+      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->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");
+      uint32_t descriptor_set = token->AsUint32();
 
-    token = tokenizer_->NextToken();
-    if (!token->IsString() || token->AsString() != "DESCRIPTOR_SET")
-      return Result("missing DESCRIPTOR_SET 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 DESCRIPTOR_SET in BIND command");
-    uint32_t descriptor_set = token->AsUint32();
+      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() && token->AsString() == "KERNEL") {
+      token = tokenizer_->NextToken();
+      if (!token->IsString())
+        return Result("missing kernel arg identifier");
 
-    token = tokenizer_->NextToken();
-    if (!token->IsString() || token->AsString() != "BINDING")
-      return Result("missing BINDING for BIND command");
+      if (token->AsString() == "ARG_NAME") {
+        token = tokenizer_->NextToken();
+        if (!token->IsString())
+          return Result("expected argument identifier");
 
-    token = tokenizer_->NextToken();
-    if (!token->IsInteger())
-      return Result("invalid value for BINDING in BIND command");
-    pipeline->AddBuffer(buffer, descriptor_set, token->AsUint32());
+        pipeline->AddBuffer(buffer, token->AsString());
+      } else if (token->AsString() == "ARG_NUMBER") {
+        token = tokenizer_->NextToken();
+        if (!token->IsInteger())
+          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");
+    }
   }
 
   return ValidateEndOfStatement("BIND command");
diff --git a/src/amberscript/parser_bind_test.cc b/src/amberscript/parser_bind_test.cc
index 7672294..5e6d87c 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) {
@@ -1149,5 +1149,181 @@ 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_NAME arg
+END)";
+
+  Parser parser;
+  Result r = parser.Parse(in);
+  ASSERT_TRUE(r.IsSuccess());
+}
+
+TEST_F(AmberScriptParserTest, BindBufferOpenCLArgNumber) {
+  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_NUMBER 0
+END)";
+
+  Parser parser;
+  Result r = parser.Parse(in);
+  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
+#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_NAME arg
+END)";
+
+  Parser parser;
+  Result r = parser.Parse(in);
+  ASSERT_FALSE(r.IsSuccess());
+  EXPECT_EQ("9: missing DESCRIPTOR_SET or KERNEL 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_NAME or ARG_NUMBER 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 KERNEL ARG_NAME
+END)";
+
+  Parser parser;
+  Result r = parser.Parse(in);
+  ASSERT_FALSE(r.IsSuccess());
+  EXPECT_EQ("10: expected argument identifier", r.Error());
+}
+
+TEST_F(AmberScriptParserTest, BindBufferOpenCLMissingArgNumber) {
+  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_NUMBER
+END)";
+
+  Parser parser;
+  Result r = parser.Parse(in);
+  ASSERT_FALSE(r.IsSuccess());
+  EXPECT_EQ("10: expected argument 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_NAME 0
+END)";
+
+  Parser parser;
+  Result r = parser.Parse(in);
+  ASSERT_FALSE(r.IsSuccess());
+  EXPECT_EQ("9: expected argument identifier", r.Error());
+}
+
+TEST_F(AmberScriptParserTest, BindBufferOpenCLArgNumberNotInteger) {
+  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 in
+END)";
+
+  Parser parser;
+  Result r = parser.Parse(in);
+  ASSERT_FALSE(r.IsSuccess());
+  EXPECT_EQ("9: expected argument number", r.Error());
+}
+
 }  // namespace amberscript
 }  // namespace amber
diff --git a/src/clspv_helper.cc b/src/clspv_helper.cc
index 056b5d5..e0d853e 100644
--- a/src/clspv_helper.cc
+++ b/src/clspv_helper.cc
@@ -14,20 +14,73 @@
 
 #include "src/clspv_helper.h"
 
+#include <utility>
+
+#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<uint32_t>* generated_binary) {
-  // TODO(alan-baker): Parse the descriptor map.
   std::vector<clspv::version0::DescriptorMapEntry> 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::POD_UBO;
+        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;
+    }
+
+    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 3672532..0df2d05 100644
--- a/src/clspv_helper.h
+++ b/src/clspv_helper.h
@@ -19,13 +19,14 @@
 #include <vector>
 
 #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<uint32_t>* generated_binary);
 
 }  // namespace clspvhelper
diff --git a/src/executor.cc b/src/executor.cc
index f74a93b..3bbec7a 100644
--- a/src/executor.cc
+++ b/src/executor.cc
@@ -37,7 +37,7 @@ Result Executor::CompileShaders(const amber::Script* script,
 
       Result r;
       std::vector<uint32_t> 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;
 
@@ -60,6 +60,9 @@ Result Executor::Execute(Engine* engine,
       return r;
 
     for (auto& pipeline : script->GetPipelines()) {
+      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 1a32a42..8f214f0 100644
--- a/src/pipeline.cc
+++ b/src/pipeline.cc
@@ -15,6 +15,7 @@
 #include "src/pipeline.h"
 
 #include <algorithm>
+#include <limits>
 #include <set>
 
 #include "src/format_parser.h"
@@ -330,4 +331,103 @@ 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;
+  info.descriptor_set = std::numeric_limits<uint32_t>::max();
+  info.binding = std::numeric_limits<uint32_t>::max();
+  info.arg_no = std::numeric_limits<uint32_t>::max();
+}
+
+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;
+  info.descriptor_set = std::numeric_limits<uint32_t>::max();
+  info.binding = std::numeric_limits<uint32_t>::max();
+}
+
+Result 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<uint32_t>::max() &&
+        info.binding == std::numeric_limits<uint32_t>::max()) {
+      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 c17a67e..8859888 100644
--- a/src/pipeline.h
+++ b/src/pipeline.h
@@ -18,6 +18,7 @@
 #include <map>
 #include <memory>
 #include <string>
+#include <unordered_map>
 #include <utility>
 #include <vector>
 
@@ -67,6 +68,34 @@ class Pipeline {
       specialization_[spec_id] = value;
     }
 
+    /// Descriptor information for an OpenCL-C shader.
+    struct DescriptorMapEntry {
+      std::string arg_name = "";
+
+      enum class Kind : int {
+        UNKNOWN,
+        SSBO,
+        UBO,
+        POD,
+        POD_UBO,
+      } kind;
+
+      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,
+                            DescriptorMapEntry&& entry) {
+      descriptor_map_[kernel].emplace_back(std::move(entry));
+    }
+    const std::unordered_map<std::string, std::vector<DescriptorMapEntry>>&
+    GetDescriptorMap() const {
+      return descriptor_map_;
+    }
+
    private:
     Shader* shader_ = nullptr;
     ShaderType shader_type_;
@@ -74,6 +103,8 @@ class Pipeline {
     std::string entry_point_;
     std::vector<uint32_t> data_;
     std::map<uint32_t, uint32_t> specialization_;
+    std::unordered_map<std::string, std::vector<DescriptorMapEntry>>
+        descriptor_map_;
   };
 
   /// Information on a buffer attached to the pipeline.
@@ -88,6 +119,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,9 +198,17 @@ 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<BufferInfo>& 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.
+  Result 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 8c61a0b..8e0ddcc 100644
--- a/src/pipeline_test.cc
+++ b/src/pipeline_test.cc
@@ -396,4 +396,90 @@ 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<Buffer>(BufferType::kStorage);
+  a_buf->SetName("buf1");
+  p.AddBuffer(a_buf.get(), "arg_a");
+
+  auto b_buf = MakeUnique<Buffer>(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);
+}
+
+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<Buffer>(BufferType::kStorage);
+  a_buf->SetName("buf1");
+  p.AddBuffer(a_buf.get(), "arg_a");
+
+  auto b_buf = MakeUnique<Buffer>(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/src/shader_compiler.cc b/src/shader_compiler.cc
index 8bf40ae..f2dc41b 100644
--- a/src/shader_compiler.cc
+++ b/src/shader_compiler.cc
@@ -51,11 +51,19 @@ ShaderCompiler::ShaderCompiler(const std::string& env) : spv_env_(env) {}
 ShaderCompiler::~ShaderCompiler() = default;
 
 std::pair<Result, std::vector<uint32_t>> 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())
+  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;
@@ -122,7 +130,7 @@ std::pair<Result, std::vector<uint32_t>> 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 +249,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<uint32_t>* 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<uint32_t>*) const {
   return {};
 }
diff --git a/src/shader_compiler.h b/src/shader_compiler.h
index 3c4b56f..a6dc8f8 100644
--- a/src/shader_compiler.h
+++ b/src/shader_compiler.h
@@ -21,6 +21,7 @@
 
 #include "amber/amber.h"
 #include "amber/result.h"
+#include "src/pipeline.h"
 #include "src/shader.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<Result, std::vector<uint32_t>> Compile(
-      const Shader* shader,
+      Pipeline::ShaderInfo* shader_info,
       const ShaderMap& shader_map) const;
 
  private:
   Result ParseHex(const std::string& data, std::vector<uint32_t>* result) const;
   Result CompileGlsl(const Shader* shader, std::vector<uint32_t>* result) const;
   Result CompileHlsl(const Shader* shader, std::vector<uint32_t>* result) const;
-  Result CompileOpenCLC(const Shader* shader,
+  Result CompileOpenCLC(Pipeline::ShaderInfo* shader,
                         std::vector<uint32_t>* result) const;
 
   std::string spv_env_;
diff --git a/src/shader_compiler_test.cc b/src/shader_compiler_test.cc
index 188a741..e33d6dd 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<uint32_t> 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<uint32_t> 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<uint32_t> 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<uint32_t> 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<uint32_t> 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<uint32_t> 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<uint32_t> 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());
@@ -220,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<uint32_t> 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<uint32_t> src_bytes = {1, 2, 3, 4, 5};
+
+  ShaderMap map;
+  map[name] = src_bytes;
+
+  ShaderCompiler sc;
+  Result r;
+  std::vector<uint32_t> 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/tests/cases/opencl_bind_buffer.amber b/tests/cases/opencl_bind_buffer.amber
new file mode 100644
index 0000000..b47739a
--- /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 KERNEL ARG_NAME in
+  BIND BUFFER out_buf KERNEL ARG_NUMBER 1
+END
+
+RUN my_pipeline 2 2 2
+
+EXPECT out_buf EQ_BUFFER in_buf
+