aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--docs/amber_script.md10
-rw-r--r--src/clspv_helper.cc111
-rw-r--r--src/clspv_helper.h1
-rw-r--r--src/executor.cc6
-rw-r--r--src/pipeline.cc92
-rw-r--r--src/pipeline.h10
-rw-r--r--src/pipeline_test.cc32
-rw-r--r--src/shader_compiler.cc9
-rw-r--r--src/shader_compiler.h10
-rw-r--r--src/shader_compiler_test.cc86
-rw-r--r--tests/cases/opencl_read_image_literal_sampler.amber69
-rwxr-xr-xtests/run_tests.py1
12 files changed, 362 insertions, 75 deletions
diff --git a/docs/amber_script.md b/docs/amber_script.md
index 03a23b0..264b85c 100644
--- a/docs/amber_script.md
+++ b/docs/amber_script.md
@@ -239,6 +239,16 @@ SAMPLER {name} \
Note: unnormalized coordinates will override MIN\_LOD and MAX\_LOD to 0.0.
+#### OpenCL Literal Samplers
+
+Literal constant samplers defined in the OpenCL program are automatically
+generated and bound to the pipeline in Amber.
+
+Note: currently the border color is always transparent black.
+
+Note: the addressing mode is used for all coordinates currently. Arrayed images
+should use `clamp_to_edge` for the array index.
+
### Pipelines
#### Pipeline type
diff --git a/src/clspv_helper.cc b/src/clspv_helper.cc
index f4c48c9..3dbb696 100644
--- a/src/clspv_helper.cc
+++ b/src/clspv_helper.cc
@@ -23,6 +23,7 @@ namespace amber {
namespace clspvhelper {
Result Compile(Pipeline::ShaderInfo* shader_info,
+ Pipeline* pipeline,
std::vector<uint32_t>* generated_binary) {
std::vector<clspv::version0::DescriptorMapEntry> entries;
const auto& src_str = shader_info->GetShader()->GetData();
@@ -36,63 +37,69 @@ Result Compile(Pipeline::ShaderInfo* shader_info,
}
for (auto& entry : entries) {
- if (entry.kind != clspv::version0::DescriptorMapEntry::KernelArg) {
- return Result(
- "Only kernel argument descriptor entries are currently supported");
+ if (entry.kind == clspv::version0::DescriptorMapEntry::Constant) {
+ return Result("Constant descriptor entries are not 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::ReadOnlyImage:
- descriptor_entry.kind =
- Pipeline::ShaderInfo::DescriptorMapEntry::Kind::RO_IMAGE;
- break;
- case clspv::ArgKind::WriteOnlyImage:
- descriptor_entry.kind =
- Pipeline::ShaderInfo::DescriptorMapEntry::Kind::WO_IMAGE;
- break;
- case clspv::ArgKind::Sampler:
- descriptor_entry.kind =
- Pipeline::ShaderInfo::DescriptorMapEntry::Kind::SAMPLER;
- break;
- case clspv::ArgKind::Local:
- // Local arguments are handled via specialization constants.
- break;
- default:
- return Result("Unsupported kernel argument descriptor entry");
- }
+ if (entry.kind == clspv::version0::DescriptorMapEntry::KernelArg) {
+ 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::ReadOnlyImage:
+ descriptor_entry.kind =
+ Pipeline::ShaderInfo::DescriptorMapEntry::Kind::RO_IMAGE;
+ break;
+ case clspv::ArgKind::WriteOnlyImage:
+ descriptor_entry.kind =
+ Pipeline::ShaderInfo::DescriptorMapEntry::Kind::WO_IMAGE;
+ break;
+ case clspv::ArgKind::Sampler:
+ descriptor_entry.kind =
+ Pipeline::ShaderInfo::DescriptorMapEntry::Kind::SAMPLER;
+ 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) {
- descriptor_entry.pod_offset = entry.kernel_arg_data.pod_offset;
- descriptor_entry.pod_arg_size = entry.kernel_arg_data.pod_arg_size;
- }
+ 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;
+ 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));
+ shader_info->AddDescriptorEntry(entry.kernel_arg_data.kernel_name,
+ std::move(descriptor_entry));
+ } else {
+ assert(entry.kind == clspv::version0::DescriptorMapEntry::Sampler);
+ // Create a new sampler info.
+ pipeline->AddSampler(entry.sampler_data.mask, entry.descriptor_set,
+ entry.binding);
+ }
}
return Result();
diff --git a/src/clspv_helper.h b/src/clspv_helper.h
index 0df2d05..82d8db3 100644
--- a/src/clspv_helper.h
+++ b/src/clspv_helper.h
@@ -27,6 +27,7 @@ namespace clspvhelper {
// Passes the OpenCL C source code to Clspv.
// Returns the generated SPIR-V binary via |generated_binary| argument.
Result Compile(Pipeline::ShaderInfo* shader_info,
+ Pipeline* pipeline,
std::vector<uint32_t>* generated_binary);
} // namespace clspvhelper
diff --git a/src/executor.cc b/src/executor.cc
index 28377f9..17c2cd2 100644
--- a/src/executor.cc
+++ b/src/executor.cc
@@ -39,8 +39,7 @@ Result Executor::CompileShaders(const amber::Script* script,
Result r;
std::vector<uint32_t> data;
- std::tie(r, data) =
- sc.Compile(pipeline->GetName(), &shader_info, shader_map);
+ std::tie(r, data) = sc.Compile(pipeline.get(), &shader_info, shader_map);
if (!r.IsSuccess())
return r;
@@ -69,6 +68,9 @@ Result Executor::Execute(Engine* engine,
r = pipeline->GenerateOpenCLPodBuffers();
if (!r.IsSuccess())
return r;
+ r = pipeline->GenerateOpenCLLiteralSamplers();
+ if (!r.IsSuccess())
+ return r;
}
for (auto& pipeline : script->GetPipelines()) {
diff --git a/src/pipeline.cc b/src/pipeline.cc
index 7a03cac..747a6c3 100644
--- a/src/pipeline.cc
+++ b/src/pipeline.cc
@@ -27,6 +27,20 @@ namespace {
const char* kDefaultColorBufferFormat = "B8G8R8A8_UNORM";
const char* kDefaultDepthBufferFormat = "D32_SFLOAT_S8_UINT";
+// OpenCL coordinates mode is bit 0
+const uint32_t kOpenCLNormalizedCoordsBit = 1;
+// OpenCL address mode bits are bits 1,2,3.
+const uint32_t kOpenCLAddressModeBits = 0xe;
+// OpenCL address mode bit values.
+const uint32_t kOpenCLAddressModeNone = 0;
+const uint32_t kOpenCLAddressModeClampToEdge = 2;
+const uint32_t kOpenCLAddressModeClamp = 4;
+const uint32_t kOpenCLAddressModeRepeat = 6;
+const uint32_t kOpenCLAddressModeMirroredRepeat = 8;
+// OpenCL filter mode bits.
+const uint32_t kOpenCLFilterModeNearestBit = 0x10;
+const uint32_t kOpenCLFilterModeLinearBit = 0x20;
+
} // namespace
const char* Pipeline::kGeneratedColorBuffer = "framebuffer";
@@ -457,6 +471,7 @@ void Pipeline::AddSampler(Sampler* sampler,
auto& info = samplers_.back();
info.descriptor_set = descriptor_set;
info.binding = binding;
+ info.mask = std::numeric_limits<uint32_t>::max();
}
void Pipeline::AddSampler(Sampler* sampler, const std::string& arg_name) {
@@ -474,6 +489,7 @@ void Pipeline::AddSampler(Sampler* sampler, const std::string& 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();
+ info.mask = std::numeric_limits<uint32_t>::max();
}
void Pipeline::AddSampler(Sampler* sampler, uint32_t arg_no) {
@@ -490,6 +506,20 @@ void Pipeline::AddSampler(Sampler* sampler, uint32_t arg_no) {
info.arg_no = arg_no;
info.descriptor_set = std::numeric_limits<uint32_t>::max();
info.binding = std::numeric_limits<uint32_t>::max();
+ info.mask = std::numeric_limits<uint32_t>::max();
+}
+
+void Pipeline::AddSampler(uint32_t mask,
+ uint32_t descriptor_set,
+ uint32_t binding) {
+ samplers_.push_back(SamplerInfo{nullptr});
+
+ auto& info = samplers_.back();
+ info.arg_name = "";
+ info.arg_no = std::numeric_limits<uint32_t>::max();
+ info.mask = mask;
+ info.descriptor_set = descriptor_set;
+ info.binding = binding;
}
Result Pipeline::UpdateOpenCLBufferBindings() {
@@ -723,4 +753,66 @@ Result Pipeline::GenerateOpenCLPodBuffers() {
return {};
}
+Result Pipeline::GenerateOpenCLLiteralSamplers() {
+ for (auto& info : samplers_) {
+ if (info.sampler || info.mask == std::numeric_limits<uint32_t>::max())
+ continue;
+
+ auto literal_sampler = MakeUnique<Sampler>();
+ literal_sampler->SetName("literal." + std::to_string(info.descriptor_set) +
+ "." + std::to_string(info.binding));
+
+ // The values for addressing modes, filtering modes and coordinate
+ // normalization are all defined in the OpenCL header.
+
+ literal_sampler->SetNormalizedCoords(info.mask &
+ kOpenCLNormalizedCoordsBit);
+
+ uint32_t addressing_bits = info.mask & kOpenCLAddressModeBits;
+ AddressMode addressing_mode = AddressMode::kUnknown;
+ if (addressing_bits == kOpenCLAddressModeNone ||
+ addressing_bits == kOpenCLAddressModeClampToEdge) {
+ // CLK_ADDRESS_NONE
+ // CLK_ADDERSS_CLAMP_TO_EDGE
+ addressing_mode = AddressMode::kClampToEdge;
+ } else if (addressing_bits == kOpenCLAddressModeClamp) {
+ // CLK_ADDRESS_CLAMP
+ addressing_mode = AddressMode::kClampToBorder;
+ } else if (addressing_bits == kOpenCLAddressModeRepeat) {
+ // CLK_ADDRESS_REPEAT
+ addressing_mode = AddressMode::kRepeat;
+ } else if (addressing_bits == kOpenCLAddressModeMirroredRepeat) {
+ // CLK_ADDRESS_MIRRORED_REPEAT
+ addressing_mode = AddressMode::kMirroredRepeat;
+ }
+ literal_sampler->SetAddressModeU(addressing_mode);
+ literal_sampler->SetAddressModeV(addressing_mode);
+ // TODO(alan-baker): If this is used with an arrayed image then W should use
+ // kClampToEdge always, but this information is not currently available.
+ literal_sampler->SetAddressModeW(addressing_mode);
+
+ // Next bit is filtering mode.
+ FilterType filtering_mode = FilterType::kUnknown;
+ if (info.mask & kOpenCLFilterModeNearestBit) {
+ filtering_mode = FilterType::kNearest;
+ } else if (info.mask & kOpenCLFilterModeLinearBit) {
+ filtering_mode = FilterType::kLinear;
+ }
+ literal_sampler->SetMagFilter(filtering_mode);
+ literal_sampler->SetMinFilter(filtering_mode);
+
+ // TODO(alan-baker): OpenCL wants the border color to be based on image
+ // channel orders which aren't accessible.
+
+ // clspv never generates multiple MIPMAP levels.
+ literal_sampler->SetMinLOD(0.0f);
+ literal_sampler->SetMaxLOD(0.0f);
+
+ opencl_literal_samplers_.push_back(std::move(literal_sampler));
+ info.sampler = opencl_literal_samplers_.back().get();
+ }
+
+ return {};
+}
+
} // namespace amber
diff --git a/src/pipeline.h b/src/pipeline.h
index 56fec41..c545a59 100644
--- a/src/pipeline.h
+++ b/src/pipeline.h
@@ -151,6 +151,7 @@ class Pipeline {
uint32_t binding = 0;
std::string arg_name = "";
uint32_t arg_no = 0;
+ uint32_t mask = 0;
};
static const char* kGeneratedColorBuffer;
@@ -262,6 +263,10 @@ class Pipeline {
void AddSampler(Sampler* sampler, const std::string& arg_name);
/// Adds |sampler| to the pieline at the given |arg_no|.
void AddSampler(Sampler* sampler, uint32_t arg_no);
+ /// Adds an entry for an OpenCL literal sampler.
+ void AddSampler(uint32_t sampler_mask,
+ uint32_t descriptor_set,
+ uint32_t binding);
/// Returns information on all samplers in this pipeline.
const std::vector<SamplerInfo>& GetSamplers() const { return samplers_; }
@@ -305,6 +310,10 @@ class Pipeline {
/// command. This should be called after all other buffers are bound.
Result GenerateOpenCLPodBuffers();
+ /// Generate the samplers necessary for OpenCL literal samplers from the
+ /// descriptor map. This should be called after all other samplers are bound.
+ Result GenerateOpenCLLiteralSamplers();
+
private:
void UpdateFramebufferSizes();
@@ -331,6 +340,7 @@ class Pipeline {
std::vector<std::unique_ptr<Buffer>> opencl_pod_buffers_;
/// Maps (descriptor set, binding) to the buffer for that binding pair.
std::map<std::pair<uint32_t, uint32_t>, Buffer*> opencl_pod_buffer_map_;
+ std::vector<std::unique_ptr<Sampler>> opencl_literal_samplers_;
};
} // namespace amber
diff --git a/src/pipeline_test.cc b/src/pipeline_test.cc
index 8bbd309..7a2b2fe 100644
--- a/src/pipeline_test.cc
+++ b/src/pipeline_test.cc
@@ -789,4 +789,36 @@ TEST_F(PipelineTest, OpenCLClone) {
EXPECT_EQ(4U, b2.buffer->ValueCount());
}
+TEST_F(PipelineTest, OpenCLGenerateLiteralSamplers) {
+ Pipeline p(PipelineType::kCompute);
+ p.SetName("my_pipeline");
+
+ p.AddSampler(16, 0, 0);
+ p.AddSampler(41, 0, 1);
+
+ auto r = p.GenerateOpenCLLiteralSamplers();
+ ASSERT_TRUE(r.IsSuccess());
+ for (auto& info : p.GetSamplers()) {
+ if (info.mask == 16) {
+ EXPECT_NE(nullptr, info.sampler);
+ EXPECT_EQ(FilterType::kNearest, info.sampler->GetMagFilter());
+ EXPECT_EQ(FilterType::kNearest, info.sampler->GetMinFilter());
+ EXPECT_EQ(AddressMode::kClampToEdge, info.sampler->GetAddressModeU());
+ EXPECT_EQ(AddressMode::kClampToEdge, info.sampler->GetAddressModeV());
+ EXPECT_EQ(AddressMode::kClampToEdge, info.sampler->GetAddressModeW());
+ EXPECT_EQ(0.0f, info.sampler->GetMinLOD());
+ EXPECT_EQ(0.0f, info.sampler->GetMaxLOD());
+ } else {
+ EXPECT_NE(nullptr, info.sampler);
+ EXPECT_EQ(FilterType::kLinear, info.sampler->GetMagFilter());
+ EXPECT_EQ(FilterType::kLinear, info.sampler->GetMinFilter());
+ EXPECT_EQ(AddressMode::kMirroredRepeat, info.sampler->GetAddressModeU());
+ EXPECT_EQ(AddressMode::kMirroredRepeat, info.sampler->GetAddressModeV());
+ EXPECT_EQ(AddressMode::kMirroredRepeat, info.sampler->GetAddressModeW());
+ EXPECT_EQ(0.0f, info.sampler->GetMinLOD());
+ EXPECT_EQ(0.0f, info.sampler->GetMaxLOD());
+ }
+ }
+}
+
} // namespace amber
diff --git a/src/shader_compiler.cc b/src/shader_compiler.cc
index f6c13ba..acb25fa 100644
--- a/src/shader_compiler.cc
+++ b/src/shader_compiler.cc
@@ -54,11 +54,12 @@ ShaderCompiler::ShaderCompiler(const std::string& env,
ShaderCompiler::~ShaderCompiler() = default;
std::pair<Result, std::vector<uint32_t>> ShaderCompiler::Compile(
- std::string pipeline_name,
+ Pipeline* pipeline,
Pipeline::ShaderInfo* shader_info,
const ShaderMap& shader_map) const {
const auto shader = shader_info->GetShader();
std::string key = shader->GetName();
+ const std::string pipeline_name = pipeline->GetName();
if (pipeline_name != "") {
key = pipeline_name + "-" + key;
}
@@ -140,7 +141,7 @@ std::pair<Result, std::vector<uint32_t>> ShaderCompiler::Compile(
#if AMBER_ENABLE_CLSPV
} else if (shader->GetFormat() == kShaderFormatOpenCLC) {
- Result r = CompileOpenCLC(shader_info, &results);
+ Result r = CompileOpenCLC(shader_info, pipeline, &results);
if (!r.IsSuccess())
return {r, {}};
#endif // AMBER_ENABLE_CLSPV
@@ -279,11 +280,13 @@ Result ShaderCompiler::CompileHlsl(const Shader*,
#if AMBER_ENABLE_CLSPV
Result ShaderCompiler::CompileOpenCLC(Pipeline::ShaderInfo* shader_info,
+ Pipeline* pipeline,
std::vector<uint32_t>* result) const {
- return clspvhelper::Compile(shader_info, result);
+ return clspvhelper::Compile(shader_info, pipeline, result);
}
#else
Result ShaderCompiler::CompileOpenCLC(Pipeline::ShaderInfo*,
+ Pipeline*,
std::vector<uint32_t>*) const {
return {};
}
diff --git a/src/shader_compiler.h b/src/shader_compiler.h
index 19f26dd..fdf7430 100644
--- a/src/shader_compiler.h
+++ b/src/shader_compiler.h
@@ -42,11 +42,12 @@ class ShaderCompiler {
/// entry in |shader_map| for that shader, then the SPIRV-Tools optimizer will
/// be invoked to produce the shader binary.
///
- /// |pipeline_name| is prefixed to shader name to distinguish between
- /// shaders used in multiple pipelines with different optimization
- /// flags.
+ /// |pipeline| is the pipeline containing |shader_info|. The name is used to
+ /// prefix shaders used in multiple pipelines with different optimization
+ /// flags. The pipeline is used in OPENCL-C compiles to create the literal
+ /// sampler bindings.
std::pair<Result, std::vector<uint32_t>> Compile(
- std::string pipeline_name,
+ Pipeline* pipeline,
Pipeline::ShaderInfo* shader_info,
const ShaderMap& shader_map) const;
@@ -55,6 +56,7 @@ class ShaderCompiler {
Result CompileGlsl(const Shader* shader, std::vector<uint32_t>* result) const;
Result CompileHlsl(const Shader* shader, std::vector<uint32_t>* result) const;
Result CompileOpenCLC(Pipeline::ShaderInfo* shader,
+ Pipeline* pipeline,
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 1c78411..d4c18d2 100644
--- a/src/shader_compiler_test.cc
+++ b/src/shader_compiler_test.cc
@@ -19,6 +19,7 @@
#include <vector>
#include "gtest/gtest.h"
+#include "src/sampler.h"
#include "src/shader_data.h"
#if AMBER_ENABLE_SHADERC
#include "shaderc/env.h"
@@ -107,7 +108,8 @@ void main() {
Result r;
std::vector<uint32_t> binary;
Pipeline::ShaderInfo shader_info(&shader, kShaderTypeCompute);
- std::tie(r, binary) = sc.Compile("", &shader_info, ShaderMap());
+ Pipeline pipeline(PipelineType::kCompute);
+ std::tie(r, binary) = sc.Compile(&pipeline, &shader_info, ShaderMap());
ASSERT_TRUE(r.IsSuccess()) << r.Error();
EXPECT_FALSE(binary.empty());
EXPECT_EQ(0x07230203, binary[0]); // Verify SPIR-V header present.
@@ -125,7 +127,8 @@ TEST_F(ShaderCompilerTest, CompilesSpirvAsm) {
Result r;
std::vector<uint32_t> binary;
Pipeline::ShaderInfo shader_info(&shader, kShaderTypeCompute);
- std::tie(r, binary) = sc.Compile("", &shader_info, ShaderMap());
+ Pipeline pipeline(PipelineType::kCompute);
+ std::tie(r, binary) = sc.Compile(&pipeline, &shader_info, ShaderMap());
ASSERT_TRUE(r.IsSuccess());
EXPECT_FALSE(binary.empty());
EXPECT_EQ(0x07230203, binary[0]); // Verify SPIR-V header present.
@@ -144,7 +147,8 @@ TEST_F(ShaderCompilerTest, InvalidSpirvHex) {
Result r;
std::vector<uint32_t> binary;
Pipeline::ShaderInfo shader_info(&shader, kShaderTypeCompute);
- std::tie(r, binary) = sc.Compile("", &shader_info, ShaderMap());
+ Pipeline pipeline(PipelineType::kCompute);
+ std::tie(r, binary) = sc.Compile(&pipeline, &shader_info, ShaderMap());
ASSERT_FALSE(r.IsSuccess());
EXPECT_EQ("Invalid shader: error: line 0: Invalid SPIR-V magic number.\n",
r.Error());
@@ -160,7 +164,8 @@ TEST_F(ShaderCompilerTest, InvalidHex) {
Result r;
std::vector<uint32_t> binary;
Pipeline::ShaderInfo shader_info(&shader, kShaderTypeCompute);
- std::tie(r, binary) = sc.Compile("", &shader_info, ShaderMap());
+ Pipeline pipeline(PipelineType::kCompute);
+ std::tie(r, binary) = sc.Compile(&pipeline, &shader_info, ShaderMap());
ASSERT_FALSE(r.IsSuccess());
EXPECT_EQ("Invalid shader: error: line 0: Invalid SPIR-V magic number.\n",
r.Error());
@@ -211,11 +216,12 @@ OpFunctionEnd
ShaderCompiler sc;
Result r;
std::vector<uint32_t> unopt_binary;
- std::tie(r, unopt_binary) = sc.Compile("", &unoptimized, ShaderMap());
+ Pipeline pipeline(PipelineType::kCompute);
+ std::tie(r, unopt_binary) = sc.Compile(&pipeline, &unoptimized, ShaderMap());
ASSERT_TRUE(r.IsSuccess());
std::vector<uint32_t> opt_binary;
- std::tie(r, opt_binary) = sc.Compile("", &optimized, ShaderMap());
+ std::tie(r, opt_binary) = sc.Compile(&pipeline, &optimized, ShaderMap());
ASSERT_TRUE(r.IsSuccess());
EXPECT_NE(opt_binary.size(), unopt_binary.size());
}
@@ -231,7 +237,8 @@ TEST_F(ShaderCompilerTest, CompilesSpirvHex) {
Result r;
std::vector<uint32_t> binary;
Pipeline::ShaderInfo shader_info(&shader, kShaderTypeCompute);
- std::tie(r, binary) = sc.Compile("", &shader_info, ShaderMap());
+ Pipeline pipeline(PipelineType::kCompute);
+ std::tie(r, binary) = sc.Compile(&pipeline, &shader_info, ShaderMap());
ASSERT_TRUE(r.IsSuccess());
EXPECT_FALSE(binary.empty());
EXPECT_EQ(0x07230203, binary[0]); // Verify SPIR-V header present.
@@ -249,7 +256,8 @@ TEST_F(ShaderCompilerTest, FailsOnInvalidShader) {
Result r;
std::vector<uint32_t> binary;
Pipeline::ShaderInfo shader_info(&shader, kShaderTypeCompute);
- std::tie(r, binary) = sc.Compile("", &shader_info, ShaderMap());
+ Pipeline pipeline(PipelineType::kCompute);
+ std::tie(r, binary) = sc.Compile(&pipeline, &shader_info, ShaderMap());
ASSERT_FALSE(r.IsSuccess());
}
@@ -274,7 +282,9 @@ TEST_F(ShaderCompilerTest, ReturnsCachedShader) {
Result r;
std::vector<uint32_t> binary;
Pipeline::ShaderInfo shader_info(&shader, kShaderTypeCompute);
- std::tie(r, binary) = sc.Compile("pipeline", &shader_info, map);
+ Pipeline pipeline(PipelineType::kCompute);
+ pipeline.SetName("pipeline");
+ std::tie(r, binary) = sc.Compile(&pipeline, &shader_info, map);
ASSERT_TRUE(r.IsSuccess()) << r.Error();
ASSERT_EQ(binary.size(), src_bytes.size());
@@ -298,7 +308,8 @@ kernel void TestShader(global int* in, global int* out) {
Result r;
std::vector<uint32_t> binary;
Pipeline::ShaderInfo shader_info(&shader, kShaderTypeCompute);
- std::tie(r, binary) = sc.Compile("", &shader_info, ShaderMap());
+ Pipeline pipeline(PipelineType::kCompute);
+ std::tie(r, binary) = sc.Compile(&pipeline, &shader_info, ShaderMap());
ASSERT_TRUE(r.IsSuccess());
EXPECT_FALSE(binary.empty());
EXPECT_EQ(0x07230203, binary[0]); // Verify SPIR-V header present.
@@ -324,7 +335,8 @@ kernel void TestShader(global int* in, global int* out) {
Result r;
std::vector<uint32_t> binary;
Pipeline::ShaderInfo shader_info(&shader, kShaderTypeCompute);
- std::tie(r, binary) = sc.Compile("", &shader_info, map);
+ Pipeline pipeline(PipelineType::kCompute);
+ std::tie(r, binary) = sc.Compile(&pipeline, &shader_info, map);
ASSERT_FALSE(r.IsSuccess());
EXPECT_TRUE(binary.empty());
}
@@ -344,7 +356,8 @@ kernel void TestShader(global int* in, global int* out, int m, int b) {
Result r;
std::vector<uint32_t> binary;
Pipeline::ShaderInfo shader_info1(&shader, kShaderTypeCompute);
- std::tie(r, binary) = sc.Compile("", &shader_info1, ShaderMap());
+ Pipeline pipeline(PipelineType::kCompute);
+ std::tie(r, binary) = sc.Compile(&pipeline, &shader_info1, ShaderMap());
ASSERT_TRUE(r.IsSuccess());
EXPECT_FALSE(binary.empty());
EXPECT_EQ(0x07230203, binary[0]); // Verify SPIR-V header present.
@@ -363,7 +376,7 @@ kernel void TestShader(global int* in, global int* out, int m, int b) {
binary.clear();
Pipeline::ShaderInfo shader_info2(&shader, kShaderTypeCompute);
shader_info2.SetCompileOptions({"-cluster-pod-kernel-args", "-pod-ubo"});
- std::tie(r, binary) = sc.Compile("", &shader_info2, ShaderMap());
+ std::tie(r, binary) = sc.Compile(&pipeline, &shader_info2, ShaderMap());
ASSERT_TRUE(r.IsSuccess());
EXPECT_FALSE(binary.empty());
EXPECT_EQ(0x07230203, binary[0]); // Verify SPIR-V header present.
@@ -398,7 +411,8 @@ kernel void TestShader(read_only image2d_t ro_image, write_only image2d_t wo_ima
Result r;
std::vector<uint32_t> binary;
Pipeline::ShaderInfo shader_info1(&shader, kShaderTypeCompute);
- std::tie(r, binary) = sc.Compile("", &shader_info1, ShaderMap());
+ Pipeline pipeline(PipelineType::kCompute);
+ std::tie(r, binary) = sc.Compile(&pipeline, &shader_info1, ShaderMap());
ASSERT_TRUE(r.IsSuccess());
EXPECT_FALSE(binary.empty());
EXPECT_EQ(0x07230203, binary[0]); // Verify SPIR-V header present.
@@ -418,6 +432,50 @@ kernel void TestShader(read_only image2d_t ro_image, write_only image2d_t wo_ima
}
}
}
+
+TEST_F(ShaderCompilerTest, ClspvLiteralSamplers) {
+ std::string data = R"(
+const sampler_t s1 = CLK_ADDRESS_NONE | CLK_FILTER_NEAREST | CLK_NORMALIZED_COORDS_FALSE;
+const sampler_t s2 = CLK_ADDRESS_MIRRORED_REPEAT | CLK_FILTER_LINEAR | CLK_NORMALIZED_COORDS_TRUE;
+
+kernel void foo(read_only image2d_t im, global float4* out) {
+ out[0] = read_imagef(im, s1, (int2)(0));
+ out[1] = read_imagef(im, s2, (int2)(0));
+}
+)";
+
+ Pipeline pipeline(PipelineType::kCompute);
+ pipeline.SetName("pipe");
+ Shader shader(kShaderTypeCompute);
+ shader.SetName("foo");
+ shader.SetFormat(kShaderFormatOpenCLC);
+ shader.SetData(data);
+
+ ShaderCompiler sc;
+ Result r;
+ std::vector<uint32_t> binary;
+ Pipeline::ShaderInfo shader_info1(&shader, kShaderTypeCompute);
+ std::tie(r, binary) = sc.Compile(&pipeline, &shader_info1, ShaderMap());
+ ASSERT_TRUE(r.IsSuccess());
+ EXPECT_FALSE(binary.empty());
+ EXPECT_EQ(0x07230203, binary[0]); // Verify SPIR-V header present.
+ bool found_s1 = false;
+ bool found_s2 = false;
+ EXPECT_EQ(0, pipeline.GetSamplers()[0].descriptor_set);
+ EXPECT_EQ(0, pipeline.GetSamplers()[1].descriptor_set);
+ EXPECT_NE(pipeline.GetSamplers()[0].binding,
+ pipeline.GetSamplers()[1].binding);
+ if (pipeline.GetSamplers()[0].mask == 0x10 ||
+ pipeline.GetSamplers()[1].mask == 0x10) {
+ found_s1 = true;
+ }
+ if (pipeline.GetSamplers()[0].mask == (0x1 | 0x8 | 0x20) ||
+ pipeline.GetSamplers()[1].mask == (0x1 | 0x8 | 0x20)) {
+ found_s2 = true;
+ }
+ EXPECT_EQ(true, found_s1);
+ EXPECT_EQ(true, found_s2);
+}
#endif // AMBER_ENABLE_CLSPV
struct ParseSpvEnvCase {
diff --git a/tests/cases/opencl_read_image_literal_sampler.amber b/tests/cases/opencl_read_image_literal_sampler.amber
new file mode 100644
index 0000000..4c72cfe
--- /dev/null
+++ b/tests/cases/opencl_read_image_literal_sampler.amber
@@ -0,0 +1,69 @@
+#!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 vertex vs PASSTHROUGH
+
+SHADER fragment fs GLSL
+#version 430
+layout(location = 0) out vec4 color_out;
+void main() {
+ color_out = vec4(1.0, 0.0, 0.0, 1.0);
+}
+END
+
+SHADER compute read_imagef OPENCL-C
+const sampler_t sampler = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_FALSE | CLK_FILTER_NEAREST;
+kernel void foo(read_only image2d_t image, global float4* out) {
+ int gid_x = get_global_id(0);
+ int gid_y = get_global_id(1);
+ int linear = 2 * gid_y + gid_x;
+ float2 coord = (float2)(gid_x, gid_y);
+ out[linear] = read_imagef(image, sampler, coord);
+}
+END
+
+IMAGE texture DATA_TYPE vec4<float> DIM_2D WIDTH 2 HEIGHT 2 FILL 0.0
+BUFFER out_buf DATA_TYPE vec4<float> DATA
+2.0 2.0 2.0 2.0
+2.0 2.0 2.0 2.0
+2.0 2.0 2.0 2.0
+2.0 2.0 2.0 2.0
+END
+SAMPLER sampler
+
+PIPELINE compute read_pipe
+ ATTACH read_imagef ENTRY_POINT foo
+ BIND BUFFER out_buf KERNEL ARG_NAME out
+ BIND BUFFER texture KERNEL ARG_NAME image
+END
+
+PIPELINE graphics fill_red
+ ATTACH vs
+ ATTACH fs
+ FRAMEBUFFER_SIZE 2 2
+ BIND BUFFER texture AS color LOCATION 0
+END
+
+CLEAR_COLOR fill_red 0 0 3 3
+CLEAR fill_red
+RUN fill_red DRAW_RECT POS 0 0 SIZE 2 2
+
+RUN read_pipe 2 2 1
+
+EXPECT out_buf IDX 0 EQ 1.0 0.0 0.0 1.0
+EXPECT out_buf IDX 16 EQ 1.0 0.0 0.0 1.0
+EXPECT out_buf IDX 32 EQ 1.0 0.0 0.0 1.0
+EXPECT out_buf IDX 48 EQ 1.0 0.0 0.0 1.0
+
diff --git a/tests/run_tests.py b/tests/run_tests.py
index e514fd7..2484be1 100755
--- a/tests/run_tests.py
+++ b/tests/run_tests.py
@@ -90,6 +90,7 @@ OPENCL_CASES = [
"opencl_c_copy.amber",
"opencl_read_and_write_image3d_rgba32i.amber",
"opencl_read_image.amber",
+ "opencl_read_image_literal_sampler.amber",
"opencl_set_arg.amber",
"opencl_write_image.amber",
]
generated by cgit v1.2.3 (git 2.25.1) at 2024-07-03 11:36:27 +0000