mirror of
https://github.com/StepanovPlaton/NeuralNetwork.git
synced 2026-04-04 04:40:40 +04:00
Check
This commit is contained in:
@@ -1,10 +1,14 @@
|
||||
#include "opencl.hpp"
|
||||
#include <CL/opencl.hpp>
|
||||
|
||||
#include "opencl.hpp"
|
||||
|
||||
#include <format>
|
||||
#include <iostream>
|
||||
#include <ostream>
|
||||
#include <string>
|
||||
#include <unordered_map>
|
||||
|
||||
template <typename T, int Dim> class Kernels {
|
||||
template <typename T> class Kernels {
|
||||
public:
|
||||
enum class Vector {
|
||||
type2 = 2,
|
||||
@@ -21,136 +25,211 @@ public:
|
||||
T_HADAMARD,
|
||||
T_MULT,
|
||||
};
|
||||
constexpr const static std::string type = typeid(T).name();
|
||||
|
||||
// TODO: get native vector size
|
||||
static Vector vector = Vector::type8;
|
||||
|
||||
private:
|
||||
static std::string unaryOperation(std::string name, std::string operation) {
|
||||
return std::format(
|
||||
constexpr std::string getTypeName() { return "unknown"; }
|
||||
Vector vector;
|
||||
std::string configuration;
|
||||
|
||||
std::string format(std::string tmp,
|
||||
std::unordered_map<std::string, std::string> args) {
|
||||
std::string result(tmp);
|
||||
for (const auto &[key, value] : args) {
|
||||
std::string placeholder = "{" + key + "}";
|
||||
size_t pos = 0;
|
||||
while ((pos = result.find(placeholder, pos)) != std::string::npos) {
|
||||
result.replace(pos, placeholder.length(), value);
|
||||
pos += value.length();
|
||||
}
|
||||
}
|
||||
return result;
|
||||
}
|
||||
|
||||
std::string unaryOperation(std::string name, std::string operation) {
|
||||
return format(
|
||||
R"(
|
||||
__kernel void {method}(__global {type}* A, int len) {{
|
||||
__kernel void {method}(__global type* A, int len) {
|
||||
int gid = get_global_id(0);
|
||||
int base = gid * {vector};
|
||||
if (base + ({vector}-1) < len) {{
|
||||
{type}{vector} data = vload{vector}(gid, A);
|
||||
vstore{vector}({operation}data, gid, A);
|
||||
}} else {{
|
||||
for (int i = 0; i < {vec_size}; i++) {{
|
||||
int base = gid * WIDTH;
|
||||
if (base + WIDTH <= len) {
|
||||
typeX data = vloadX(gid, A);
|
||||
vstoreX({operation}data, gid, A);
|
||||
} else {
|
||||
for (int i = 0; i < WIDTH; i++) {
|
||||
int idx = base + i;
|
||||
if (idx < len) A[idx] = {operation}A[idx];
|
||||
}}
|
||||
}}
|
||||
}}
|
||||
)",
|
||||
std::make_format_args(std::make_pair("method", name),
|
||||
std::make_pair("vector", vector),
|
||||
std::make_pair("type", type),
|
||||
std::make_pair("operation", operation)));
|
||||
}
|
||||
}
|
||||
})",
|
||||
{{"method", name}, {"operation", operation}});
|
||||
}
|
||||
|
||||
static std::string scalarOperation(std::string name, std::string operation) {
|
||||
return std::format(
|
||||
std::string scalarOperation(std::string name, std::string operation) {
|
||||
return format(
|
||||
R"(
|
||||
__kernel void {method}(__global {type}* A, int len, {type} scalar) {{
|
||||
__kernel void {method}(__global type* A, int len, type scalar) {
|
||||
int gid = get_global_id(0);
|
||||
int base = gid * {vector};
|
||||
if (base + ({vector}-1) < len) {{
|
||||
{type}{vector} data = vload{vector}(gid, A);
|
||||
int base = gid * WIDTH;
|
||||
if (base + WIDTH <= len) {
|
||||
typeX data = vloadX(gid, A);
|
||||
data = data {operation} scalar;
|
||||
vstore{vector}(data, gid, A);
|
||||
}} else {{
|
||||
for (int i = 0; i < {vec_size}; i++) {{
|
||||
vstoreX(data, gid, A);
|
||||
} else {
|
||||
for (int i = 0; i < WIDTH; i++) {
|
||||
int idx = base + i;
|
||||
if (idx < len) A[idx] = A[idx] {operation} scalar;
|
||||
}}
|
||||
}}
|
||||
}}
|
||||
)",
|
||||
std::make_format_args(std::make_pair("method", name),
|
||||
std::make_pair("vector", vector),
|
||||
std::make_pair("type", type),
|
||||
std::make_pair("operation", operation)));
|
||||
}
|
||||
}
|
||||
})",
|
||||
{{"method", name}, {"operation", operation}});
|
||||
}
|
||||
|
||||
static std::string binaryOperation(std::string name, std::string operation) {
|
||||
return std::format(
|
||||
std::string binaryOperation(std::string name, std::string operation) {
|
||||
return format(
|
||||
R"(
|
||||
__kernel void {method}(__global {type}* A, __global {type}* B, int len) {{
|
||||
__kernel void {method}(__global type* A, __global type* B, int len) {
|
||||
int gid = get_global_id(0);
|
||||
int base = gid * {vector};
|
||||
if (base + ({vector}-1) < len) {{
|
||||
{type}{vector} dataA = vload{vector}(gid, A);
|
||||
{type}{vector} dataB = vload{vector}(gid, B);
|
||||
vstore{vector}(dataA {operation} dataB, gid, A);
|
||||
}} else {{
|
||||
for (int i = 0; i < {vector}; i++) {{
|
||||
int base = gid * WIDTH;
|
||||
if (base + WIDTH <= len) {
|
||||
typeX dataA = vloadX(gid, A);
|
||||
typeX dataB = vloadX(gid, B);
|
||||
vstoreX(dataA {operation} dataB, gid, A);
|
||||
} else {
|
||||
for (int i = 0; i < WIDTH; i++) {
|
||||
int idx = base + i;
|
||||
if (idx < len) A[idx] = A[idx] {operation} B[idx];
|
||||
}}
|
||||
}}
|
||||
}}
|
||||
)",
|
||||
std::make_format_args(std::make_pair("method", name),
|
||||
std::make_pair("vector", vector),
|
||||
std::make_pair("type", type),
|
||||
std::make_pair("operation", operation)));
|
||||
}
|
||||
}
|
||||
})",
|
||||
{{"method", name}, {"operation", operation}});
|
||||
}
|
||||
|
||||
static std::unordered_map<Method, std::tuple<std::string, std::string>>
|
||||
programs = {
|
||||
{Method::POSITIVE, {unaryOperation("positive", "+"), "positive"}},
|
||||
{Method::NEGATIVE, {unaryOperation("negative", "-")}, "negative"},
|
||||
std::string matrixMult(std::string name) {
|
||||
return format(
|
||||
R"(
|
||||
#define TILE_SIZE WIDTH*4
|
||||
__kernel void mult(const __global typeX* A,
|
||||
const __global typeX* B,
|
||||
__global typeX* C, const int M, const int N, const int K) {
|
||||
const int row = get_local_id(0);
|
||||
const int col = get_local_id(1);
|
||||
const int globalRow = (TILE_SIZE/WIDTH)*get_group_id(0) + row;
|
||||
const int globalCol = TILE_SIZE*get_group_id(1) + col;
|
||||
__local typeX Asub[TILE_SIZE][TILE_SIZE/WIDTH];
|
||||
__local typeX Bsub[TILE_SIZE][TILE_SIZE/WIDTH];
|
||||
typeX acc = 0;
|
||||
const int numTiles = K/TILE_SIZE;
|
||||
for (int tile = 0; tile < numTiles; tile++) {
|
||||
const int tiledRow = (TILE_SIZE/WIDTH)*tile + row;
|
||||
const int tiledCol = TILE_SIZE*tile + col;
|
||||
Asub[col][row] = A[tiledCol*(M/WIDTH) + globalRow];
|
||||
Bsub[col][row] = B[globalCol*(K/WIDTH) + tiledRow];
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
typeX vecA, vecB;
|
||||
type valB;
|
||||
for (int k = 0; k < TILE_SIZE/WIDTH; k++) {
|
||||
vecB = Bsub[col][k];
|
||||
for (int w = 0; w < WIDTH; w++) {
|
||||
vecA = Asub[WIDTH*k + w][row];
|
||||
valB = vecB[w];
|
||||
for (int i = 0; i < WIDTH; i++)
|
||||
acc[i] += vecA[i] * valB;
|
||||
}
|
||||
}
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
}
|
||||
C[globalCol*(M/WIDTH) + globalRow] = acc;
|
||||
}
|
||||
)",
|
||||
{{"method", name}});
|
||||
}
|
||||
|
||||
{Method::S_ADD, {scalarOperation("add", "+")}, "add"},
|
||||
{Method::S_MULT, {scalarOperation("mult", "*")}, "mult"},
|
||||
std::unordered_map<Method, std::tuple<std::string, std::string>> programs = {
|
||||
{Method::POSITIVE, {unaryOperation("positive", "+"), "positive"}},
|
||||
{Method::NEGATIVE, {unaryOperation("negative", "-"), "negative"}},
|
||||
|
||||
{Method::T_ADD, {binaryOperation("add", "+")}, "add"},
|
||||
{Method::T_HADAMARD,
|
||||
{binaryOperation("hadamard_mult", "*")},
|
||||
"hadamard_mult"},
|
||||
{Method::T_MULT, {"", "mult"}},
|
||||
{Method::S_ADD, {scalarOperation("add", "+"), "add"}},
|
||||
{Method::S_MULT, {scalarOperation("mult", "*"), "mult"}},
|
||||
|
||||
{Method::T_ADD, {binaryOperation("add", "+"), "add"}},
|
||||
{Method::T_HADAMARD,
|
||||
{binaryOperation("hadamard_mult", "*"), "hadamard_mult"}},
|
||||
|
||||
{Method::T_MULT, {matrixMult("mult"), "mult"}},
|
||||
};
|
||||
|
||||
static inline std::unordered_map<Method, cl::Program> compiledPrograms;
|
||||
static inline std::mutex compileMutex;
|
||||
std::unordered_map<Method, cl::Program> compiledPrograms;
|
||||
|
||||
public:
|
||||
static cl::Kernel create(Method method) {
|
||||
std::lock_guard<std::mutex> lock(compileMutex);
|
||||
Kernels(Vector vec = Vector::type4) : vector(vec) {
|
||||
std::string extensions = openCL.getDevice().getInfo<CL_DEVICE_EXTENSIONS>();
|
||||
if (extensions.find("cl_khr_fp16") != std::string::npos)
|
||||
configuration = R"(
|
||||
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
|
||||
typedef half _half;
|
||||
typedef half2 _half2;
|
||||
typedef half4 _half4;
|
||||
typedef half8 _half8;
|
||||
typedef half16 _half16;
|
||||
)";
|
||||
else
|
||||
configuration = R"(
|
||||
typedef float _half;
|
||||
typedef float2 _half2;
|
||||
typedef float4 _half4;
|
||||
typedef float8 _half8;
|
||||
typedef float16 _half16;
|
||||
)";
|
||||
configuration += format(
|
||||
R"(
|
||||
typedef {type} type;
|
||||
typedef {type}{vector} typeX;
|
||||
#define WIDTH {vector}
|
||||
#define vloadX vload{vector}
|
||||
#define vstoreX vstore{vector}
|
||||
)",
|
||||
{{"type", getTypeName()}, {"vector", std::to_string((int)vector)}});
|
||||
|
||||
auto cache = compiledPrograms.find(method);
|
||||
if (cache != compiledPrograms.end()) {
|
||||
const auto &programName = std::get<1>(programs[method]);
|
||||
return cl::Kernel(cache->second, programName.c_str());
|
||||
}
|
||||
|
||||
auto program = programs.find(method);
|
||||
if (program == programs.end())
|
||||
throw std::runtime_error("Unknown method: " +
|
||||
std::to_string(static_cast<int>(method)));
|
||||
const auto &[sourceCode, kernelName] = program->second;
|
||||
|
||||
try {
|
||||
cl::Program::Sources sources;
|
||||
sources.push_back({sourceCode.c_str(), sourceCode.length()});
|
||||
cl::Program program(openCL.getContext(), sources);
|
||||
program.build({openCL.getDevice()});
|
||||
compiledPrograms[method] = program;
|
||||
return cl::Kernel(program, kernelName.c_str());
|
||||
|
||||
} catch (const cl::Error &e) {
|
||||
if (e.err() == CL_BUILD_PROGRAM_FAILURE) {
|
||||
cl::Program program(openCL.getContext(),
|
||||
{sourceCode.c_str(), sourceCode.length()});
|
||||
auto buildInfo =
|
||||
program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(openCL.getDevice());
|
||||
throw std::runtime_error(
|
||||
"OpenCL compilation failed: " + std::string(e.what()) +
|
||||
"\nBuild log:\n" + buildInfo);
|
||||
for (const auto &[method, programInfo] : programs) {
|
||||
const auto &[sourceCode, kernelName] = programInfo;
|
||||
if (!sourceCode.empty()) {
|
||||
cl::Program program(openCL.getContext(), configuration + sourceCode);
|
||||
try {
|
||||
program.build({openCL.getDevice()});
|
||||
compiledPrograms[method] = program;
|
||||
} catch (const cl::Error &e) {
|
||||
std::cerr << "OpenCL compilation error for method "
|
||||
<< static_cast<int>(method) << ": " << e.what()
|
||||
<< std::endl;
|
||||
std::string buildLog =
|
||||
program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(openCL.getDevice());
|
||||
std::cerr << "Build log for method " << static_cast<int>(method)
|
||||
<< ":" << std::endl;
|
||||
std::cerr << buildLog << std::endl;
|
||||
}
|
||||
}
|
||||
throw std::runtime_error("OpenCL error: " + std::string(e.what()));
|
||||
}
|
||||
}
|
||||
|
||||
cl::Kernel create(Method method) {
|
||||
auto it = compiledPrograms.find(method);
|
||||
if (it == compiledPrograms.end())
|
||||
throw std::runtime_error("Program for method not found or not compiled");
|
||||
const auto &kernelName = std::get<1>(programs[method]);
|
||||
return cl::Kernel(it->second, kernelName.c_str());
|
||||
}
|
||||
};
|
||||
|
||||
#define SPECIALIZE_KERNELS_TYPE(type, name) \
|
||||
template <> constexpr std::string Kernels<type>::getTypeName() { \
|
||||
return name; \
|
||||
}
|
||||
SPECIALIZE_KERNELS_TYPE(char, "char")
|
||||
SPECIALIZE_KERNELS_TYPE(short, "short")
|
||||
SPECIALIZE_KERNELS_TYPE(int, "int")
|
||||
SPECIALIZE_KERNELS_TYPE(long, "long")
|
||||
SPECIALIZE_KERNELS_TYPE(float, "float")
|
||||
SPECIALIZE_KERNELS_TYPE(double, "double")
|
||||
|
||||
typedef cl_half half;
|
||||
SPECIALIZE_KERNELS_TYPE(half, "_half")
|
||||
Reference in New Issue
Block a user