ClickHouse/src/Common/BitonicSort.h

222 lines
8.5 KiB
C++
Raw Normal View History

2020-05-15 00:01:14 +00:00
#pragma once
#include <string.h>
#if !defined(__APPLE__) && !defined(__FreeBSD__)
#include <malloc.h>
#endif
#ifdef __APPLE__
#include <OpenCL/opencl.h>
#else
#include <CL/cl.h>
#endif
#include <ext/bit_cast.h>
#include <Core/Types.h>
#include <Core/Defines.h>
#include <Common/PODArray.h>
#include <Columns/ColumnsCommon.h>
2020-05-18 10:26:23 +00:00
#include "oclBasics.h"
2020-05-15 00:01:14 +00:00
#include "bitonicSortKernels.cl"
class BitonicSort
{
public:
2020-05-18 21:41:23 +00:00
using KernelType = OCL::KernelType;
enum Types
{
KernelInt8 = 0,
KernelUInt8,
KernelInt16,
KernelUInt16,
KernelInt32,
KernelUInt32,
KernelInt64,
KernelUInt64,
KernelMax
};
2020-05-15 00:01:14 +00:00
static BitonicSort & getInstance()
{
static BitonicSort instance = BitonicSort();
return instance;
}
/// Sorts given array in specified order. Returns `true` if given sequence was sorted, `false` otherwise.
template <typename T>
2020-05-18 21:41:23 +00:00
bool sort(const DB::PaddedPODArray<T> & data, DB::IColumn::Permutation & res, cl_uint sort_ascending [[maybe_unused]]) const
2020-05-15 00:01:14 +00:00
{
2020-05-18 21:41:23 +00:00
if constexpr (
std::is_same_v<T, Int8> ||
std::is_same_v<T, UInt8> ||
std::is_same_v<T, Int16> ||
std::is_same_v<T, UInt16> ||
std::is_same_v<T, Int32> ||
std::is_same_v<T, UInt32> ||
std::is_same_v<T, Int64> ||
std::is_same_v<T, UInt64>)
{
size_t data_size = data.size();
2020-05-15 00:01:14 +00:00
2020-05-18 21:41:23 +00:00
/// Getting the nearest power of 2.
size_t power = 8;
while (power < data_size)
power <<= 1;
2020-05-15 00:01:14 +00:00
2020-05-18 21:41:23 +00:00
/// Allocates more space for additional stubs to be added if needed.
std::vector<T> pairs_content(power);
std::vector<UInt32> pairs_indices(power);
2020-05-15 00:01:14 +00:00
2020-05-18 21:41:23 +00:00
memcpy(&pairs_content[0], &data[0], sizeof(T) * data_size);
for (UInt32 i = 0; i < data_size; ++i)
pairs_indices[i] = i;
2020-05-15 00:01:14 +00:00
2020-05-18 21:41:23 +00:00
fillWithStubs(pairs_content.data(), pairs_indices.data(), data_size, power - data_size, sort_ascending);
sort(pairs_content.data(), pairs_indices.data(), power, sort_ascending);
2020-05-15 00:01:14 +00:00
2020-05-18 21:41:23 +00:00
for (size_t i = 0, shift = 0; i < power; ++i)
2020-05-15 00:01:14 +00:00
{
2020-05-18 21:41:23 +00:00
if (pairs_indices[i] >= data_size)
{
++shift;
continue;
}
res[i - shift] = pairs_indices[i];
2020-05-15 00:01:14 +00:00
}
2020-05-18 21:41:23 +00:00
return true;
2020-05-15 00:01:14 +00:00
}
2020-05-18 21:41:23 +00:00
return false;
2020-05-15 00:01:14 +00:00
}
/// Creating a configuration instance with making all OpenCl required variables
/// such as device, platform, context, queue, program and kernel.
void configure()
{
OCL::Settings settings = OCL::Settings(1, nullptr, 1, nullptr, 1, 0);
cl_platform_id platform = OCL::getPlatformID(settings);
cl_device_id device = OCL::getDeviceID(platform, settings);
cl_context gpu_context = OCL::makeContext(device, settings);
2020-05-18 21:41:23 +00:00
cl_command_queue command_queue = OCL::makeCommandQueue<2>(device, gpu_context, settings);
2020-05-15 00:01:14 +00:00
cl_program program = OCL::makeProgram(bitonic_sort_kernels, gpu_context, device, settings);
/// Creating kernels for each specified data type.
cl_int error = 0;
2020-05-18 21:41:23 +00:00
kernels.resize(KernelMax);
2020-05-15 00:01:14 +00:00
2020-05-18 21:41:23 +00:00
kernels[KernelInt8] = std::shared_ptr<KernelType>(clCreateKernel(program, "bitonicSort_char", &error), clReleaseKernel);
2020-05-15 00:01:14 +00:00
OCL::checkError(error);
2020-05-18 21:41:23 +00:00
kernels[KernelUInt8] = std::shared_ptr<KernelType>(clCreateKernel(program, "bitonicSort_uchar", &error), clReleaseKernel);
OCL::checkError(error);
2020-05-15 00:01:14 +00:00
2020-05-18 21:41:23 +00:00
kernels[KernelInt16] = std::shared_ptr<KernelType>(clCreateKernel(program, "bitonicSort_short", &error), clReleaseKernel);
OCL::checkError(error);
2020-05-15 00:01:14 +00:00
2020-05-18 21:41:23 +00:00
kernels[KernelUInt16] = std::shared_ptr<KernelType>(clCreateKernel(program, "bitonicSort_ushort", &error), clReleaseKernel);
OCL::checkError(error);
2020-05-15 00:01:14 +00:00
2020-05-18 21:41:23 +00:00
kernels[KernelInt32] = std::shared_ptr<KernelType>(clCreateKernel(program, "bitonicSort_int", &error), clReleaseKernel);
OCL::checkError(error);
2020-05-15 00:01:14 +00:00
2020-05-18 21:41:23 +00:00
kernels[KernelUInt32] = std::shared_ptr<KernelType>(clCreateKernel(program, "bitonicSort_uint", &error), clReleaseKernel);
OCL::checkError(error);
2020-05-15 00:01:14 +00:00
2020-05-18 21:41:23 +00:00
kernels[KernelInt64] = std::shared_ptr<KernelType>(clCreateKernel(program, "bitonicSort_long", &error), clReleaseKernel);
OCL::checkError(error);
2020-05-15 00:01:14 +00:00
2020-05-18 21:41:23 +00:00
kernels[KernelUInt64] = std::shared_ptr<KernelType>(clCreateKernel(program, "bitonicSort_ulong", &error), clReleaseKernel);
OCL::checkError(error);
2020-05-15 00:01:14 +00:00
2020-05-18 21:41:23 +00:00
configuration = std::shared_ptr<OCL::Configuration>(new OCL::Configuration(device, gpu_context, command_queue, program));
2020-05-15 00:01:14 +00:00
}
2020-05-18 21:41:23 +00:00
private:
/// Dictionary with kernels for each type from list: uchar, char, ushort, short, uint, int, ulong and long.
std::vector<std::shared_ptr<KernelType>> kernels;
/// Current configuration with core OpenCL instances.
std::shared_ptr<OCL::Configuration> configuration = nullptr;
2020-05-15 00:01:14 +00:00
2020-05-18 21:41:23 +00:00
cl_kernel getKernel(Int8) const { return kernels[KernelInt8].get(); }
cl_kernel getKernel(UInt8) const { return kernels[KernelUInt8].get(); }
cl_kernel getKernel(Int16) const { return kernels[KernelInt16].get(); }
cl_kernel getKernel(UInt16) const { return kernels[KernelUInt16].get(); }
cl_kernel getKernel(Int32) const { return kernels[KernelInt32].get(); }
cl_kernel getKernel(UInt32) const { return kernels[KernelUInt32].get(); }
cl_kernel getKernel(Int64) const { return kernels[KernelInt64].get(); }
cl_kernel getKernel(UInt64) const { return kernels[KernelUInt64].get(); }
2020-05-15 00:01:14 +00:00
/// Sorts p_input inplace with indices. Works only with arrays which size equals to power of two.
template <class T>
2020-05-18 21:41:23 +00:00
void sort(T * p_input, cl_uint * indices, cl_int array_size, cl_uint sort_ascending) const
2020-05-15 00:01:14 +00:00
{
2020-05-18 21:41:23 +00:00
cl_kernel kernel = getKernel(T(0));
2020-05-15 00:01:14 +00:00
cl_int error = CL_SUCCESS;
cl_int num_stages = 0;
for (cl_int temp = array_size; temp > 2; temp >>= 1)
num_stages++;
/// Creating OpenCL buffers using input arrays memory.
cl_mem cl_input_buffer = OCL::createBuffer<T>(p_input, array_size, configuration.get()->context());
cl_mem cl_indices_buffer = OCL::createBuffer<cl_uint>(indices, array_size, configuration.get()->context());
configureKernel<cl_mem>(kernel, 0, static_cast<void *>(&cl_input_buffer));
configureKernel<cl_mem>(kernel, 1, static_cast<void *>(&cl_indices_buffer));
configureKernel<cl_uint>(kernel, 4, static_cast<void *>(&sort_ascending));
for (cl_int stage = 0; stage < num_stages; stage++)
{
configureKernel<cl_uint>(kernel, 2, static_cast<void *>(&stage));
for (cl_int pass_of_stage = stage; pass_of_stage >= 0; pass_of_stage--)
{
configureKernel<cl_uint>(kernel, 3, static_cast<void *>(&pass_of_stage));
/// Setting work-item dimensions.
size_t gsize = array_size / (2 * 4);
size_t global_work_size[1] = {pass_of_stage ? gsize : gsize << 1 }; // number of quad items in input array
/// Executing kernel.
error = clEnqueueNDRangeKernel(configuration.get()->commandQueue(), kernel, 1, nullptr,
global_work_size, nullptr, 0, nullptr, nullptr);
OCL::checkError(error);
}
}
/// Syncs all threads.
OCL::finishCommandQueue(configuration.get()->commandQueue());
OCL::releaseData(p_input, array_size, cl_input_buffer, configuration.get()->commandQueue());
OCL::releaseData(indices, array_size, cl_indices_buffer, configuration.get()->commandQueue());
}
template <class T>
2020-05-18 21:41:23 +00:00
void configureKernel(cl_kernel kernel, int number_of_argument, void * source) const
2020-05-15 00:01:14 +00:00
{
cl_int error = clSetKernelArg(kernel, number_of_argument, sizeof(T), source);
OCL::checkError(error);
}
/// Fills given sequences from `arraySize` index with `numberOfStubs` values.
template <class T>
2020-05-18 21:41:23 +00:00
void fillWithStubs(T * p_input, cl_uint * indices, cl_int array_size, cl_int number_of_stubs, cl_uint sort_ascending) const
2020-05-15 00:01:14 +00:00
{
2020-05-18 21:41:23 +00:00
T value = sort_ascending ? std::numeric_limits<T>::max() : std::numeric_limits<T>::min();
2020-05-15 00:01:14 +00:00
for (cl_int index = 0; index < number_of_stubs; ++index)
{
p_input[array_size + index] = value;
indices[array_size + index] = array_size + index;
}
}
2020-05-18 21:41:23 +00:00
BitonicSort() = default;
BitonicSort(BitonicSort const &) = delete;
void operator = (BitonicSort const &) = delete;
2020-05-15 00:01:14 +00:00
};