diff options
| author | alan-baker <alanbaker@google.com> | 2020-11-04 17:00:05 -0500 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2020-11-04 17:00:05 -0500 |
| commit | 8915c4f9a04e703d5273939f25d8fe08010c2b43 (patch) | |
| tree | c426b3aebba3cc132cc3f26ed118b696bc465919 | |
| parent | 760076b7c8bf43d6c6cc4928f5129afca16e895e (diff) | |
| download | amber-8915c4f9a04e703d5273939f25d8fe08010c2b43.tar.gz | |
Add passthrough support for region offset (#922)
* In order to support OpenCL C 2.0 kernels that use get_global_id() pass
0s for the region offset push constant
* Allow use of the region offset push constant reflection instruction
| -rw-r--r-- | src/clspv_helper.cc | 10 | ||||
| -rw-r--r-- | src/pipeline.cc | 6 | ||||
| -rw-r--r-- | src/pipeline.h | 1 |
3 files changed, 16 insertions, 1 deletions
diff --git a/src/clspv_helper.cc b/src/clspv_helper.cc index a6c4d50..59b8737 100644 --- a/src/clspv_helper.cc +++ b/src/clspv_helper.cc @@ -180,9 +180,17 @@ spv_result_t ParseExtendedInst(ReflectionHelper* helper, helper->shader_info->AddPushConstant(std::move(push_constant)); break; } + case NonSemanticClspvReflectionPushConstantRegionOffset: { + auto offset_id = inst->words[inst->operands[4].offset]; + auto size_id = inst->words[inst->operands[5].offset]; + Pipeline::ShaderInfo::PushConstant push_constant{ + Pipeline::ShaderInfo::PushConstant::PushConstantType::kRegionOffset, + helper->constants[offset_id], helper->constants[size_id]}; + helper->shader_info->AddPushConstant(std::move(push_constant)); + break; + } case NonSemanticClspvReflectionPushConstantEnqueuedLocalSize: case NonSemanticClspvReflectionPushConstantGlobalSize: - case NonSemanticClspvReflectionPushConstantRegionOffset: case NonSemanticClspvReflectionPushConstantNumWorkgroups: case NonSemanticClspvReflectionPushConstantRegionGroupOffset: helper->error_message = "Unsupported push constant"; diff --git a/src/pipeline.cc b/src/pipeline.cc index 7df2401..3e5353a 100644 --- a/src/pipeline.cc +++ b/src/pipeline.cc @@ -1015,6 +1015,12 @@ Result Pipeline::GenerateOpenCLPushConstants() { bytes[base + 1] = 0; bytes[base + 2] = 0; break; + case Pipeline::ShaderInfo::PushConstant::PushConstantType::kRegionOffset: + // Region offsets are not currently supported. + bytes[base] = 0; + bytes[base + 1] = 0; + bytes[base + 2] = 0; + break; } memcpy(buf->ValuePtr()->data() + pc.offset, bytes.data(), bytes.size() * sizeof(uint32_t)); diff --git a/src/pipeline.h b/src/pipeline.h index 212255b..ddeb2a4 100644 --- a/src/pipeline.h +++ b/src/pipeline.h @@ -149,6 +149,7 @@ class Pipeline { enum class PushConstantType { kDimensions = 0, kGlobalOffset, + kRegionOffset, }; PushConstantType type; uint32_t offset = 0; |