/* oclscale.c
Copyright (c) 2003-2017 HandBrake Team
This file is part of the HandBrake source code
Homepage: .
It may be used under the terms of the GNU General Public License v2.
For full terms see the file COPYING file or visit http://www.gnu.org/licenses/gpl-2.0.html
Authors: Peng Gao
Li Cao
*/
#include
#include "common.h"
#include "opencl.h"
#include "openclwrapper.h"
#define FILTER_LEN 4
#define _A -0.5f
cl_float cubic(cl_float x)
{
if (x < 0)
x = -x;
if (x < 1)
return (_A + 2.0f) * (x * x * x) - (_A + 3.0f) * (x * x) + 0 + 1;
else if (x < 2)
return (_A) * (x * x * x) - (5.0f * _A) * (x * x) + (8.0f * _A) * x - (4.0f * _A);
else
return 0;
}
cl_float *hb_bicubic_weights(cl_float scale, int length)
{
cl_float *weights = (cl_float*) malloc(length * sizeof(cl_float) * 4);
int i; // C rocks
cl_float *out = weights;
for (i = 0; i < length; ++i)
{
cl_float x = i / scale;
cl_float dx = x - (int)x;
*out++ = cubic(-dx - 1.0f);
*out++ = cubic(-dx);
*out++ = cubic(-dx + 1.0f);
*out++ = cubic(-dx + 2.0f);
}
return weights;
}
int setupScaleWeights(cl_float xscale, cl_float yscale, int width, int height, hb_oclscale_t *os, KernelEnv *kenv);
/**
* executive scale using opencl
* get filter args
* create output buffer
* create horizontal filter buffer
* create vertical filter buffer
* create kernels
*/
int hb_ocl_scale_func( void **data, KernelEnv *kenv )
{
cl_int status;
cl_mem in_buf = data[0];
cl_mem out_buf = data[1];
int crop_top = (intptr_t)data[2];
int crop_bottom = (intptr_t)data[3];
int crop_left = (intptr_t)data[4];
int crop_right = (intptr_t)data[5];
cl_int in_frame_w = (intptr_t)data[6];
cl_int in_frame_h = (intptr_t)data[7];
cl_int out_frame_w = (intptr_t)data[8];
cl_int out_frame_h = (intptr_t)data[9];
hb_oclscale_t *os = data[10];
hb_buffer_t *in = data[11];
hb_buffer_t *out = data[12];
if (hb_ocl == NULL)
{
hb_error("hb_ocl_scale_func: OpenCL support not available");
return 0;
}
if (os->initialized == 0)
{
hb_log( "Scaling With OpenCL" );
if (kenv->isAMD != 0)
hb_log( "Using Zero Copy");
// create the block kernel
cl_int status;
os->m_kernel = hb_ocl->clCreateKernel(kenv->program, "frame_scale", &status);
os->initialized = 1;
}
{
// Use the new kernel
cl_event events[5];
int eventCount = 0;
if (kenv->isAMD == 0) {
status = hb_ocl->clEnqueueUnmapMemObject(kenv->command_queue,
in->cl.buffer, in->data, 0,
NULL, &events[eventCount++]);
status = hb_ocl->clEnqueueUnmapMemObject(kenv->command_queue,
out->cl.buffer, out->data, 0,
NULL, &events[eventCount++]);
}
cl_int srcPlaneOffset0 = in->plane[0].data - in->data;
cl_int srcPlaneOffset1 = in->plane[1].data - in->data;
cl_int srcPlaneOffset2 = in->plane[2].data - in->data;
cl_int srcRowWords0 = in->plane[0].stride;
cl_int srcRowWords1 = in->plane[1].stride;
cl_int srcRowWords2 = in->plane[2].stride;
cl_int dstPlaneOffset0 = out->plane[0].data - out->data;
cl_int dstPlaneOffset1 = out->plane[1].data - out->data;
cl_int dstPlaneOffset2 = out->plane[2].data - out->data;
cl_int dstRowWords0 = out->plane[0].stride;
cl_int dstRowWords1 = out->plane[1].stride;
cl_int dstRowWords2 = out->plane[2].stride;
if (crop_top != 0 || crop_bottom != 0 || crop_left != 0 || crop_right != 0) {
srcPlaneOffset0 += crop_left + crop_top * srcRowWords0;
srcPlaneOffset1 += crop_left / 2 + (crop_top / 2) * srcRowWords1;
srcPlaneOffset2 += crop_left / 2 + (crop_top / 2) * srcRowWords2;
in_frame_w = in_frame_w - crop_right - crop_left;
in_frame_h = in_frame_h - crop_bottom - crop_top;
}
cl_float xscale = (out_frame_w * 1.0f) / in_frame_w;
cl_float yscale = (out_frame_h * 1.0f) / in_frame_h;
setupScaleWeights(xscale, yscale, out_frame_w, out_frame_h, os, kenv);
HB_OCL_CHECK(hb_ocl->clSetKernelArg, os->m_kernel, 0, sizeof(cl_mem), &out_buf);
HB_OCL_CHECK(hb_ocl->clSetKernelArg, os->m_kernel, 1, sizeof(cl_mem), &in_buf);
HB_OCL_CHECK(hb_ocl->clSetKernelArg, os->m_kernel, 2, sizeof(cl_float), &xscale);
HB_OCL_CHECK(hb_ocl->clSetKernelArg, os->m_kernel, 3, sizeof(cl_float), &yscale);
HB_OCL_CHECK(hb_ocl->clSetKernelArg, os->m_kernel, 4, sizeof(cl_int), &srcPlaneOffset0);
HB_OCL_CHECK(hb_ocl->clSetKernelArg, os->m_kernel, 5, sizeof(cl_int), &srcPlaneOffset1);
HB_OCL_CHECK(hb_ocl->clSetKernelArg, os->m_kernel, 6, sizeof(cl_int), &srcPlaneOffset2);
HB_OCL_CHECK(hb_ocl->clSetKernelArg, os->m_kernel, 7, sizeof(cl_int), &dstPlaneOffset0);
HB_OCL_CHECK(hb_ocl->clSetKernelArg, os->m_kernel, 8, sizeof(cl_int), &dstPlaneOffset1);
HB_OCL_CHECK(hb_ocl->clSetKernelArg, os->m_kernel, 9, sizeof(cl_int), &dstPlaneOffset2);
HB_OCL_CHECK(hb_ocl->clSetKernelArg, os->m_kernel, 10, sizeof(cl_int), &srcRowWords0);
HB_OCL_CHECK(hb_ocl->clSetKernelArg, os->m_kernel, 11, sizeof(cl_int), &srcRowWords1);
HB_OCL_CHECK(hb_ocl->clSetKernelArg, os->m_kernel, 12, sizeof(cl_int), &srcRowWords2);
HB_OCL_CHECK(hb_ocl->clSetKernelArg, os->m_kernel, 13, sizeof(cl_int), &dstRowWords0);
HB_OCL_CHECK(hb_ocl->clSetKernelArg, os->m_kernel, 14, sizeof(cl_int), &dstRowWords1);
HB_OCL_CHECK(hb_ocl->clSetKernelArg, os->m_kernel, 15, sizeof(cl_int), &dstRowWords2);
HB_OCL_CHECK(hb_ocl->clSetKernelArg, os->m_kernel, 16, sizeof(cl_int), &in_frame_w);
HB_OCL_CHECK(hb_ocl->clSetKernelArg, os->m_kernel, 17, sizeof(cl_int), &in_frame_h);
HB_OCL_CHECK(hb_ocl->clSetKernelArg, os->m_kernel, 18, sizeof(cl_int), &out_frame_w);
HB_OCL_CHECK(hb_ocl->clSetKernelArg, os->m_kernel, 19, sizeof(cl_int), &out_frame_h);
HB_OCL_CHECK(hb_ocl->clSetKernelArg, os->m_kernel, 20, sizeof(cl_mem), &os->bicubic_x_weights);
HB_OCL_CHECK(hb_ocl->clSetKernelArg, os->m_kernel, 21, sizeof(cl_mem), &os->bicubic_y_weights);
size_t workOffset[] = { 0, 0, 0 };
size_t globalWorkSize[] = { 1, 1, 1 };
size_t localWorkSize[] = { 1, 1, 1 };
int xgroups = (out_frame_w + 63) / 64;
int ygroups = (out_frame_h + 15) / 16;
localWorkSize[0] = 64;
localWorkSize[1] = 1;
localWorkSize[2] = 1;
globalWorkSize[0] = xgroups * 64;
globalWorkSize[1] = ygroups;
globalWorkSize[2] = 3;
HB_OCL_CHECK(hb_ocl->clEnqueueNDRangeKernel, kenv->command_queue,
os->m_kernel, 3, workOffset, globalWorkSize, localWorkSize,
eventCount, eventCount == 0 ? NULL : &events[0], &events[eventCount]);
++eventCount;
if (kenv->isAMD == 0) {
in->data = hb_ocl->clEnqueueMapBuffer(kenv->command_queue, in->cl.buffer,
CL_FALSE, CL_MAP_READ|CL_MAP_WRITE,
0, in->alloc,
eventCount ? 1 : 0,
eventCount ? &events[eventCount - 1] : NULL,
&events[eventCount], &status);
out->data = hb_ocl->clEnqueueMapBuffer(kenv->command_queue, out->cl.buffer,
CL_FALSE, CL_MAP_READ|CL_MAP_WRITE,
0, out->alloc,
eventCount ? 1 : 0,
eventCount ? &events[eventCount - 1] : NULL,
&events[eventCount + 1], &status);
eventCount += 2;
}
hb_ocl->clFlush(kenv->command_queue);
hb_ocl->clWaitForEvents(eventCount, &events[0]);
int i;
for (i = 0; i < eventCount; ++i)
{
hb_ocl->clReleaseEvent(events[i]);
}
}
return 1;
}
int setupScaleWeights(cl_float xscale, cl_float yscale, int width, int height, hb_oclscale_t *os, KernelEnv *kenv)
{
cl_int status;
if (hb_ocl == NULL)
{
hb_error("setupScaleWeights: OpenCL support not available");
return 1;
}
if (os->xscale != xscale || os->width < width)
{
cl_float *xweights = hb_bicubic_weights(xscale, width);
HB_OCL_BUF_FREE (hb_ocl, os->bicubic_x_weights);
HB_OCL_BUF_CREATE(hb_ocl, os->bicubic_x_weights, CL_MEM_READ_ONLY,
sizeof(cl_float) * width * 4);
HB_OCL_CHECK(hb_ocl->clEnqueueWriteBuffer, kenv->command_queue, os->bicubic_x_weights,
CL_TRUE, 0, sizeof(cl_float) * width * 4, xweights, 0, NULL, NULL);
os->width = width;
os->xscale = xscale;
free(xweights);
}
if ((os->yscale != yscale) || (os->height < height))
{
cl_float *yweights = hb_bicubic_weights(yscale, height);
HB_OCL_BUF_FREE (hb_ocl, os->bicubic_y_weights);
HB_OCL_BUF_CREATE(hb_ocl, os->bicubic_y_weights, CL_MEM_READ_ONLY,
sizeof(cl_float) * height * 4);
HB_OCL_CHECK(hb_ocl->clEnqueueWriteBuffer, kenv->command_queue, os->bicubic_y_weights,
CL_TRUE, 0, sizeof(cl_float) * height * 4, yweights, 0, NULL, NULL);
os->height = height;
os->yscale = yscale;
free(yweights);
}
return 0;
}
/**
* function describe: this function is used to scaling video frame. it uses the gausi scaling algorithm
* parameter:
* inputFrameBuffer: the source video frame opencl buffer
* outputdata: the destination video frame buffer
* inputWidth: the width of the source video frame
* inputHeight: the height of the source video frame
* outputWidth: the width of destination video frame
* outputHeight: the height of destination video frame
*/
static int s_scale_init_flag = 0;
int do_scale_init()
{
if ( s_scale_init_flag==0 )
{
int st = hb_register_kernel_wrapper( "frame_scale", hb_ocl_scale_func );
if( !st )
{
hb_log( "register kernel[%s] failed", "frame_scale" );
return 0;
}
s_scale_init_flag++;
}
return 1;
}
int hb_ocl_scale(hb_buffer_t *in, hb_buffer_t *out, int *crop, hb_oclscale_t *os)
{
void *data[13];
if (do_scale_init() == 0)
return 0;
data[0] = in->cl.buffer;
data[1] = out->cl.buffer;
data[2] = (void*)(intptr_t)(crop[0]);
data[3] = (void*)(intptr_t)(crop[1]);
data[4] = (void*)(intptr_t)(crop[2]);
data[5] = (void*)(intptr_t)(crop[3]);
data[6] = (void*)(intptr_t)(in->f.width);
data[7] = (void*)(intptr_t)(in->f.height);
data[8] = (void*)(intptr_t)(out->f.width);
data[9] = (void*)(intptr_t)(out->f.height);
data[10] = os;
data[11] = in;
data[12] = out;
if( !hb_run_kernel( "frame_scale", data ) )
hb_log( "run kernel[%s] failed", "frame_scale" );
return 0;
}