Google Git
Sign in
android/platform/external/clvk/refs/heads/android17-release/./src/program.cpp
blob: 29c5c57940b95a82d70787d3e8c832ec9cd5c184 [file] [edit]
// Copyright 2018 The clvk authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#include <algorithm>
#include <cstring>
#include <filesystem>
#include <iostream>
#include <map>
#include <sstream>
#include <system_error>
#include <utility>
#include <vector>
#include <vulkan/vulkan.h>
#include "clspv/Sampler.h"
#include "utils.hpp"
#ifdef CLSPV_ONLINE_COMPILER
#ifdef ENABLE_SPIRV_IL
#include "LLVMSPIRVLib.h"
#include "llvm/Bitcode/BitcodeWriter.h"
#include "llvm/Support/CommandLine.h" // FIXME(#380) remove
#include "llvm/Support/raw_ostream.h"
#endif
#include "clspv/Compiler.h"
#endif
#include "spirv-tools/linker.hpp"
#include "spirv-tools/optimizer.hpp"
#include "spirv/unified1/NonSemanticClspvReflection.h"
#include "spirv/unified1/spirv.hpp"
#include "config.hpp"
#include "init.hpp"
#include "log.hpp"
#include "program.hpp"
#include "tracing.hpp"
struct membuf : public std::streambuf {
membuf(const unsigned char* begin, const unsigned char* end) {
auto sbegin =
reinterpret_cast<char*>(const_cast<unsigned char*>(begin));
auto send = reinterpret_cast<char*>(const_cast<unsigned char*>(end));
setg(sbegin, sbegin, send);
}
membuf(unsigned char* begin, unsigned char* end) {
auto sbegin = reinterpret_cast<char*>(begin);
auto send = reinterpret_cast<char*>(end);
setp(sbegin, send);
}
pos_type seekoff(off_type off, std::ios_base::seekdir dir,
std::ios_base::openmode which) override {
UNUSED(which);
char* whence = eback();
if (dir == std::ios_base::cur) {
whence = gptr();
} else if (dir == std::ios_base::end) {
whence = egptr();
}
char* to = whence + off;
if (to >= eback() && to <= egptr()) {
setg(eback(), to, egptr());
return gptr() - eback();
}
return -1;
}
pos_type seekpos(pos_type pos, std::ios_base::openmode which) override {
UNUSED(which);
char* to = eback() + pos;
if (to >= eback() && to <= egptr()) {
setg(eback(), to, egptr());
return gptr() - eback();
}
return -1;
}
};
struct reflection_parse_data {
uint32_t uint_id = 0;
std::unordered_map<uint32_t, uint32_t> constants;
std::unordered_map<uint32_t, std::string> strings;
spir_binary* binary;
std::unordered_map<uint32_t, kernel_argument_info> arg_infos;
};
spv_result_t parse_reflection(void* user_data,
const spv_parsed_instruction_t* inst) {
// Helper function to map instruction to argument type.
auto inst_to_arg_kind = [](uint32_t inst) {
switch (static_cast<NonSemanticClspvReflectionInstructions>(inst)) {
case NonSemanticClspvReflectionArgumentStorageBuffer:
return kernel_argument_kind::buffer;
case NonSemanticClspvReflectionArgumentUniform:
return kernel_argument_kind::buffer_ubo;
case NonSemanticClspvReflectionArgumentPodStorageBuffer:
return kernel_argument_kind::pod;
case NonSemanticClspvReflectionArgumentPodUniform:
return kernel_argument_kind::pod_ubo;
case NonSemanticClspvReflectionArgumentPodPushConstant:
return kernel_argument_kind::pod_pushconstant;
case NonSemanticClspvReflectionArgumentPointerUniform:
return kernel_argument_kind::pointer_ubo;
case NonSemanticClspvReflectionArgumentPointerPushConstant:
return kernel_argument_kind::pointer_pushconstant;
case NonSemanticClspvReflectionArgumentSampledImage:
return kernel_argument_kind::sampled_image;
case NonSemanticClspvReflectionArgumentStorageImage:
return kernel_argument_kind::storage_image;
case NonSemanticClspvReflectionArgumentStorageTexelBuffer:
return kernel_argument_kind::storage_texel_buffer;
case NonSemanticClspvReflectionArgumentUniformTexelBuffer:
return kernel_argument_kind::uniform_texel_buffer;
case NonSemanticClspvReflectionArgumentSampler:
return kernel_argument_kind::sampler;
case NonSemanticClspvReflectionArgumentWorkgroup:
return kernel_argument_kind::local;
default:
cvk_error_fn("Unhandled reflection instruction for arg kind");
break;
}
return kernel_argument_kind::buffer;
};
// Helper function to map instruction to push constant type.
auto inst_to_push_constant = [](uint32_t inst) {
switch (static_cast<NonSemanticClspvReflectionInstructions>(inst)) {
case NonSemanticClspvReflectionPushConstantGlobalOffset:
return pushconstant::global_offset;
case NonSemanticClspvReflectionPushConstantEnqueuedLocalSize:
return pushconstant::enqueued_local_size;
case NonSemanticClspvReflectionPushConstantGlobalSize:
return pushconstant::global_size;
case NonSemanticClspvReflectionPushConstantRegionOffset:
return pushconstant::region_offset;
case NonSemanticClspvReflectionPushConstantNumWorkgroups:
return pushconstant::num_workgroups;
case NonSemanticClspvReflectionPushConstantRegionGroupOffset:
return pushconstant::region_group_offset;
case NonSemanticClspvReflectionImageArgumentInfoChannelOrderPushConstant:
return pushconstant::image_metadata;
case NonSemanticClspvReflectionImageArgumentInfoChannelDataTypePushConstant:
return pushconstant::image_metadata;
case NonSemanticClspvReflectionConstantDataPointerPushConstant:
return pushconstant::module_constants_pointer;
case NonSemanticClspvReflectionPrintfBufferPointerPushConstant:
return pushconstant::printf_buffer_pointer;
case NonSemanticClspvReflectionNormalizedSamplerMaskPushConstant:
return pushconstant::normalized_sampler_mask;
default:
cvk_error_fn("Unhandled reflection instruction for push constant");
break;
}
return pushconstant::global_offset;
};
auto* parse_data = reinterpret_cast<reflection_parse_data*>(user_data);
switch (inst->opcode) {
case spv::OpTypeInt:
if (inst->words[2] == 32 && inst->words[3] == 0) {
parse_data->uint_id = inst->result_id;
}
break;
case spv::OpConstant:
if (inst->words[1] == parse_data->uint_id) {
parse_data->constants[inst->result_id] = inst->words[3];
}
break;
case spv::OpString:
parse_data->strings[inst->result_id] =
std::string(reinterpret_cast<const char*>(&inst->words[2]));
break;
case spv::OpExtInst:
if (inst->ext_inst_type ==
SPV_EXT_INST_TYPE_NONSEMANTIC_CLSPVREFLECTION) {
auto ext_inst = inst->words[4];
switch (ext_inst) {
case NonSemanticClspvReflectionKernel: {
// Record the kernel name.
const auto& name = parse_data->strings[inst->words[6]];
const auto& num_args = parse_data->constants[inst->words[7]];
const auto& flags = parse_data->constants[inst->words[8]];
const auto& attributes = parse_data->strings[inst->words[9]];
parse_data->strings[inst->result_id] = name;
parse_data->binary->add_kernel(name, num_args, attributes,
flags);
break;
}
case NonSemanticClspvReflectionArgumentInfo: {
// Record the argument info.
kernel_argument_info info;
info.name = parse_data->strings[inst->words[5]];
if (inst->num_operands > 6) {
info.type_name = parse_data->strings[inst->words[6]];
info.address_qualifier =
parse_data->constants[inst->words[7]];
info.access_qualifier =
parse_data->constants[inst->words[8]];
info.type_qualifier = parse_data->constants[inst->words[9]];
info.extended_valid = true;
}
parse_data->arg_infos[inst->result_id] = info;
break;
}
case NonSemanticClspvReflectionNormalizedSamplerMaskPushConstant: {
auto kernel = parse_data->strings[inst->words[5]];
auto ordinal = parse_data->constants[inst->words[6]];
auto offset = parse_data->constants[inst->words[7]];
auto size = parse_data->constants[inst->words[8]];
parse_data->binary->add_sampler_metadata(kernel, ordinal,
offset);
auto pc = inst_to_push_constant(ext_inst);
parse_data->binary->add_push_constant(pc, {offset, size});
break;
}
case NonSemanticClspvReflectionImageArgumentInfoChannelOrderPushConstant: {
auto kernel = parse_data->strings[inst->words[5]];
auto ordinal = parse_data->constants[inst->words[6]];
auto offset = parse_data->constants[inst->words[7]];
auto size = parse_data->constants[inst->words[8]];
parse_data->binary->add_image_channel_order_metadata(
kernel, ordinal, offset);
auto pc = inst_to_push_constant(ext_inst);
parse_data->binary->add_push_constant(pc, {offset, size});
break;
}
case NonSemanticClspvReflectionImageArgumentInfoChannelDataTypePushConstant: {
auto kernel = parse_data->strings[inst->words[5]];
auto ordinal = parse_data->constants[inst->words[6]];
auto offset = parse_data->constants[inst->words[7]];
auto size = parse_data->constants[inst->words[8]];
parse_data->binary->add_image_channel_data_type_metadata(
kernel, ordinal, offset);
auto pc = inst_to_push_constant(ext_inst);
parse_data->binary->add_push_constant(pc, {offset, size});
break;
}
case NonSemanticClspvReflectionArgumentStorageBuffer:
case NonSemanticClspvReflectionArgumentUniform:
case NonSemanticClspvReflectionArgumentSampledImage:
case NonSemanticClspvReflectionArgumentStorageImage:
case NonSemanticClspvReflectionArgumentStorageTexelBuffer:
case NonSemanticClspvReflectionArgumentUniformTexelBuffer:
case NonSemanticClspvReflectionArgumentSampler: {
// These arguments have descriptor set, binding and an optional
// arg info.
auto kernel = parse_data->strings[inst->words[5]];
auto ordinal = parse_data->constants[inst->words[6]];
auto descriptor_set = parse_data->constants[inst->words[7]];
if (descriptor_set >= spir_binary::MAX_DESCRIPTOR_SETS)
return SPV_ERROR_INVALID_DATA;
auto binding = parse_data->constants[inst->words[8]];
kernel_argument_info arg_info;
if (inst->num_operands == 9) {
arg_info = parse_data->arg_infos[inst->words[9]];
}
auto kind = inst_to_arg_kind(ext_inst);
kernel_argument arg = {arg_info, ordinal, descriptor_set,
binding, 0, 0,
kind, 0, 0};
parse_data->binary->add_kernel_argument(kernel, std::move(arg));
break;
}
case NonSemanticClspvReflectionArgumentPodStorageBuffer:
case NonSemanticClspvReflectionArgumentPodUniform:
case NonSemanticClspvReflectionArgumentPointerUniform: {
// These arguments have descriptor set, binding, offset, size
// and an optional arg info.
auto kernel = parse_data->strings[inst->words[5]];
auto ordinal = parse_data->constants[inst->words[6]];
auto descriptor_set = parse_data->constants[inst->words[7]];
if (descriptor_set >= spir_binary::MAX_DESCRIPTOR_SETS)
return SPV_ERROR_INVALID_DATA;
auto binding = parse_data->constants[inst->words[8]];
auto offset = parse_data->constants[inst->words[9]];
auto size = parse_data->constants[inst->words[10]];
kernel_argument_info arg_info;
if (inst->num_operands == 11) {
arg_info = parse_data->arg_infos[inst->words[11]];
}
auto kind = inst_to_arg_kind(ext_inst);
kernel_argument arg = {arg_info, ordinal, descriptor_set,
binding, offset, size,
kind, 0, 0};
parse_data->binary->add_kernel_argument(kernel, std::move(arg));
break;
}
case NonSemanticClspvReflectionArgumentPodPushConstant:
case NonSemanticClspvReflectionArgumentPointerPushConstant: {
// These arguments have offset, size and an optional arg info.
auto kernel = parse_data->strings[inst->words[5]];
auto ordinal = parse_data->constants[inst->words[6]];
auto offset = parse_data->constants[inst->words[7]];
auto size = parse_data->constants[inst->words[8]];
kernel_argument_info arg_info;
if (inst->num_operands == 9) {
arg_info = parse_data->arg_infos[inst->words[9]];
}
auto kind = inst_to_arg_kind(ext_inst);
kernel_argument arg = {arg_info, ordinal, 0, 0, offset,
size, kind, 0, 0};
parse_data->binary->add_kernel_argument(kernel, std::move(arg));
break;
}
case NonSemanticClspvReflectionWorkgroupVariableSize: {
auto size = parse_data->constants[inst->words[6]];
parse_data->binary->add_workgroup_variable_size(size);
break;
}
case NonSemanticClspvReflectionArgumentWorkgroup: {
// These arguments have spec id, elem size and an optional arg
// info.
auto kernel = parse_data->strings[inst->words[5]];
auto ordinal = parse_data->constants[inst->words[6]];
auto spec_id = parse_data->constants[inst->words[7]];
auto size = parse_data->constants[inst->words[8]];
kernel_argument_info arg_info;
if (inst->num_operands == 9) {
arg_info = parse_data->arg_infos[inst->words[9]];
}
auto kind = inst_to_arg_kind(ext_inst);
kernel_argument arg = {arg_info, ordinal, 0, 0, 0,
0, kind, spec_id, size};
parse_data->binary->add_kernel_argument(kernel, std::move(arg));
break;
}
case NonSemanticClspvReflectionSpecConstantWorkgroupSize: {
// Reflection encodes all three spec ids in a single
// instruction.
auto x_id = parse_data->constants[inst->words[5]];
auto y_id = parse_data->constants[inst->words[6]];
auto z_id = parse_data->constants[inst->words[7]];
parse_data->binary->add_spec_constant(
spec_constant::workgroup_size_x, x_id);
parse_data->binary->add_spec_constant(
spec_constant::workgroup_size_y, y_id);
parse_data->binary->add_spec_constant(
spec_constant::workgroup_size_z, z_id);
break;
}
case NonSemanticClspvReflectionSpecConstantGlobalOffset: {
// Reflection encodes all three spec ids in a single
// instruction.
auto x_id = parse_data->constants[inst->words[5]];
auto y_id = parse_data->constants[inst->words[6]];
auto z_id = parse_data->constants[inst->words[7]];
parse_data->binary->add_spec_constant(
spec_constant::global_offset_x, x_id);
parse_data->binary->add_spec_constant(
spec_constant::global_offset_y, y_id);
parse_data->binary->add_spec_constant(
spec_constant::global_offset_z, z_id);
break;
}
case NonSemanticClspvReflectionSpecConstantWorkDim: {
auto dim_id = parse_data->constants[inst->words[5]];
parse_data->binary->add_spec_constant(spec_constant::work_dim,
dim_id);
break;
}
case NonSemanticClspvReflectionSpecConstantSubgroupMaxSize: {
auto size_id = parse_data->constants[inst->words[5]];
parse_data->binary->add_spec_constant(
spec_constant::subgroup_max_size, size_id);
break;
}
case NonSemanticClspvReflectionPushConstantGlobalOffset:
case NonSemanticClspvReflectionPushConstantEnqueuedLocalSize:
case NonSemanticClspvReflectionPushConstantGlobalSize:
case NonSemanticClspvReflectionPushConstantRegionOffset:
case NonSemanticClspvReflectionPushConstantNumWorkgroups:
case NonSemanticClspvReflectionPushConstantRegionGroupOffset: {
auto offset = parse_data->constants[inst->words[5]];
auto size = parse_data->constants[inst->words[6]];
auto pc = inst_to_push_constant(ext_inst);
parse_data->binary->add_push_constant(pc, {offset, size});
break;
}
case NonSemanticClspvReflectionLiteralSampler: {
// Track descriptor set and binding. Decode the sampler mask.
auto descriptor_set = parse_data->constants[inst->words[5]];
if (descriptor_set >= spir_binary::MAX_DESCRIPTOR_SETS)
return SPV_ERROR_INVALID_DATA;
auto binding = parse_data->constants[inst->words[6]];
auto mask = parse_data->constants[inst->words[7]];
uint32_t coords = mask & clspv::kSamplerNormalizedCoordsMask;
bool normalized_coords =
coords == clspv::CLK_NORMALIZED_COORDS_TRUE;
cl_addressing_mode addressing;
switch (mask & clspv::kSamplerAddressMask) {
case clspv::CLK_ADDRESS_NONE:
default:
addressing = CL_ADDRESS_NONE;
break;
case clspv::CLK_ADDRESS_CLAMP_TO_EDGE:
addressing = CL_ADDRESS_CLAMP_TO_EDGE;
break;
case clspv::CLK_ADDRESS_CLAMP:
addressing = CL_ADDRESS_CLAMP;
break;
case clspv::CLK_ADDRESS_MIRRORED_REPEAT:
addressing = CL_ADDRESS_MIRRORED_REPEAT;
break;
case clspv::CLK_ADDRESS_REPEAT:
addressing = CL_ADDRESS_REPEAT;
break;
}
cl_filter_mode filter;
switch (mask & clspv::kSamplerFilterMask) {
case clspv::CLK_FILTER_NEAREST:
default:
filter = CL_FILTER_NEAREST;
break;
case clspv::CLK_FILTER_LINEAR:
filter = CL_FILTER_LINEAR;
break;
}
parse_data->binary->add_literal_sampler(
{descriptor_set, binding, normalized_coords, addressing,
filter});
break;
}
case NonSemanticClspvReflectionPropertyRequiredWorkgroupSize: {
auto kernel = parse_data->strings[inst->words[5]];
auto x = parse_data->constants[inst->words[6]];
auto y = parse_data->constants[inst->words[7]];
auto z = parse_data->constants[inst->words[8]];
parse_data->binary->set_required_work_group_size(kernel, x, y,
z);
break;
}
case NonSemanticClspvReflectionConstantDataStorageBuffer:
case NonSemanticClspvReflectionConstantDataPointerPushConstant: {
auto char2int = [](char c) {
if (c >= '0' && c <= '9')
return c - '0';
if (c >= 'A' && c <= 'F')
return c - 'A' + 10;
if (c >= 'a' && c <= 'f')
return c - 'a' + 10;
return 0;
};
auto hex2bin = [&char2int](const char* str, char* bin) {
while (str[0] && str[1]) {
*(bin++) = char2int(str[0]) * 16 + char2int(str[1]);
str += 2;
}
};
auto data = parse_data->strings[inst->words[7]];
if (data.size() & 1) {
cvk_error_fn("invalid constant data buffer string (odd "
"number of digits)");
return SPV_ERROR_INVALID_DATA;
}
constant_data_buffer_info binfo{};
auto data_size = data.size() / 2;
binfo.data.resize(data_size);
hex2bin(data.c_str(), binfo.data.data());
if (ext_inst ==
NonSemanticClspvReflectionConstantDataStorageBuffer) {
binfo.type = module_buffer_type::storage_buffer;
binfo.set = parse_data->constants[inst->words[5]];
if (binfo.set >= spir_binary::MAX_DESCRIPTOR_SETS)
return SPV_ERROR_INVALID_DATA;
binfo.binding = parse_data->constants[inst->words[6]];
} else {
binfo.type = module_buffer_type::pointer_push_constant;
binfo.pc_offset = parse_data->constants[inst->words[5]];
parse_data->binary->add_push_constant(
pushconstant::module_constants_pointer,
{binfo.pc_offset, 8u});
}
parse_data->binary->set_constant_data_buffer(binfo);
break;
}
case NonSemanticClspvReflectionPrintfInfo: {
uint32_t printf_id = parse_data->constants[inst->words[5]];
std::string printf_string = parse_data->strings[inst->words[6]];
std::vector<uint32_t> printf_arg_sizes;
for (int i = 6; i < inst->num_operands; i++) {
printf_arg_sizes.push_back(
parse_data
->constants[inst->words[inst->operands[i].offset]]);
}
parse_data->binary->add_printf_descriptor(
{printf_id, printf_string, printf_arg_sizes});
break;
}
case NonSemanticClspvReflectionPrintfBufferStorageBuffer: {
printf_buffer_desc_info binfo;
binfo.type = module_buffer_type::storage_buffer;
binfo.set = parse_data->constants[inst->words[5]];
binfo.binding = parse_data->constants[inst->words[6]];
binfo.size = parse_data->constants[inst->words[7]];
parse_data->binary->set_printf_buffer_info(binfo);
break;
}
case NonSemanticClspvReflectionPrintfBufferPointerPushConstant: {
printf_buffer_desc_info binfo;
binfo.type = module_buffer_type::pointer_push_constant;
binfo.pc_offset = parse_data->constants[inst->words[5]];
binfo.size = parse_data->constants[inst->words[7]];
parse_data->binary->set_printf_buffer_info(binfo);
parse_data->binary->add_push_constant(
pushconstant::printf_buffer_pointer, {binfo.pc_offset, 8u});
break;
}
default:
return SPV_ERROR_INVALID_DATA;
}
}
break;
default:
break;
}
return SPV_SUCCESS;
}
bool spir_binary::load(const char* fname) {
std::ifstream ifile;
ifile.open(fname, std::ios::in | std::ios::binary);
if (!ifile.is_open()) {
cvk_error("Failed to open %s", fname);
return false;
}
ifile.seekg(0, std::ios::end);
uint32_t size = ifile.tellg();
ifile.seekg(0, std::ios::beg);
return load(ifile, size);
}
bool spir_binary::load(std::istream& istream, uint32_t size) {
m_code.assign(size / SPIR_WORD_SIZE, 0);
istream.read(reinterpret_cast<char*>(m_code.data()), size);
if (!istream.good()) {
cvk_warn("Failed to load SPIR-V (size: %u)", size);
return false;
}
return true;
}
bool spir_binary::read(const unsigned char* src, size_t size) {
m_loaded_from_binary = true;
membuf bufview(src, src + size);
std::istream istream(&bufview);
return load(istream, size);
}
bool spir_binary::save(std::ostream& ostream) const {
ostream.write(reinterpret_cast<const char*>(m_code.data()), size());
return ostream.good();
}
bool spir_binary::save(const char* fname) const {
std::ofstream ofile;
ofile.open(fname, std::ios::out | std::ios::binary);
if (!ofile.is_open()) {
return false;
}
return save(ofile);
}
size_t spir_binary::size() const { return m_code.size() * SPIR_WORD_SIZE; }
bool spir_binary::write(unsigned char* dst) const {
membuf bufview(dst, dst + size());
std::ostream ostream(&bufview);
return save(ostream);
}
void spir_binary::use(std::vector<uint32_t>&& src) { m_code = std::move(src); }
void spir_binary::set_target_env(spv_target_env env) {
spvContextDestroy(m_context);
m_context = spvContextCreate(env);
}
bool spir_binary::validate(const spirv_validation_options& val_options) const {
spv_diagnostic diag;
spv_validator_options options = spvValidatorOptionsCreate();
spvValidatorOptionsSetUniformBufferStandardLayout(
options, val_options.uniform_buffer_std_layout);
spv_const_binary_t binary{m_code.data(), m_code.size()};
spv_result_t res =
spvValidateWithOptions(m_context, options, &binary, &diag);
spvDiagnosticPrint(diag);
spvDiagnosticDestroy(diag);
spvValidatorOptionsDestroy(options);
return res == SPV_SUCCESS;
}
bool spir_binary::strip_reflection(std::vector<uint32_t>* stripped) {
const spvtools::MessageConsumer consumer =
[](spv_message_level_t level, const char*,
const spv_position_t& position, const char* message) {
#define msgtpl "spvtools says '%s' at position %zu"
switch (level) {
case SPV_MSG_FATAL:
case SPV_MSG_INTERNAL_ERROR:
case SPV_MSG_ERROR:
cvk_error(msgtpl, message, position.index);
break;
case SPV_MSG_WARNING:
cvk_warn(msgtpl, message, position.index);
break;
case SPV_MSG_INFO:
cvk_info(msgtpl, message, position.index);
break;
case SPV_MSG_DEBUG:
cvk_debug(msgtpl, message, position.index);
break;
}
#undef msgtpl
};
#if COMPILER_AVAILABLE || USING_SWIFTSHADER
spvtools::Optimizer opt(m_target_env);
opt.SetMessageConsumer(consumer);
opt.RegisterPass(spvtools::CreateStripReflectInfoPass());
spvtools::OptimizerOptions options;
options.set_run_validator(false);
if (!opt.Run(m_code.data(), m_code.size(), stripped, options)) {
return false;
}
#else
*stripped = m_code;
#endif
return true;
}
bool spir_binary::load_descriptor_map() {
reflection_parse_data parse_data;
parse_data.binary = this;
// TODO: The parser assumes a valid SPIR-V module, but validation is not
// run until later.
auto result =
spvBinaryParse(m_context, &parse_data, m_code.data(), m_code.size(),
nullptr, parse_reflection, nullptr);
if (result != SPV_SUCCESS) {
cvk_error_fn("Parsing SPIR-V module reflection failed: %d", result);
return false;
}
return true;
}
bool spir_binary::get_capabilities(
std::vector<spv::Capability>& capabilities) const {
// Callback for receiving parsed instructions.
// The `user_data` parameter will be a pointer to the vector of
// capabilities (we cannot use a lambda capture for this as it prevents the
// lambda from being able to be converted to a function pointer).
auto parse_inst = [](void* user_data,
const spv_parsed_instruction_t* inst) {
// Stop parsing at first instruction that is not an OpCapability.
if (inst->opcode != spv::Op::OpCapability) {
return SPV_END_OF_STREAM;
}
// Add the capability to the list.
uint32_t capability = inst->words[inst->operands[0].offset];
auto capabilities =
reinterpret_cast<std::vector<spv::Capability>*>(user_data);
capabilities->push_back(static_cast<spv::Capability>(capability));
return SPV_SUCCESS;
};
// Parse the SPIR-V binary to build the list of required capabilities.
spv_result_t result =
spvBinaryParse(m_context, &capabilities, m_code.data(), m_code.size(),
nullptr, parse_inst, nullptr);
if (result != SPV_SUCCESS && result != SPV_END_OF_STREAM) {
cvk_error_fn("Parsing SPIR-V module failed: %d", result);
return false;
}
return true;
}
namespace {
#if COMPILER_AVAILABLE
bool save_cstring_to_file(
const std::string& fname, const char* data, size_t size,
std::ios_base::openmode open_mode = std::ios_base::out) {
std::filesystem::path ofname(fname);
if (ofname.has_parent_path()) {
std::error_code error;
std::filesystem::create_directories(ofname.parent_path(), error);
if (error) {
cvk_error_fn("create_directories failed: '%s' (%u)",
error.message().c_str(), error.value());
return false;
}
}
std::ofstream ofile{fname, open_mode};
if (!ofile.is_open()) {
return false;
}
ofile.write(data, size);
ofile.close();
return ofile.good();
}
bool save_string_to_file(const std::string& fname, const std::string& text) {
return save_cstring_to_file(fname, text.c_str(), text.size());
}
#ifndef CLSPV_ONLINE_COMPILER
bool save_il_to_file(const std::string& fname, const std::vector<uint8_t>& il) {
return save_cstring_to_file(fname, reinterpret_cast<const char*>(il.data()),
il.size(), std::ios::binary);
}
#endif // CLSPV_ONLINE_COMPILER
#endif // COMPILER_AVAILABLE
struct temp_folder_deletion {
temp_folder_deletion(const std::string& path) : m_path(path) {}
~temp_folder_deletion() {
if (!config.keep_temporaries && !m_path.empty())
std::filesystem::remove_all(m_path.c_str());
}
private:
std::string m_path;
};
enum class spirv_validation_level
{
skip,
warn,
error,
};
bool validate_binary(spir_binary const& binary,
spirv_validation_options const& val_options) {
spirv_validation_level level = spirv_validation_level::error;
if (config.spirv_validation.set) {
if (config.spirv_validation == 0) {
level = spirv_validation_level::skip;
} else if (config.spirv_validation == 1) {
level = spirv_validation_level::warn;
}
}
if (level == spirv_validation_level::skip) {
cvk_info("Skipping validation of SPIR-V binary.");
return true;
}
if (binary.validate(val_options)) {
cvk_info("SPIR-V binary is valid.");
return true;
}
if (level == spirv_validation_level::warn) {
cvk_warn("SPIR-V binary is invalid.");
return true;
}
cvk_error("SPIR-V binary is invalid.");
return false;
}
const uint32_t clvk_binary_magic =
0x6B766C63; // "clvk" in ASCII in little-endian
const uint32_t clvk_binary_version = 1;
struct clvk_binary_header {
uint32_t magic;
uint32_t version;
uint32_t binary_type;
};
#define COPY_WORD(dst, src) \
do { \
((unsigned char*)dst)[0] = ((unsigned char*)(src))[0]; \
((unsigned char*)dst)[1] = ((unsigned char*)(src))[1]; \
((unsigned char*)dst)[2] = ((unsigned char*)(src))[2]; \
((unsigned char*)dst)[3] = ((unsigned char*)(src))[3]; \
} while (0)
} // namespace
bool cvk_program::read_llvm_bitcode(const unsigned char* src, size_t size) {
if (size >= 4 &&
(src[0] == 'B' && src[1] == 'C' && src[2] == 0xc0 && src[3] == 0xde)) {
m_ir.resize(size);
memcpy(m_ir.data(), src, size);
m_dev_status[m_context->device()] = CL_BUILD_SUCCESS;
return true;
}
return false;
}
void cvk_program::write_binary_header(unsigned char* dst) const {
struct clvk_binary_header* header = (struct clvk_binary_header*)dst;
COPY_WORD(&header->magic, &clvk_binary_magic);
COPY_WORD(&header->version, &clvk_binary_version);
COPY_WORD(&header->binary_type, &m_binary_type);
}
cl_program_binary_type cvk_program::read_binary_header(const unsigned char* src,
size_t size) {
struct clvk_binary_header* header = (struct clvk_binary_header*)src;
if (size < sizeof(header)) {
return CL_PROGRAM_BINARY_TYPE_NONE;
}
uint32_t magic, version, binary_type;
COPY_WORD(&magic, &header->magic);
COPY_WORD(&version, &header->version);
COPY_WORD(&binary_type, &header->binary_type);
if (magic != clvk_binary_magic) {
cvk_info_fn("magic not found");
return CL_PROGRAM_BINARY_TYPE_NONE;
}
if (version != clvk_binary_version) {
cvk_warn_fn("wrong version");
return CL_PROGRAM_BINARY_TYPE_NONE;
}
return binary_type;
}
bool cvk_program::read(const unsigned char* src, size_t size) {
bool success = false;
auto binary_type = read_binary_header(src, size);
// if the binary does not have a clvk binary header, let's try to read
// it first as a llvm ir buffer, then as a vulkan spirv buffer.
if (binary_type == CL_PROGRAM_BINARY_TYPE_NONE) {
cvk_info_fn("no clvk binary header found, looking for llvm bitcode");
if (read_llvm_bitcode(src, size)) {
m_binary_type = CL_PROGRAM_BINARY_TYPE_COMPILED_OBJECT;
cvk_info_fn("llvm bitcode compiled object found");
return true;
}
cvk_info_fn("llvm bitcode not found, looking for Vulkan SPIR-V binary");
if (m_binary.read(src, size)) {
m_binary_type = CL_PROGRAM_BINARY_TYPE_EXECUTABLE;
cvk_info_fn("Vulkan SPIR-V binary executable found");
return true;
}
cvk_error_fn("unable to read binary");
return success;
}
auto header_size = sizeof(struct clvk_binary_header);
src += header_size;
size -= header_size;
switch (binary_type) {
case CL_PROGRAM_BINARY_TYPE_LIBRARY:
case CL_PROGRAM_BINARY_TYPE_COMPILED_OBJECT:
success = read_llvm_bitcode(src, size);
break;
case CL_PROGRAM_BINARY_TYPE_EXECUTABLE:
success = m_binary.read(src, size);
break;
}
if (success) {
m_binary_type = binary_type;
}
return success;
}
bool cvk_program::write(unsigned char* dst) const {
write_binary_header(dst);
dst += sizeof(struct clvk_binary_header);
switch (m_binary_type) {
case CL_PROGRAM_BINARY_TYPE_LIBRARY:
case CL_PROGRAM_BINARY_TYPE_COMPILED_OBJECT:
memcpy(dst, m_ir.data(), m_ir.size());
return true;
case CL_PROGRAM_BINARY_TYPE_EXECUTABLE:
return m_binary.write(dst);
}
return false;
}
size_t cvk_program::binary_size() const {
auto header_size = sizeof(struct clvk_binary_header);
switch (m_binary_type) {
case CL_PROGRAM_BINARY_TYPE_LIBRARY:
case CL_PROGRAM_BINARY_TYPE_COMPILED_OBJECT:
return header_size + m_ir.size();
case CL_PROGRAM_BINARY_TYPE_EXECUTABLE:
return header_size + m_binary.size();
}
return 0;
}
std::string cvk_program::prepare_build_options(const cvk_device* device) const {
// Strip off a few options we can't handle
std::string options;
if (m_build_options.size() > 0) {
options += " ";
options += m_build_options;
}
std::vector<std::pair<std::string, std::string>> option_substitutions = {
// FIXME The 1.2 conformance tests shouldn't pass this option.
// It doesn't exist after OpenCL 1.0.
{"-cl-strict-aliasing", ""},
// clspv require entrypoint inlining for OpenCL 2.0 and OpenCL 3.0 (for
// generic addrspace for example).
{"-cl-std=CL2.0", "-cl-std=CL2.0 -inline-entry-points"},
{"-cl-std=CL3.0", "-cl-std=CL3.0 -inline-entry-points"},
{"-create-library", ""},
};
for (auto& subst : option_substitutions) {
size_t loc = options.find(subst.first);
if (loc != std::string::npos) {
options.replace(loc, subst.first.length(), subst.second);
}
}
// Check for some options we need and add them if not present.
std::string necessary_options[] = {
"-cl-single-precision-constant",
"-cl-kernel-arg-info",
};
for (std::string option : necessary_options) {
if (options.find(option) == std::string::npos) {
options += " " + option + " ";
}
}
// The device sets up some compiler options based on its capabilities.
options += " " + device->get_device_specific_compile_options() + " ";
// Features
if (options.find("-cl-std=CL3.0") != std::string::npos) {
auto features = device->opencl_c_features();
if (!features.empty()) {
options += "-enable-feature-macros=";
for (auto& feature : features) {
options += feature.name;
options += ',';
}
options.back() = ' '; // replace the final comma
}
}
auto buff_size = m_context->get_printf_buffersize();
options += " -enable-printf ";
options += " -printf-buffer-size=" + std::to_string(buff_size) + " ";
#if COMPILER_AVAILABLE
options += " " + config.clspv_options() + " ";
#endif
// split options into a vector
std::istringstream iss(options);
std::vector<std::string> vector_options;
std::string token;
while (std::getline(iss, token, ' ')) {
vector_options.push_back(token);
}
// loop through the options and quote the ones that need it
std::string quoted_options;
for (size_t i = 0; i < vector_options.size(); i++) {
if (vector_options[i].empty()) {
continue;
}
if (vector_options[i].find("-") == 0) {
quoted_options += vector_options[i];
} else {
quoted_options += "\"" + vector_options[i] + "\"";
}
quoted_options += " ";
}
return quoted_options;
}
cl_int cvk_program::parse_user_spec_constants() {
#if COMPILER_AVAILABLE && ENABLE_SPIRV_IL
#ifndef CLSPV_ONLINE_COMPILER
// We'll need to go through the whole temp folder rigamarole to query the
// spec constant info with the command line tool.
std::filesystem::path tmp_prefix(config.compiler_temp_dir());
std::filesystem::path tmp_suffix("clvk-XXXXXX");
std::string tmp_template = (tmp_prefix / tmp_suffix).string();
const char* tmp = cvk_mkdtemp(tmp_template);
if (tmp == nullptr) {
cvk_error_fn("Could not create temporary folder \"%s\"",
tmp_template.c_str());
return CL_INVALID_VALUE;
}
std::string tmp_folder = tmp;
cvk_info("Created temporary folder \"%s\"", tmp_folder.c_str());
std::string llvmspirv_input_file = tmp_folder + "/source.spv";
if (!save_il_to_file(llvmspirv_input_file, m_il)) {
cvk_error_fn("Couldn't save IL to file!");
return CL_INVALID_VALUE;
}
std::string cmd_spv{config.llvmspirv_bin()};
cmd_spv += " --spec-const-info ";
cmd_spv += llvmspirv_input_file;
std::string output = "";
cvk_exec(cmd_spv, &output);
auto output_stream = std::istringstream(output);
for (std::string line; std::getline(output_stream, line);) {
if (line.find("Spec const id") == std::string::npos) {
continue;
}
auto delim = line.find(",");
auto id_string = line.substr(0, delim);
line.erase(0, delim + 1);
id_string = id_string.substr(id_string.find("=") + 1);
auto id = static_cast<uint32_t>(atoi(id_string.c_str()));
delim = line.find(",");
auto size_string = line.substr(0, delim);
line.erase(0, delim + 1);
size_string = size_string.substr(size_string.find("=") + 1);
auto size = static_cast<uint32_t>(atoi(size_string.c_str()));
// We need to consume the "=" and the " " here since we aren't
// converting the value from a string.
auto type_string = line.substr(line.find("=") + 2);
m_user_spec_constants.emplace(
id, user_spec_constant_data{type_string, size});
}
std::filesystem::remove_all(tmp_folder.c_str());
return CL_SUCCESS;
#else
auto m_il_start = reinterpret_cast<const unsigned char*>(m_il.data());
membuf m_il_buf(m_il_start, m_il_start + m_il.size());
std::istream m_il_stream(&m_il_buf);
SPIRV::TranslatorOpts translator_opts;
// We don't need llvm-spirv to validate extensions for us.
translator_opts.enableAllExtensions();
static std::mutex llvmspirv_compile_mutex;
std::lock_guard<std::mutex> llvmspirv_compile_lock(llvmspirv_compile_mutex);
std::vector<llvm::SpecConstInfoTy> spec_const_info;
if (!getSpecConstInfo(m_il_stream, spec_const_info)) {
cvk_error_fn("Failed to parse spec constants");
return CL_INVALID_VALUE;
}
for (const auto& spec_const : spec_const_info) {
m_user_spec_constants.emplace(
spec_const.ID,
user_spec_constant_data{spec_const.Type, spec_const.Size});
}
return CL_SUCCESS;
#endif // CLSPV_ONLINE_COMPILER
#else
#if !COMPILER_AVAILABLE
cvk_error_fn("Could not parse user spec constants because clvk has been "
"built with CLVK_COMPILER_AVAILABLE=OFF");
#elif !ENABLE_SPIRV_IL
cvk_error_fn("Could not parse user spec constants because clvk has been "
"built with CLVK_ENABLE_SPIRV_IL=OFF");
#endif
return CL_INVALID_OPERATION;
#endif // COMPILER_AVAILABLE
}
#if COMPILER_AVAILABLE
#ifndef CLSPV_ONLINE_COMPILER
cl_build_status cvk_program::do_build_inner_offline(bool build_to_ir,
bool build_from_il,
std::string& build_options,
std::string& tmp_folder) {
TRACE_FUNCTION("build_to_ir", build_to_ir, "build_from_il", build_from_il,
"build_options", TRACE_STRING(build_options.c_str()));
// Compose clspv command-line
std::string cmd{config.clspv_path};
cmd += " ";
std::string clspv_input_file{tmp_folder + "/source"};
// Save input program to a file
if (build_from_il) {
#ifndef ENABLE_SPIRV_IL
cvk_error_fn("Could not build from il because clvk has been built with "
"CLVK_ENABLE_SPIRV_IL=OFF");
return CL_BUILD_ERROR;
#else // ENABLE_SPIRV_IL
std::string llvmspirv_input_file{tmp_folder + "/source.spv"};
clspv_input_file += ".bc";
if (!save_il_to_file(llvmspirv_input_file, m_il)) {
cvk_error_fn("Couldn't save IL to file!");
return CL_BUILD_ERROR;
}
// Compose llvm-spirv command-line
std::string cmd_spv{config.llvmspirv_bin};
if (!m_user_spec_constants.empty()) {
std::string spec_constant_flag = " --spec-const=";
for (const auto& spec_const : m_user_spec_constants) {
if (!spec_const.second.set) {
continue;
}
const auto& type_string = spec_const.second.type;
spec_constant_flag +=
std::to_string(spec_const.first) + ":" + type_string + ":";
if (type_string.find("i32") != std::string::npos) {
spec_constant_flag +=
std::to_string(spec_const.second.data.i32);
} else if (type_string.find("i16") != std::string::npos) {
spec_constant_flag +=
std::to_string(spec_const.second.data.i16);
} else if (type_string.find("i8") != std::string::npos ||
type_string.find("i1") != std::string::npos) {
spec_constant_flag +=
std::to_string(spec_const.second.data.i8);
} else if (type_string.find("i64") != std::string::npos) {
spec_constant_flag +=
std::to_string(spec_const.second.data.i64);
} else if (type_string.find("f16") != std::string::npos) {
// At most we should need seven (0xFFFF\0)
char buf[7];
std::snprintf(buf, 7, "0x%04X", spec_const.second.data.i16);
spec_constant_flag += std::string(buf);
} else if (type_string.find("f32") != std::string::npos) {
float spec_value;
std::memcpy(&spec_value, &spec_const.second.data.i32,
spec_const.second.size);
spec_constant_flag += std::to_string(spec_value);
} else if (type_string.find("f64") != std::string::npos) {
double spec_value;
std::memcpy(&spec_value, &spec_const.second.data.i64,
spec_const.second.size);
spec_constant_flag += std::to_string(spec_value);
}
spec_constant_flag += " ";
}
cmd_spv += spec_constant_flag;
}
cmd_spv += " -r ";
cmd_spv += " -o ";
cmd_spv += clspv_input_file;
cmd_spv += " ";
cmd_spv += llvmspirv_input_file;
// Call the translator
int status = cvk_exec(cmd_spv);
if (status != 0) {
cvk_error_fn("failed to translate SPIR-V to LLVM IR");
return CL_BUILD_ERROR;
}
cmd += clspv_input_file;
cmd += " ";
#endif // ENABLE_SPIRV_IL
} else if (m_operation == build_operation::link) {
for (auto input_program : m_input_programs) {
if (input_program->m_binary_type !=
CL_PROGRAM_BINARY_TYPE_COMPILED_OBJECT &&
input_program->m_binary_type !=
CL_PROGRAM_BINARY_TYPE_LIBRARY) {
return CL_BUILD_ERROR;
}
std::string input_file = clspv_input_file + "_" +
std::to_string((uintptr_t)input_program) +
".bc";
if (!save_il_to_file(input_file, input_program->m_ir)) {
cvk_error_fn("Couldn't save source to file!");
return CL_BUILD_ERROR;
}
cmd += input_file;
cmd += " ";
}
} else {
if (m_source.empty() && !m_ir.empty()) {
clspv_input_file += ".bc";
if (!save_il_to_file(clspv_input_file, m_ir)) {
cvk_error_fn("Couldn't save source to file!");
return CL_BUILD_ERROR;
}
cmd += clspv_input_file;
cmd += " ";
} else {
clspv_input_file += ".cl";
if (!save_string_to_file(clspv_input_file, m_source)) {
cvk_error_fn("Couldn't save source to file!");
return CL_BUILD_ERROR;
}
cmd += clspv_input_file;
cmd += " ";
}
}
std::string clspv_output_file{tmp_folder + "/compiled"};
if (build_to_ir) {
clspv_output_file += ".bc";
} else {
clspv_output_file += ".spv";
}
cmd += build_options;
cmd += " -o ";
cmd += clspv_output_file;
// Call clspv
int status = cvk_exec(cmd, &m_build_log);
if (status != 0) {
cvk_error_fn("failed to compile the program");
cvk_debug_fn("%s", m_build_log.c_str());
return CL_BUILD_ERROR;
}
// Load output from clspv
if (build_to_ir) {
std::ifstream stream(clspv_output_file,
std::ios::in | std::ios::binary);
m_ir.assign((std::istreambuf_iterator<char>(stream)),
std::istreambuf_iterator<char>());
if (!stream.good()) {
return CL_BUILD_ERROR;
}
} else {
const char* filename = clspv_output_file.c_str();
if (!m_binary.load(filename)) {
cvk_error("Could not load SPIR-V binary from \"%s\"", filename);
return CL_BUILD_ERROR;
}
}
cvk_info("Loaded %s binary from \"%s\", size = %zu %s",
build_to_ir ? "IR" : "SPIR-V", clspv_output_file.c_str(),
build_to_ir ? m_ir.size() : m_binary.code().size(),
build_to_ir ? "bytes" : "words");
return CL_BUILD_SUCCESS;
}
#else // #ifndef CLSPV_ONLINE_COMPILER
cl_build_status cvk_program::do_build_inner_online(bool build_to_ir,
bool build_from_il,
std::string& build_options) {
TRACE_FUNCTION("build_to_ir", build_to_ir, "build_from_il", build_from_il,
"build_options", TRACE_STRING(build_options.c_str()));
cvk_info_fn("build_from_il %u - build_to_ir %u", build_from_il,
build_to_ir);
if (build_from_il) {
#ifndef ENABLE_SPIRV_IL
cvk_error_fn("Could not build from il because clvk has been built with "
"CLVK_ENABLE_SPIRV_IL=OFF");
return CL_BUILD_ERROR;
#else // ENABLE_SPIRV_IL
llvm::LLVMContext llvm_context;
llvm::Module* llvm_module;
std::string err;
auto m_il_start = reinterpret_cast<const unsigned char*>(m_il.data());
membuf m_il_buf(m_il_start, m_il_start + m_il.size());
std::istream m_il_stream(&m_il_buf);
SPIRV::TranslatorOpts translator_opts;
// We don't need llvm-spirv to validate extensions for us.
translator_opts.enableAllExtensions();
for (const auto& spec_const : m_user_spec_constants) {
if (!spec_const.second.set) {
continue;
}
uint64_t spec_const_value = 0;
if (spec_const.second.type == "i1" ||
spec_const.second.type == "i8") {
spec_const_value = spec_const.second.data.i8;
} else if (spec_const.second.type == "i16" ||
spec_const.second.type == "f16") {
spec_const_value = spec_const.second.data.i16;
} else if (spec_const.second.type == "i32" ||
spec_const.second.type == "f32") {
spec_const_value = spec_const.second.data.i32;
} else if (spec_const.second.type == "i64" ||
spec_const.second.type == "f64") {
spec_const_value = spec_const.second.data.i64;
}
translator_opts.setSpecConst(spec_const.first, spec_const_value);
}
// llvm-spirv is based on LLVM. LLVM options parsing is done using
// global variable that are not thread safe. Thus, we need to lock call
// to llvm-spirv in order to ensure a thread safe execution.
{
static std::mutex llvmspirv_compile_mutex;
std::lock_guard<std::mutex> llvmspirv_compile_lock(
llvmspirv_compile_mutex);
std::vector<const char*> llvmArgv{
"llvm-spirv(online)",
};
llvm::cl::ResetAllOptionOccurrences();
llvm::cl::ParseCommandLineOptions(llvmArgv.size(), llvmArgv.data());
if (!llvm::readSpirv(llvm_context, translator_opts, m_il_stream,
llvm_module, err)) {
cvk_error_fn("Fails to load SPIR-V as LLVM Module: %s",
err.c_str());
return CL_BUILD_ERROR;
}
}
m_source.clear();
llvm::raw_string_ostream spirv_stream(m_source);
llvm::WriteBitcodeToFile(*llvm_module, spirv_stream);
#endif // ENABLE_SPIRV_IL
}
cvk_info("About to compile \"%s\"", build_options.c_str());
int status;
// clspv is based on LLVM. LLVM options parsing is done using global
// variable that are not thread safe. Thus, we need to lock call to clspv in
// order to ensure a thread safe execution.
{
static std::mutex clspv_compile_mutex;
std::lock_guard<std::mutex> clspv_compile_lock(clspv_compile_mutex);
std::vector<std::string> programs;
if (m_operation == build_operation::link) {
for (auto input_program : m_input_programs) {
if (input_program->m_binary_type !=
CL_PROGRAM_BINARY_TYPE_COMPILED_OBJECT &&
input_program->m_binary_type !=
CL_PROGRAM_BINARY_TYPE_LIBRARY) {
return CL_BUILD_ERROR;
}
programs.emplace_back(std::string{input_program->m_ir.begin(),
input_program->m_ir.end()});
}
} else {
if (m_source.empty() && !m_ir.empty()) {
programs.emplace_back(std::string{m_ir.begin(), m_ir.end()});
} else {
programs.emplace_back(m_source);
}
}
m_build_log.clear();
if (build_to_ir) {
std::vector<uint32_t> ir;
status = clspv::CompileFromSourcesString(programs, build_options,
&ir, &m_build_log);
m_ir.clear();
auto size = ir.size() * sizeof(uint32_t);
m_ir.resize(size);
memcpy(m_ir.data(), ir.data(), size);
} else {
status = clspv::CompileFromSourcesString(
programs, build_options, m_binary.raw_binary(), &m_build_log);
}
}
if (status != 0) {
cvk_error_fn("failed to compile the program");
cvk_debug_fn("%s", m_build_log.c_str());
return CL_BUILD_ERROR;
}
return CL_BUILD_SUCCESS;
}
#endif // #ifndef CLSPV_ONLINE_COMPILER
#endif // #if COMPILER_AVAILABLE
cl_build_status cvk_program::do_build_inner(const cvk_device* device) {
#if !COMPILER_AVAILABLE
UNUSED(device);
#else
// m_il and m_source cannot be set together.
// They can both be empty when loading from binary.
CVK_ASSERT(m_il.empty() || m_source.empty());
bool build_from_il =
m_il.size() > 0 && m_operation != build_operation::link;
bool use_tmp_folder = true;
#ifdef CLSPV_ONLINE_COMPILER
use_tmp_folder =
m_operation == build_operation::compile && m_num_input_programs > 0;
#endif
std::string tmp_folder;
if (use_tmp_folder) {
// Create temporary folder
std::filesystem::path tmp_prefix(config.compiler_temp_dir());
std::filesystem::path tmp_suffix("clvk-XXXXXX");
std::string tmp_template = (tmp_prefix / tmp_suffix).string();
const char* tmp = cvk_mkdtemp(tmp_template);
if (tmp == nullptr) {
cvk_error_fn("Could not create temporary folder \"%s\"",
tmp_template.c_str());
return CL_BUILD_ERROR;
}
tmp_folder = tmp;
cvk_info("Created temporary folder \"%s\"", tmp_folder.c_str());
}
temp_folder_deletion temp(tmp_folder);
// Prepare build options
bool create_library =
m_build_options.find("create-library") != std::string::npos;
auto build_options = prepare_build_options(device);
// Add options to specify input/output types
if (m_source.empty() || m_operation == build_operation::link) {
build_options += " -x ir ";
}
bool build_to_ir =
m_operation == build_operation::compile || create_library;
if (build_to_ir) {
build_options += " --output-format=bc ";
}
// Save headers
if (use_tmp_folder && m_operation == build_operation::compile) {
build_options += "-I" + tmp_folder;
for (cl_uint i = 0; i < m_num_input_programs; i++) {
std::string fname{tmp_folder + "/" + m_header_include_names[i]};
if (!save_string_to_file(fname, m_input_programs[i]->source())) {
cvk_error_fn("Couldn't save header to file!");
return CL_BUILD_ERROR;
}
}
}
cl_build_status build_status;
#ifdef CLSPV_ONLINE_COMPILER
build_status =
do_build_inner_online(build_to_ir, build_from_il, build_options);
#else
build_status = do_build_inner_offline(build_to_ir, build_from_il,
build_options, tmp_folder);
#endif // CLSPV_ONLINE_COMPILER
if (build_status != CL_BUILD_SUCCESS) {
return build_status;
}
// Select operation
if (m_operation == build_operation::compile) {
m_binary_type = CL_PROGRAM_BINARY_TYPE_COMPILED_OBJECT;
} else if (m_operation == build_operation::link && create_library) {
m_binary_type = CL_PROGRAM_BINARY_TYPE_LIBRARY;
} else {
m_binary_type = CL_PROGRAM_BINARY_TYPE_EXECUTABLE;
}
#endif // #if !COMPILER_AVAILABLE
return CL_BUILD_SUCCESS;
}
void cvk_program::prepare_push_constant_range() {
auto& pcs = m_binary.push_constants();
uint32_t min_offset = UINT32_MAX;
uint32_t max_offset = 0, max_offset_size = 0;
for (auto& pc_pcd : pcs) {
auto pcd = pc_pcd.second;
min_offset = std::min(min_offset, pcd.offset);
if (pcd.offset >= max_offset) {
max_offset = pcd.offset;
max_offset_size = pcd.size;
}
}
m_push_constant_range = {VK_SHADER_STAGE_COMPUTE_BIT, min_offset,
max_offset + max_offset_size};
}
bool cvk_program::check_capabilities(const cvk_device* device) {
// Get list of required SPIR-V capabilities.
std::vector<spv::Capability> capabilities;
if (!m_binary.get_capabilities(capabilities)) {
cvk_error("Failed to get required SPIR-V capabilities.");
return false;
}
// Check that each required capability is supported by the device.
for (auto c : capabilities) {
cvk_info_fn("Program requires SPIR-V capability %d (%s).", c,
spirv_capability_to_string(c));
if (!device->supports_capability(c)
#ifdef CLVK_UNIT_TESTING_ENABLED
|| config.force_check_capabilities_error()
#endif
) {
std::stringstream error_message;
error_message << "Device does not support SPIR-V capability " << c
<< " (" << spirv_capability_to_string(c) << ").";
cvk_error_fn("%s", error_message.str().c_str());
m_build_log += error_message.str();
return false;
}
}
return true;
}
void cvk_program::do_build() {
// Destroy entry points from previous build
m_entry_points.clear();
auto device = m_context->device();
if (m_operation != build_operation::build_binary) {
cl_build_status status = do_build_inner(device);
if ((m_binary_type != CL_PROGRAM_BINARY_TYPE_EXECUTABLE) ||
(status != CL_BUILD_SUCCESS)) {
complete_operation(device, status);
return;
}
}
// Load descriptor map
if (!m_binary.load_descriptor_map()) {
cvk_error("Could not load descriptor map for SPIR-V binary.");
complete_operation(device, CL_BUILD_ERROR);
return;
}
if (!create_module_constant_data_buffer()) {
complete_operation(device, CL_BUILD_ERROR);
return;
}
prepare_push_constant_range();
bool cache_hit =
device->get_pipeline_cache(m_binary.code(), m_pipeline_cache);
if (m_pipeline_cache == VK_NULL_HANDLE) {
complete_operation(device, CL_BUILD_ERROR);
return;
}
if (!cache_hit) {
// Validate
// TODO validate with different rules depending on the binary type
if (m_binary_type == CL_PROGRAM_BINARY_TYPE_EXECUTABLE) {
spirv_validation_options validation_options{};
validation_options.uniform_buffer_std_layout =
m_context->device()->supports_ubo_stdlayout();
if (!validate_binary(m_binary, validation_options)) {
complete_operation(device, CL_BUILD_ERROR);
return;
}
}
}
// Check capabilities against the device.
if ((m_binary_type == CL_PROGRAM_BINARY_TYPE_EXECUTABLE) &&
!config.skip_spirv_capability_check && !check_capabilities(device)) {
cvk_error("Missing support for required SPIR-V capabilities.");
complete_operation(device, CL_BUILD_ERROR);
return;
}
// Create literal samplers
for (auto const& desc : literal_sampler_descs()) {
auto sampler =
cvk_sampler::create(context(), desc.normalized_coords,
desc.addressing_mode, desc.filter_mode);
if (sampler == nullptr) {
complete_operation(device, CL_BUILD_ERROR);
return;
}
m_literal_samplers.emplace_back(sampler);
}
// Strip the reflection information if non-semantic info is not supported
// by the Vulkan implementation. This stripped binary is stored separately
// from |m_binary| because clvk needs to be able to provide the binary with
// reflection information for clGetProgramInfo.
const uint32_t* spir_data = m_binary.spir_data();
size_t spir_size = m_binary.spir_size();
const bool should_strip_reflection =
!device->is_vulkan_extension_enabled(
VK_KHR_SHADER_NON_SEMANTIC_INFO_EXTENSION_NAME)
#ifdef USING_SWIFTSHADER
|| true
#endif
;
if (should_strip_reflection) {
if (!m_binary.strip_reflection(&m_stripped_binary)) {
cvk_error_fn("couldn't strip reflection from SPIR-V module");
complete_operation(device, CL_BUILD_ERROR);
return;
}
spir_data = m_stripped_binary.data();
spir_size = m_stripped_binary.size() * sizeof(uint32_t);
}
// Create a shader module
VkDevice dev = device->vulkan_device();
VkShaderModuleCreateInfo moduleCreateInfo = {
VK_STRUCTURE_TYPE_SHADER_MODULE_CREATE_INFO, // sType
nullptr, // pNext
0, // flags
spir_size, // codeSize
spir_data // pCode
};
VkResult res =
vkCreateShaderModule(dev, &moduleCreateInfo, nullptr, &m_shader_module);
if (res != VK_SUCCESS) {
cvk_error("vkCreateShaderModule returned %d", res);
complete_operation(device, CL_BUILD_ERROR);
return;
}
complete_operation(device, CL_BUILD_SUCCESS);
}
cl_int cvk_program::build(build_operation operation, cl_uint num_devices,
const cl_device_id* device_list, const char* options,
cl_uint num_input_programs,
const cl_program* input_programs,
const char** header_include_names,
cvk_program_callback cb, void* data) {
std::lock_guard<std::mutex> lock(m_lock);
// Check if there is already a build in progress
// TODO: Allow concurrent builds targeting different devices
if (std::count_if(m_dev_status.begin(), m_dev_status.end(),
[](auto& status) {
return status.second == CL_BUILD_IN_PROGRESS;
})) {
return CL_INVALID_OPERATION;
}
retain();
for (cl_uint i = 0; i < num_input_programs; i++) {
cvk_program* iprog =
const_cast<cvk_program*>(icd_downcast(input_programs[i]));
iprog->retain();
m_input_programs.push_back(iprog);
if (header_include_names != nullptr) {
m_header_include_names.push_back(header_include_names[i]);
}
}
// Mark build in-progress and save devices
if (num_devices == 0) {
m_num_devices = 1u;
m_dev_status[m_context->device()] = CL_BUILD_IN_PROGRESS;
} else {
m_num_devices = num_devices;
for (cl_uint i = 0; i < num_devices; i++) {
m_dev_status[icd_downcast(device_list[i])] = CL_BUILD_IN_PROGRESS;
}
}
if (options != nullptr) {
m_build_options = options;
}
m_num_input_programs = num_input_programs;
m_operation = operation;
m_operation_callback = cb;
m_operation_callback_data = data;
cl_int ret = CL_SUCCESS;
bool build_in_separate_thread = config.build_in_separate_thread() || cb;
bool wait_for_completion = !cb;
if (build_in_separate_thread) {
// Kick off build
m_thread = std::make_unique<std::thread>(
&cvk_program::do_build_in_separate_thread, this);
if (!wait_for_completion) {
m_thread->detach();
}
} else {
do_build();
}
if (wait_for_completion) {
if (build_in_separate_thread) {
CVK_ASSERT(m_thread->joinable());
m_thread->join();
}
if (build_status() != CL_BUILD_SUCCESS) {
switch (operation) {
case build_operation::link:
ret = CL_LINK_PROGRAM_FAILURE;
break;
case build_operation::build:
case build_operation::build_binary:
ret = CL_BUILD_PROGRAM_FAILURE;
break;
case build_operation::compile:
ret = CL_COMPILE_PROGRAM_FAILURE;
break;
}
CVK_ASSERT(ret != CL_SUCCESS);
}
}
return ret;
}
cvk_entry_point::cvk_entry_point(cvk_device* dev, cvk_program* program,
const std::string& name)
: m_device(dev), m_context(program->context()), m_program(program),
m_name(name), m_pod_descriptor_type(VK_DESCRIPTOR_TYPE_MAX_ENUM),
m_pod_buffer_size(0u), m_has_pod_arguments(false),
m_has_pod_buffer_arguments(false), m_sampler_metadata(nullptr),
m_image_metadata(nullptr), m_descriptor_pool(VK_NULL_HANDLE),
m_pipeline_layout(VK_NULL_HANDLE), m_nb_descriptor_set_allocated(0),
m_first_allocation_failure(true) {
TRACE_CNT_VAR_INIT(descriptor_set_allocated_counter,
"clvk-entry_point_" + std::to_string((uintptr_t)this),
dev->track());
TRACE_CNT(descriptor_set_allocated_counter, 0);
}
std::shared_ptr<cvk_entry_point>
cvk_program::get_entry_point(std::string& name, cl_int* errcode_ret) {
std::lock_guard<std::mutex> lock(m_lock);
// Check for existing entry point in cache
if (m_entry_points.count(name)) {
*errcode_ret = CL_SUCCESS;
return m_entry_points.at(name);
}
// Create and initialize entry point
std::shared_ptr<cvk_entry_point> entry_point =
std::make_shared<cvk_entry_point>(m_context->device(), this, name);
*errcode_ret = entry_point->init();
if (*errcode_ret != CL_SUCCESS) {
return nullptr;
}
// Add to cache for reuse by other kernels
m_entry_points.insert({name, entry_point});
return entry_point;
}
bool cvk_entry_point::build_descriptor_set_layout(
const std::vector<VkDescriptorSetLayoutBinding>& bindings) {
VkDescriptorSetLayoutCreateInfo createInfo = {
VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO, nullptr,
0, // flags
static_cast<uint32_t>(bindings.size()), // bindingCount
bindings.data() // pBindings
};
VkResult res;
if (bindings.size() > 0) {
VkDescriptorSetLayout setLayout;
res = vkCreateDescriptorSetLayout(m_device->vulkan_device(),
&createInfo, 0, &setLayout);
if (res != VK_SUCCESS) {
cvk_error("Could not create descriptor set layout");
return false;
}
m_descriptor_set_layouts.push_back(setLayout);
}
return true;
}
bool cvk_entry_point::build_descriptor_sets_layout_bindings_for_arguments(
binding_stat_map& smap, uint32_t& num_resource_slots) {
bool pod_found = false;
uint32_t highest_binding = 0;
std::vector<VkDescriptorSetLayoutBinding> layoutBindings;
for (auto& arg : m_args) {
VkDescriptorType dt = VK_DESCRIPTOR_TYPE_MAX_ENUM;
switch (arg.kind) {
case kernel_argument_kind::buffer:
dt = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER;
break;
case kernel_argument_kind::buffer_ubo:
dt = VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER;
break;
case kernel_argument_kind::sampled_image:
dt = VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE;
break;
case kernel_argument_kind::storage_image:
dt = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE;
break;
case kernel_argument_kind::storage_texel_buffer:
dt = VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER;
break;
case kernel_argument_kind::uniform_texel_buffer:
dt = VK_DESCRIPTOR_TYPE_UNIFORM_TEXEL_BUFFER;
break;
case kernel_argument_kind::sampler:
dt = VK_DESCRIPTOR_TYPE_SAMPLER;
break;
case kernel_argument_kind::local:
continue;
case kernel_argument_kind::pod:
case kernel_argument_kind::pod_ubo:
case kernel_argument_kind::pointer_ubo:
if (!pod_found) {
if (arg.kind == kernel_argument_kind::pod) {
dt = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER;
} else if (arg.kind == kernel_argument_kind::pod_ubo ||
arg.kind == kernel_argument_kind::pointer_ubo) {
dt = VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER;
}
m_pod_descriptor_type = dt;
pod_found = true;
} else {
continue;
}
break;
case kernel_argument_kind::pod_pushconstant:
case kernel_argument_kind::pointer_pushconstant:
case kernel_argument_kind::unused:
continue;
}
VkDescriptorSetLayoutBinding binding = {
arg.binding, // binding
dt, // descriptorType
1, // decriptorCount
VK_SHADER_STAGE_COMPUTE_BIT, // stageFlags
nullptr // pImmutableSamplers
};
highest_binding = std::max(arg.binding, highest_binding);
layoutBindings.push_back(binding);
smap[binding.descriptorType]++;
}
num_resource_slots = highest_binding + 1;
if (!build_descriptor_set_layout(layoutBindings)) {
return false;
}
return true;
}
bool cvk_entry_point::
build_descriptor_sets_layout_bindings_for_literal_samplers(
binding_stat_map& smap) {
std::vector<VkDescriptorSetLayoutBinding> layoutBindings;
for (auto& desc : m_program->literal_sampler_descs()) {
VkDescriptorSetLayoutBinding binding = {
desc.binding, // binding
VK_DESCRIPTOR_TYPE_SAMPLER, // descriptorType
1, // decriptorCount
VK_SHADER_STAGE_COMPUTE_BIT, // stageFlags
nullptr // pImmutableSamplers
};
layoutBindings.push_back(binding);
smap[binding.descriptorType]++;
}
if (!build_descriptor_set_layout(layoutBindings)) {
return false;
}
return true;
}
bool cvk_entry_point::
build_descriptor_sets_layout_bindings_for_program_scope_buffers(
binding_stat_map& smap) {
std::vector<VkDescriptorSetLayoutBinding> layoutBindings;
if (m_program->module_constant_data_buffer() != nullptr) {
auto info = m_program->module_constant_data_buffer_info();
// If the program scope buffer isn't passed as a storage buffer (i.e.
// it is passed a pointer push constant), there is nothing to bind here
if (info->type != module_buffer_type::storage_buffer) {
return true;
}
VkDescriptorSetLayoutBinding binding = {
info->binding, // binding
VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, // descriptorType
1, // decriptorCount
VK_SHADER_STAGE_COMPUTE_BIT, // stageFlags
nullptr // pImmutableSamplers
};
layoutBindings.push_back(binding);
smap[binding.descriptorType]++;
}
if (!build_descriptor_set_layout(layoutBindings)) {
return false;
}
return true;
}
bool cvk_entry_point::build_descriptor_sets_layout_bindings_for_printf_buffer(
binding_stat_map& smap) {
std::vector<VkDescriptorSetLayoutBinding> layoutBindings;
if (m_program->printf_buffer_info().size > 0 && (uses_printf())) {
auto info = m_program->printf_buffer_info();
if (info.type != module_buffer_type::storage_buffer) {
return true;
}
VkDescriptorSetLayoutBinding binding = {
info.binding, // binding
VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, // descriptorType
1, // decriptorCount
VK_SHADER_STAGE_COMPUTE_BIT, // stageFlags
nullptr // pImmutableSamplers
};
layoutBindings.push_back(binding);
smap[binding.descriptorType]++;
}
if (!build_descriptor_set_layout(layoutBindings)) {
return false;
}
return true;
}
cl_int cvk_entry_point::init() {
VkResult res;
// Get the image metadata for the this entry point
if (auto* md = m_program->image_metadata(m_name)) {
m_image_metadata = md;
}
// Get the sampler metadata for this entry point
if (auto* md = m_program->sampler_metadata(m_name)) {
m_sampler_metadata = md;
}
// Get a pointer to the arguments from the program
auto args = m_program->args_for_kernel(m_name);
if (args == nullptr) {
cvk_error("Kernel %s doesn't exist in program", m_name.c_str());
return CL_INVALID_KERNEL_NAME;
}
// Store a sorted copy of the arguments
m_args = *args;
std::sort(
m_args.begin(), m_args.end(),
[](kernel_argument a, kernel_argument b) { return a.pos < b.pos; });
// Create Descriptor Sets Layout
std::unordered_map<VkDescriptorType, uint32_t> bindingTypes;
if (!build_descriptor_sets_layout_bindings_for_literal_samplers(
bindingTypes)) {
return CL_INVALID_VALUE;
}
if (!build_descriptor_sets_layout_bindings_for_arguments(
bindingTypes, m_num_resource_slots)) {
return CL_INVALID_VALUE;
}
if (!build_descriptor_sets_layout_bindings_for_program_scope_buffers(
bindingTypes)) {
return CL_INVALID_VALUE;
}
if (!build_descriptor_sets_layout_bindings_for_printf_buffer(
bindingTypes)) {
return CL_INVALID_VALUE;
}
// Do we have POD arguments?
for (auto& arg : m_args) {
if (arg.is_pod()) {
m_has_pod_arguments = true;
if (arg.is_pod_buffer()) {
m_has_pod_buffer_arguments = true;
}
}
}
// Calculate POD buffer size and update the push constant range.
VkPushConstantRange push_constant_range = m_program->push_constant_range();
if (m_has_pod_arguments) {
// Check we know the POD buffer's descriptor type
if (m_has_pod_buffer_arguments &&
m_pod_descriptor_type == VK_DESCRIPTOR_TYPE_MAX_ENUM) {
return CL_INVALID_PROGRAM;
}
// Find how big the POD buffer should be
uint32_t max_offset = 0;
uint32_t max_offset_arg_size = 0;
for (auto& arg : m_args) {
if (arg.is_pod()) {
if (arg.offset >= max_offset) {
max_offset = arg.offset;
max_offset_arg_size = arg.size;
}
if (!arg.is_pod_buffer()) {
if (arg.offset < push_constant_range.offset) {
push_constant_range.offset = arg.offset;
}
if (arg.offset + arg.size >
push_constant_range.offset + push_constant_range.size) {
push_constant_range.size =
arg.offset + arg.size - push_constant_range.offset;
}
}
}
}
m_pod_buffer_size = max_offset + max_offset_arg_size;
m_pod_buffer_size = round_up(m_pod_buffer_size, 4);
}
// Take the size of image & sampler metadata into account for the pod buffer
// size
{
uint32_t max_offset = 0;
if (m_image_metadata) {
// Find how big the POD buffer should be
for (const auto& md : *m_image_metadata) {
auto order_offset = md.second.order_offset;
auto data_type_offset = md.second.data_type_offset;
if (md.second.has_valid_order()) {
max_offset = std::max(order_offset, max_offset);
push_constant_range.offset =
std::min(order_offset, push_constant_range.offset);
if (order_offset + sizeof(uint32_t) >
push_constant_range.offset + push_constant_range.size) {
push_constant_range.size = order_offset +
sizeof(uint32_t) -
push_constant_range.offset;
}
}
if (md.second.has_valid_data_type()) {
max_offset = std::max(data_type_offset, max_offset);
push_constant_range.offset =
std::min(data_type_offset, push_constant_range.offset);
if (data_type_offset + sizeof(uint32_t) >
push_constant_range.offset + push_constant_range.size) {
push_constant_range.size = data_type_offset +
sizeof(uint32_t) -
push_constant_range.offset;
}
}
}
}
if (m_sampler_metadata) {
for (const auto& md : *m_sampler_metadata) {
auto offset = md.second;
max_offset = std::max(offset, max_offset);
push_constant_range.offset =
std::min(offset, push_constant_range.offset);
if (offset + sizeof(uint32_t) >
push_constant_range.offset + push_constant_range.size) {
push_constant_range.size =
offset + sizeof(uint32_t) - push_constant_range.offset;
}
}
}
if (max_offset + sizeof(uint32_t) > m_pod_buffer_size) {
m_pod_buffer_size = round_up(max_offset + sizeof(uint32_t), 4);
}
}
// Don't pass the range at pipeline layout creation time if no push
// constants are used
uint32_t num_push_constant_ranges = 1;
if (push_constant_range.offset == UINT32_MAX) {
num_push_constant_ranges = 0;
}
// The size of the range must be a multiple of 4, round up to guarantee this
push_constant_range.size = round_up(push_constant_range.size, 4);
// Its offset must be a multiple of 4, round down to guarantee this
push_constant_range.offset &= ~0x3U;
if (push_constant_range.size >
m_context->device()->vulkan_max_push_constants_size()) {
cvk_error("Not enough space for push constants");
return CL_INVALID_VALUE;
}
// Create pipeline layout
cvk_debug("about to create pipeline layout, number of descriptor set "
"layouts: %zu",
m_descriptor_set_layouts.size());
VkPipelineLayoutCreateInfo pipelineLayoutCreateInfo = {
VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO,
0,
0,
static_cast<uint32_t>(m_descriptor_set_layouts.size()),
m_descriptor_set_layouts.data(),
num_push_constant_ranges,
&push_constant_range};
res = vkCreatePipelineLayout(m_device->vulkan_device(),
&pipelineLayoutCreateInfo, 0,
&m_pipeline_layout);
if (res != VK_SUCCESS) {
cvk_error("Could not create pipeline layout.");
return CL_INVALID_VALUE;
}
// Determine number and types of bindings
std::vector<VkDescriptorPoolSize> poolSizes(bindingTypes.size());
int bidx = 0;
for (auto& bt : bindingTypes) {
poolSizes[bidx].type = bt.first;
poolSizes[bidx].descriptorCount = bt.second * MAX_INSTANCES;
bidx++;
}
// Create descriptor pool
if (poolSizes.size() > 0) {
VkDescriptorPoolCreateInfo descriptorPoolCreateInfo = {
VK_STRUCTURE_TYPE_DESCRIPTOR_POOL_CREATE_INFO,
nullptr,
VK_DESCRIPTOR_POOL_CREATE_FREE_DESCRIPTOR_SET_BIT, // flags
MAX_INSTANCES * spir_binary::MAX_DESCRIPTOR_SETS, // maxSets
static_cast<uint32_t>(poolSizes.size()), // poolSizeCount
poolSizes.data(), // pPoolSizes
};
res = vkCreateDescriptorPool(m_device->vulkan_device(),
&descriptorPoolCreateInfo, 0,
&m_descriptor_pool);
if (res != VK_SUCCESS) {
cvk_error("Could not create descriptor pool.");
return CL_INVALID_VALUE;
}
}
return CL_SUCCESS;
}
VkPipeline
cvk_entry_point::create_pipeline(const cvk_spec_constant_map& spec_constants) {
std::lock_guard<std::mutex> lock(m_pipeline_cache_lock);
// Check for a cached pipeline using the same specialization constants
if (m_pipelines.count(spec_constants)) {
VkPipeline pipeline = m_pipelines.at(spec_constants);
cvk_info("reusing pipeline %p for kernel %s", pipeline, m_name.c_str());
return pipeline;
}
std::vector<VkSpecializationMapEntry> mapEntries;
std::vector<uint32_t> specConstantData;
uint32_t constantDataOffset = 0;
for (auto& spec_const : spec_constants) {
VkSpecializationMapEntry entry = {spec_const.first, constantDataOffset,
sizeof(uint32_t)};
mapEntries.push_back(entry);
specConstantData.push_back(spec_const.second);
constantDataOffset += sizeof(uint32_t);
}
VkSpecializationInfo specializationInfo = {
static_cast<uint32_t>(mapEntries.size()),
mapEntries.data(),
specConstantData.size() * sizeof(uint32_t),
specConstantData.data(),
};
void* pipelineShaderStageCreateInfoPNext = nullptr;
VkPipelineShaderStageRequiredSubgroupSizeCreateInfo
reqdSubgroupSizeCreateInfo;
if (m_device->supports_subgroup_size_selection()) {
auto reqdSubgroupSize = m_program->required_sub_group_size(m_name);
if (reqdSubgroupSize > m_device->max_sub_group_size() ||
reqdSubgroupSize < m_device->min_sub_group_size()) {
if (reqdSubgroupSize != 0) {
cvk_error_fn("required subgroup size '%u' for '%s' is out of "
"the supported range [%u, %u]",
reqdSubgroupSize, m_name.c_str(),
m_device->min_sub_group_size(),
m_device->max_sub_group_size());
return VK_NULL_HANDLE;
}
reqdSubgroupSize = m_device->sub_group_size();
}
reqdSubgroupSizeCreateInfo.sType =
VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_REQUIRED_SUBGROUP_SIZE_CREATE_INFO;
reqdSubgroupSizeCreateInfo.pNext = nullptr;
reqdSubgroupSizeCreateInfo.requiredSubgroupSize = reqdSubgroupSize;
pipelineShaderStageCreateInfoPNext = &reqdSubgroupSizeCreateInfo;
}
const VkComputePipelineCreateInfo createInfo = {
VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO, // sType
nullptr, // pNext
0, // flags
{
VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO, // sType
pipelineShaderStageCreateInfoPNext, // pNext
0, // flags
VK_SHADER_STAGE_COMPUTE_BIT, // stage
m_program->shader_module(), // module
m_name.c_str(),
&specializationInfo // pSpecializationInfo
}, // stage
m_pipeline_layout, // layout
VK_NULL_HANDLE, // basePipelineHandle
0 // basePipelineIndex
};
VkPipeline pipeline;
VkResult res = vkCreateComputePipelines(m_device->vulkan_device(),
m_program->pipeline_cache(), 1,
&createInfo, nullptr, &pipeline);
if (res != VK_SUCCESS) {
cvk_error_fn("Could not create compute pipeline for kernel %s: %s",
vulkan_error_string(res), m_name.c_str());
return VK_NULL_HANDLE;
}
// Add to pipeline cache
m_pipelines[spec_constants] = pipeline;
cvk_info("created pipeline %p for kernel %s", pipeline, m_name.c_str());
return pipeline;
}
bool cvk_entry_point::allocate_descriptor_sets(VkDescriptorSet* ds) {
TRACE_FUNCTION();
if (m_descriptor_set_layouts.size() == 0) {
return true;
}
std::lock_guard<std::mutex> lock(m_descriptor_pool_lock);
#if CLVK_UNIT_TESTING_ENABLED
if (config.force_descriptor_set_allocation_failure() &&
m_nb_descriptor_set_allocated + m_descriptor_set_layouts.size() >
config.max_entry_points_instances()) {
return false;
}
#endif
VkDescriptorSetAllocateInfo descriptorSetAllocateInfo = {
VK_STRUCTURE_TYPE_DESCRIPTOR_SET_ALLOCATE_INFO, nullptr,
m_descriptor_pool,
static_cast<uint32_t>(
m_descriptor_set_layouts.size()), // descriptorSetCount
m_descriptor_set_layouts.data()};
VkResult res = vkAllocateDescriptorSets(m_device->vulkan_device(),
&descriptorSetAllocateInfo, ds);
if (res != VK_SUCCESS) {
if (config.enqueue_command_retry_sleep_us == UINT32_MAX) {
cvk_error_fn("could not allocate descriptor sets: %s",
vulkan_error_string(res));
} else if (m_first_allocation_failure) {
cvk_warn_fn(
"could not allocate descriptor sets: %s, retry in %u us",
vulkan_error_string(res),
config.enqueue_command_retry_sleep_us());
} else {
cvk_info_fn("could not allocate descriptor sets: %s",
vulkan_error_string(res));
}
m_first_allocation_failure = false;
return false;
}
m_nb_descriptor_set_allocated += m_descriptor_set_layouts.size();
TRACE_CNT(descriptor_set_allocated_counter, m_nb_descriptor_set_allocated);
return true;
}
std::unique_ptr<cvk_buffer> cvk_entry_point::allocate_pod_buffer() {
cl_int err;
auto buffer =
cvk_buffer::create(m_context, 0, m_pod_buffer_size, nullptr, &err);
if (err != CL_SUCCESS) {
return nullptr;
}
return buffer;
}
bool cvk_entry_point::uses_printf() const {
return m_program->kernel_flags(m_name) &
NonSemanticClspvReflectionMayUsePrintf;
}