From e4455322d8b85a0e296fbdb19db8fb9226df6268 Mon Sep 17 00:00:00 2001 From: Liangliang He Date: Thu, 2 Nov 2017 17:46:32 +0800 Subject: [PATCH] Add opencl kernel with image example --- mace/core/runtime/opencl/opencl_runtime.cc | 6 +- mace/core/runtime/opencl/opencl_wrapper.cc | 38 +++++- mace/kernels/opencl/cl/common.h | 2 + mace/kernels/opencl/cl/conv_2d_1x1.cl | 123 +++++++++++++---- mace/kernels/opencl/conv_2d_opencl_1x1.cc | 152 +++++++++++++++++---- 5 files changed, 258 insertions(+), 63 deletions(-) create mode 100644 mace/kernels/opencl/cl/common.h diff --git a/mace/core/runtime/opencl/opencl_runtime.cc b/mace/core/runtime/opencl/opencl_runtime.cc index b85b9e48..26bc60d8 100644 --- a/mace/core/runtime/opencl/opencl_runtime.cc +++ b/mace/core/runtime/opencl/opencl_runtime.cc @@ -26,6 +26,7 @@ bool ReadSourceFile(const std::string &filename, std::string *content) { std::string line; while (std::getline(ifs, line)) { *content += line; + *content += "\n"; } ifs.close(); return true; @@ -66,14 +67,15 @@ bool BuildProgram(OpenCLRuntime *runtime, *program = cl::Program(runtime->context(), sources); std::string build_options = "-Werror -cl-mad-enable -cl-fast-relaxed-math -I" + path; // TODO(heliangliang) -cl-unsafe-math-optimizations -cl-fast-relaxed-math - if (program->build({runtime->device()}, build_options.c_str()) != CL_SUCCESS) { + cl_int ret = program->build({runtime->device()}, build_options.c_str()); + if (ret != CL_SUCCESS) { if (program->getBuildInfo(runtime->device()) == CL_BUILD_ERROR) { std::string build_log = program->getBuildInfo(runtime->device()); LOG(INFO) << "Program build log: " << build_log; } - LOG(FATAL) << "Build program failed"; + LOG(FATAL) << "Build program failed: " << ret; } return true; diff --git a/mace/core/runtime/opencl/opencl_wrapper.cc b/mace/core/runtime/opencl/opencl_wrapper.cc index 49f92934..ad8a98f8 100644 --- a/mace/core/runtime/opencl/opencl_wrapper.cc +++ b/mace/core/runtime/opencl/opencl_wrapper.cc @@ -126,6 +126,12 @@ class OpenCLLibraryImpl final { using clRetainKernelFunc = cl_int (*)(cl_kernel kernel); using clCreateBufferFunc = cl_mem (*)(cl_context, cl_mem_flags, size_t, void *, cl_int *); + using clCreateImageFunc = cl_mem (*)(cl_context, + cl_mem_flags, + const cl_image_format *, + const cl_image_desc *, + void *, + cl_int *); using clCreateProgramWithSourceFunc = cl_program (*)( cl_context, cl_uint, const char **, const size_t *, cl_int *); using clReleaseKernelFunc = cl_int (*)(cl_kernel kernel); @@ -136,8 +142,12 @@ class OpenCLLibraryImpl final { using clRetainDeviceFunc = cl_int (*)(cl_device_id); using clReleaseDeviceFunc = cl_int (*)(cl_device_id); using clRetainEventFunc = cl_int (*)(cl_event); - using clGetKernelWorkGroupInfoFunc = - cl_int (*)(cl_kernel, cl_device_id, cl_kernel_work_group_info, size_t, void *, size_t *); + using clGetKernelWorkGroupInfoFunc = cl_int (*)(cl_kernel, + cl_device_id, + cl_kernel_work_group_info, + size_t, + void *, + size_t *); #define DEFINE_FUNC_PTR(func) func##Func func = nullptr @@ -149,6 +159,7 @@ class OpenCLLibraryImpl final { DEFINE_FUNC_PTR(clReleaseKernel); DEFINE_FUNC_PTR(clCreateProgramWithSource); DEFINE_FUNC_PTR(clCreateBuffer); + DEFINE_FUNC_PTR(clCreateImage); DEFINE_FUNC_PTR(clRetainKernel); DEFINE_FUNC_PTR(clCreateKernel); DEFINE_FUNC_PTR(clGetProgramInfo); @@ -269,6 +280,7 @@ void *OpenCLLibraryImpl::LoadFromPath(const std::string &path) { ASSIGN_FROM_DLSYM(clReleaseKernel); ASSIGN_FROM_DLSYM(clCreateProgramWithSource); ASSIGN_FROM_DLSYM(clCreateBuffer); + ASSIGN_FROM_DLSYM(clCreateImage); ASSIGN_FROM_DLSYM(clRetainKernel); ASSIGN_FROM_DLSYM(clCreateKernel); ASSIGN_FROM_DLSYM(clGetProgramInfo); @@ -708,6 +720,24 @@ cl_mem clCreateBuffer(cl_context context, } } +cl_mem clCreateImage(cl_context context, + cl_mem_flags flags, + const cl_image_format *image_format, + const cl_image_desc *image_desc, + void *host_ptr, + cl_int *errcode_ret) { + auto func = mace::OpenCLLibraryImpl::Get().clCreateImage; + if (func != nullptr) { + return func(context, flags, image_format, image_desc, host_ptr, + errcode_ret); + } else { + if (errcode_ret != nullptr) { + *errcode_ret = CL_OUT_OF_RESOURCES; + } + return nullptr; + } +} + cl_program clCreateProgramWithSource(cl_context context, cl_uint count, const char **strings, @@ -795,8 +825,8 @@ cl_int clGetKernelWorkGroupInfo(cl_kernel kernel, size_t *param_value_size_ret) { auto func = mace::OpenCLLibraryImpl::Get().clGetKernelWorkGroupInfo; if (func != nullptr) { - return func(kernel, device, param_name, param_value_size, - param_value, param_value_size_ret); + return func(kernel, device, param_name, param_value_size, param_value, + param_value_size_ret); } else { return CL_OUT_OF_RESOURCES; } diff --git a/mace/kernels/opencl/cl/common.h b/mace/kernels/opencl/cl/common.h new file mode 100644 index 00000000..a6be3c53 --- /dev/null +++ b/mace/kernels/opencl/cl/common.h @@ -0,0 +1,2 @@ +#pragma OPENCL EXTENSION cl_khr_fp16 : enable +#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable diff --git a/mace/kernels/opencl/cl/conv_2d_1x1.cl b/mace/kernels/opencl/cl/conv_2d_1x1.cl index 7b856f89..dc5b1f81 100644 --- a/mace/kernels/opencl/cl/conv_2d_1x1.cl +++ b/mace/kernels/opencl/cl/conv_2d_1x1.cl @@ -1,21 +1,22 @@ -void kernel conv_2d_1x1_naive(global const float *input, /* n, c, h, w */ - global const float *filter, /* o, i, kh, kw */ - global const float *bias, /* o */ - global float *output, /* n, c, h, w */ - private const int input_channels) { +#include + +__kernel void conv_2d_1x1_naive(__global const float *input, /* n, c, h, w */ + __global const float *filter, /* o, i, kh, kw */ + __global const float *bias, /* o */ + __global float *output, /* n, c, h, w */ + __private const int in_chan_num) { const int batch = get_global_id(0); const int channel = get_global_id(1); const int channels = get_global_size(1); const int pixel = get_global_id(2); const int pixels = get_global_size(2); - float *output_ptr = output + (batch * channels + channel) * pixels; output_ptr[pixel] = bias[channel]; - for (int inc = 0; inc < input_channels; ++inc) { - const float *input_ptr = input + (batch * input_channels + inc) * pixels + pixel; - const float weights = filter[channel * input_channels + inc]; + for (int inc = 0; inc < in_chan_num; ++inc) { + const float *input_ptr = input + (batch * in_chan_num + inc) * pixels + pixel; + const float weights = filter[channel * in_chan_num + inc]; float in = input_ptr[0]; float out = output_ptr[0]; out += in * weights; @@ -23,17 +24,19 @@ void kernel conv_2d_1x1_naive(global const float *input, /* n, c, h, w */ } } -void kernel conv_2d_1x1_v2(global const float *input, /* n, c, h, w */ - global const float *filter, /* o, i, kh, kw */ - global const float *bias, /* o */ - global float *output, /* n, c, h, w */ - private const int in_chan_num, - private const int out_chan_num, - private const int pixel_num) { +__kernel void conv_2d_1x1_v2(__global const float *input, /* n, c, h, w */ + __global const float *filter, /* o, i, kh, kw */ + __global const float *bias, /* o */ + __global float *output, /* n, c, h, w */ + __private const int in_chan_num, + __private const int out_chan_num, + __private const int pixel_num) { int batch = get_global_id(0); int out_chan_blk = get_global_id(1); int out_pixel_blk = get_global_id(2); + __private float output_slice[4 * 4]; + const int out_chan_begin = out_chan_blk * 4; const int out_chan_end = min(out_chan_begin + 4, out_chan_num); const int out_pixel_begin = out_pixel_blk * 4; @@ -41,21 +44,23 @@ void kernel conv_2d_1x1_v2(global const float *input, /* n, c, h, w */ const int in_offset = batch * in_chan_num * pixel_num; const int out_offset = batch * out_chan_num * pixel_num; + const float *input_base = input + in_offset + out_pixel_begin; float *output_base = output + out_offset + out_pixel_begin; - int pixels = out_pixel_end - out_pixel_begin; + int out_chan_len = out_chan_end - out_chan_begin; + int pixel_len = out_pixel_end - out_pixel_begin; for (int out_chan = out_chan_begin; out_chan < out_chan_end; ++out_chan) { float bias_value = bias[out_chan]; - float *output_ptr = output_base + out_chan * pixel_num; - for (int p = 0; p < pixels; ++p) { - output_ptr[p] = bias_value; + int out_chan_offset = out_chan - out_chan_begin; + for (int p = 0; p < pixel_len; ++p) { + output_slice[out_chan_offset * 4 + p] = bias_value; } } int in_chan = 0; - if (pixels == 4) { + if (pixel_len == 4) { for (; in_chan + 3 < in_chan_num; in_chan += 4) { const float *input_ptr = input_base + in_chan * pixel_num; int out_chan = out_chan_begin; @@ -66,30 +71,32 @@ void kernel conv_2d_1x1_v2(global const float *input, /* n, c, h, w */ float4 in1 = vload4(0, input_ptr + pixel_num); float4 in2 = vload4(0, input_ptr + 2 * pixel_num); float4 in3 = vload4(0, input_ptr + 3 * pixel_num); + #pragma unroll for (int oc = 0; oc < 4; ++oc) { float4 weights = vload4(0, filter_ptr + oc * in_chan_num); - float4 out = vload4(0, output_ptr + oc * pixel_num); + float4 out = vload4(oc, output_slice); out += in0 * weights.x; out += in1 * weights.y; out += in2 * weights.z; out += in3 * weights.w; - vstore4(out, 0, output_ptr + oc * pixel_num); + vstore4(out, oc, output_slice); } } for (; out_chan < out_chan_end; ++out_chan) { const float* filter_ptr = filter + out_chan * in_chan_num + in_chan; float *output_ptr = output_base + out_chan * pixel_num; + int out_chan_offset = out_chan - out_chan_begin; float4 weights = vload4(0, filter_ptr); float4 in0 = vload4(0, input_ptr); float4 in1 = vload4(0, input_ptr + pixel_num); float4 in2 = vload4(0, input_ptr + 2 * pixel_num); float4 in3 = vload4(0, input_ptr + 3 * pixel_num); - float4 out = vload4(0, output_ptr); + float4 out = vload4(out_chan_offset, output_slice); out += in0 * weights.x; out += in1 * weights.y; out += in2 * weights.z; out += in3 * weights.w; - vstore4(out, 0, output_ptr); + vstore4(out, out_chan_offset, output_slice); } } } @@ -99,13 +106,71 @@ void kernel conv_2d_1x1_v2(global const float *input, /* n, c, h, w */ for (int out_chan = out_chan_begin; out_chan < out_chan_end; ++out_chan) { float weights = filter[out_chan * in_chan_num + in_chan]; float *output_ptr = output_base + out_chan * pixel_num; + int out_chan_offset = out_chan - out_chan_begin; - for (int p = 0; p < pixels; ++p) { + for (int p = 0; p < pixel_len; ++p) { float in = input_ptr[p]; - float out = output_ptr[p]; - out += in * weights; - output_ptr[p] = out; + output_slice[out_chan_offset * 4 + p] += in * weights; } } } + + for (int out_chan_offset = 0; out_chan_offset < out_chan_len; ++out_chan_offset) { + int out_chan = out_chan_begin + out_chan_offset; + float *output_ptr = output_base + out_chan * pixel_num; + if (pixel_len == 4) { + float4 out = vload4(out_chan_offset, output_slice); + vstore4(out, 0, output_ptr); + } else { + int offset = out_chan_offset << 2; + for (int p = 0; p < pixel_len; ++p) { + output_ptr[p] = output_slice[offset + p]; + } + } + } +} + +/* FIXME this is incomplete */ +__kernel void conv_2d_1x1_v3(__read_only image3d_t input, /* n, c/4, h, w, 4 */ + __global const float *filter, /* o, i, kh, kw */ + __global const float *bias, /* o */ + __write_only image3d_t output, /* n, c/4, h, w, 4 */ + __private const int batch_num, + __private const int in_chan_num, + __private const int out_chan_num, + __private const int height, + __private const int width) { + int out_chan_blk = get_global_id(0); + int h = get_global_id(1); + int w = get_global_id(2); + + + int in_chan_blk_num = (in_chan_num + 3) / 4; + int out_chan_blk_num = (out_chan_num + 3) / 4; + + const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST; + + for (int batch = 0; batch < batch_num; ++batch) { + float4 bias_value = vload4(out_chan_blk, bias); + __private float4 out = bias_value; + + for (int in_chan_blk = 0; in_chan_blk < in_chan_blk_num; ++in_chan_blk) { + int in_d = batch * in_chan_blk_num + in_chan_blk; + float4 in = read_imagef(input, sampler, (int4)(in_d, h, w, 0)); + + const float *filter_base = filter + (out_chan_blk << 2) * in_chan_num; + float4 weights = vload4(in_chan_blk, filter_base); + out.x += dot(in, weights); + weights = vload4(in_chan_blk, filter_base + in_chan_num); + out.y += dot(in, weights); + weights = vload4(in_chan_blk, filter_base + in_chan_num * 2); + out.z += dot(in, weights); + weights = vload4(in_chan_blk, filter_base + in_chan_num * 3); + out.w += dot(in, weights); + } + + int out_d = batch * out_chan_blk_num + out_chan_blk; + int4 out_coord = (int4)(out_d, h, w, 0); + write_imagef(output, out_coord, out); + } } diff --git a/mace/kernels/opencl/conv_2d_opencl_1x1.cc b/mace/kernels/opencl/conv_2d_opencl_1x1.cc index 6f4bd21d..130ca4b7 100644 --- a/mace/kernels/opencl/conv_2d_opencl_1x1.cc +++ b/mace/kernels/opencl/conv_2d_opencl_1x1.cc @@ -2,9 +2,10 @@ // Copyright (c) 2017 XiaoMi All rights reserved. // +#include "mace/kernels/conv_2d.h" #include "mace/core/common.h" +#include "mace/core/runtime/opencl/cl2_header.h" #include "mace/core/runtime/opencl/opencl_runtime.h" -#include "mace/kernels/conv_2d.h" #include "mace/utils/utils.h" namespace mace { @@ -22,22 +23,22 @@ void Conv1x1Naive(const Tensor *input, auto runtime = OpenCLRuntime::Get(); auto program = runtime->program(); - auto conv_2d = cl::KernelFunctor(program, "conv_2d_1x1_naive"); + auto conv_2d = + cl::KernelFunctor(program, "conv_2d_1x1_naive"); const index_t pixels = height * width; cl_int error; - conv_2d(cl::EnqueueArgs(runtime->command_queue(), - cl::NDRange(static_cast(batch), - static_cast(channels), - static_cast(pixels)), - cl::NDRange(1, 1, 128)), + conv_2d(cl::EnqueueArgs( + runtime->command_queue(), + cl::NDRange(static_cast(batch), static_cast(channels), + static_cast(pixels)), + cl::NDRange(1, 1, 128)), *(static_cast(input->buffer())), *(static_cast(filter->buffer())), *(static_cast(bias->buffer())), *(static_cast(output->buffer())), - static_cast(input_channels), - error); + static_cast(input_channels), error); MACE_CHECK(error == CL_SUCCESS); } @@ -53,31 +54,126 @@ void Conv1x1V2(const Tensor *input, auto runtime = OpenCLRuntime::Get(); auto program = runtime->program(); - auto conv_2d = cl::KernelFunctor(program, "conv_2d_1x1_v2"); const index_t pixels = height * width; const index_t channel_blocks = (channels + 3) / 4; const index_t pixel_blocks = (pixels + 3) / 4; + // TODO KernelFunctor has an extra clReleaseCommandQueue due to a copy + // TODO check wired clReleaseCommandQueue latency + // The KernelFunctor can cause segment faults in cb_retain_event + auto conv_2d_kernel = cl::Kernel(program, "conv_2d_1x1_v2"); + const uint32_t kwg_size = runtime->GetKernelMaxWorkGroupSize(conv_2d_kernel); + uint32_t idx = 0; + conv_2d_kernel.setArg(idx++, + *(static_cast(input->buffer()))); + conv_2d_kernel.setArg(idx++, + *(static_cast(filter->buffer()))); + conv_2d_kernel.setArg(idx++, + *(static_cast(bias->buffer()))); + conv_2d_kernel.setArg(idx++, *(static_cast(output->buffer()))); + conv_2d_kernel.setArg(idx++, static_cast(input_channels)); + conv_2d_kernel.setArg(idx++, static_cast(channels)); + conv_2d_kernel.setArg(idx++, static_cast(pixels)); + + auto command_queue = runtime->command_queue(); + cl_int error = command_queue.enqueueNDRangeKernel( + conv_2d_kernel, cl::NullRange, + cl::NDRange(static_cast(batch), static_cast(channel_blocks), + static_cast(pixel_blocks)), + cl::NDRange(1, 2, kwg_size / 2)); + MACE_CHECK(error == CL_SUCCESS, error); +} + +void Conv1x1V3(const Tensor *input, + const Tensor *filter, + const Tensor *bias, + Tensor *output) { + const index_t batch = output->shape()[0]; + const index_t channels = output->shape()[1]; + const index_t height = output->shape()[2]; + const index_t width = output->shape()[3]; + const index_t input_channels = input->shape()[1]; + + auto runtime = OpenCLRuntime::Get(); + auto program = runtime->program(); + + const index_t pixels = height * width; + const index_t pixel_blocks = (pixels + 3) / 4; + + const index_t channel_blocks = (channels + 3) / 4; + const index_t input_channel_blocks = (input_channels + 3) / 4; + + // FIXME temp hacking + static std::map input_image_map; + static std::map output_image_map; + cl::Image3D input_image; + cl::Image3D output_image; + auto input_iter = + input_image_map.find(reinterpret_cast(input->buffer())); + if (input_iter != input_image_map.end()) { + input_image = input_iter->second; + } else { + // The batch dimension is collapsed with channel + cl_int error; + cl::Image3D image = + cl::Image3D(OpenCLRuntime::Get()->context(), + CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, + cl::ImageFormat(CL_RGBA, CL_FLOAT), height, width, + batch * input_channel_blocks, 0, 0, nullptr, &error); + MACE_CHECK(error == CL_SUCCESS); + input_image = image; + input_image_map.clear(); + input_image_map.emplace(reinterpret_cast(input->buffer()), + image); + } + auto output_iter = + output_image_map.find(reinterpret_cast(output->buffer())); + if (output_iter != output_image_map.end()) { + output_image = output_iter->second; + } else { + cl_int error; + cl::Image3D image = + cl::Image3D(OpenCLRuntime::Get()->context(), + CL_MEM_WRITE_ONLY | CL_MEM_ALLOC_HOST_PTR, + cl::ImageFormat(CL_RGBA, CL_FLOAT), height, width, + batch * channel_blocks, 0, 0, nullptr, &error); + MACE_CHECK(error == CL_SUCCESS); + output_image = image; + output_image_map.clear(); + output_image_map.emplace(reinterpret_cast(output->buffer()), + image); + } + + auto conv_2d_kernel = cl::Kernel(program, "conv_2d_1x1_v3"); + const uint32_t kwg_size = runtime->GetKernelMaxWorkGroupSize(conv_2d_kernel); + + uint32_t idx = 0; + conv_2d_kernel.setArg(idx++, input_image); + conv_2d_kernel.setArg(idx++, + *(static_cast(filter->buffer()))); + conv_2d_kernel.setArg(idx++, + *(static_cast(bias->buffer()))); + conv_2d_kernel.setArg(idx++, output_image); + conv_2d_kernel.setArg(idx++, static_cast(batch)); + conv_2d_kernel.setArg(idx++, static_cast(input_channels)); + conv_2d_kernel.setArg(idx++, static_cast(channels)); + conv_2d_kernel.setArg(idx++, static_cast(height)); + conv_2d_kernel.setArg(idx++, static_cast(width)); + + auto command_queue = runtime->command_queue(); cl_int error; - conv_2d(cl::EnqueueArgs(runtime->command_queue(), - cl::NDRange(static_cast(batch), - static_cast(channel_blocks), - static_cast(pixel_blocks)), - cl::NDRange(1, 1, 256)), - *(static_cast(input->buffer())), - *(static_cast(filter->buffer())), - *(static_cast(bias->buffer())), - *(static_cast(output->buffer())), - static_cast(input_channels), - static_cast(channels), - static_cast(pixels), - error); - MACE_CHECK(error == CL_SUCCESS); + error = command_queue.enqueueNDRangeKernel( + conv_2d_kernel, cl::NullRange, + cl::NDRange(static_cast(channel_blocks), static_cast(height), + static_cast(width)), + cl::NDRange(1, 2, kwg_size / 2)); + MACE_CHECK(error == CL_SUCCESS, error); } -extern void Conv2dOpenclK1x1S1(const Tensor *input, const Tensor *filter, - const Tensor *bias, Tensor *output) { +extern void Conv2dOpenclK1x1S1(const Tensor *input, + const Tensor *filter, + const Tensor *bias, + Tensor *output) { const index_t batch = output->shape()[0]; const index_t height = output->shape()[2]; const index_t width = output->shape()[3]; -- GitLab