[deb_ffmpeg.git] / ffmpeg / cmdutils_opencl.c

/*
 * Copyright (C) 2013 Lenny Wang
 *
 * This file is part of FFmpeg.
 *
 * FFmpeg is free software; you can redistribute it and/or
 * modify it under the terms of the GNU Lesser General Public
 * License as published by the Free Software Foundation; either
 * version 2.1 of the License, or (at your option) any later version.
 *
 * FFmpeg is distributed in the hope that it will be useful,
 * but WITHOUT ANY WARRANTY; without even the implied warranty of
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
 * Lesser General Public License for more details.
 *
 * You should have received a copy of the GNU Lesser General Public
 * License along with FFmpeg; if not, write to the Free Software
 * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
 */

#include "libavutil/opt.h"
#include "libavutil/time.h"
#include "libavutil/log.h"
#include "libavutil/opencl.h"
#include "cmdutils.h"

typedef struct {
    int platform_idx;
    int device_idx;
    char device_name[64];
    int64_t runtime;
} OpenCLDeviceBenchmark;

const char *ocl_bench_source = AV_OPENCL_KERNEL(
inline unsigned char clip_uint8(int a)
{
    if (a & (~0xFF))
        return (-a)>>31;
    else
        return a;
}

kernel void unsharp_bench(
                    global unsigned char *src,
                    global unsigned char *dst,
                    global int *mask,
                    int width,
                    int height)
{
    int i, j, local_idx, lc_idx, sum = 0;
    int2 thread_idx, block_idx, global_idx, lm_idx;
    thread_idx.x = get_local_id(0);
    thread_idx.y = get_local_id(1);
    block_idx.x = get_group_id(0);
    block_idx.y = get_group_id(1);
    global_idx.x = get_global_id(0);
    global_idx.y = get_global_id(1);
    local uchar data[32][32];
    local int lc[128];

    for (i = 0; i <= 1; i++) {
        lm_idx.y = -8 + (block_idx.y + i) * 16 + thread_idx.y;
        lm_idx.y = lm_idx.y < 0 ? 0 : lm_idx.y;
        lm_idx.y = lm_idx.y >= height ? height - 1: lm_idx.y;
        for (j = 0; j <= 1; j++) {
            lm_idx.x = -8 + (block_idx.x + j) * 16 + thread_idx.x;
            lm_idx.x = lm_idx.x < 0 ? 0 : lm_idx.x;
            lm_idx.x = lm_idx.x >= width ? width - 1: lm_idx.x;
            data[i*16 + thread_idx.y][j*16 + thread_idx.x] = src[lm_idx.y*width + lm_idx.x];
        }
    }
    local_idx = thread_idx.y*16 + thread_idx.x;
    if (local_idx < 128)
        lc[local_idx] = mask[local_idx];
    barrier(CLK_LOCAL_MEM_FENCE);

    \n#pragma unroll\n
    for (i = -4; i <= 4; i++) {
        lm_idx.y = 8 + i + thread_idx.y;
        \n#pragma unroll\n
        for (j = -4; j <= 4; j++) {
            lm_idx.x = 8 + j + thread_idx.x;
            lc_idx = (i + 4)*8 + j + 4;
            sum += (int)data[lm_idx.y][lm_idx.x] * lc[lc_idx];
        }
    }
    int temp = (int)data[thread_idx.y + 8][thread_idx.x + 8];
    int res = temp + (((temp - (int)((sum + 1<<15) >> 16))) >> 16);
    if (global_idx.x < width && global_idx.y < height)
        dst[global_idx.x + global_idx.y*width] = clip_uint8(res);
}
);

#define OCLCHECK(method, ... )                                                 \
do {                                                                           \
    status = method(__VA_ARGS__);                                              \
    if (status != CL_SUCCESS) {                                                \
        av_log(NULL, AV_LOG_ERROR, # method " error '%s'\n",                   \
               av_opencl_errstr(status));                                      \
        ret = AVERROR_EXTERNAL;                                                \
        goto end;                                                              \
    }                                                                          \
} while (0)

#define CREATEBUF(out, flags, size)                                            \
do {                                                                           \
    out = clCreateBuffer(ext_opencl_env->context, flags, size, NULL, &status); \
    if (status != CL_SUCCESS) {                                                \
        av_log(NULL, AV_LOG_ERROR, "Could not create OpenCL buffer\n");        \
        ret = AVERROR_EXTERNAL;                                                \
        goto end;                                                              \
    }                                                                          \
} while (0)

static void fill_rand_int(int *data, int n)
{
    int i;
    srand(av_gettime());
    for (i = 0; i < n; i++)
        data[i] = rand();
}

#define OPENCL_NB_ITER 5
static int64_t run_opencl_bench(AVOpenCLExternalEnv *ext_opencl_env)
{
    int i, arg = 0, width = 1920, height = 1088;
    int64_t start, ret = 0;
    cl_int status;
    size_t kernel_len;
    char *inbuf;
    int *mask;
    int buf_size = width * height * sizeof(char);
    int mask_size = sizeof(uint32_t) * 128;

    cl_mem cl_mask, cl_inbuf, cl_outbuf;
    cl_kernel kernel = NULL;
    cl_program program = NULL;
    size_t local_work_size_2d[2] = {16, 16};
    size_t global_work_size_2d[2] = {(size_t)width, (size_t)height};

    if (!(inbuf = av_malloc(buf_size)) || !(mask = av_malloc(mask_size))) {
        av_log(NULL, AV_LOG_ERROR, "Out of memory\n");
        ret = AVERROR(ENOMEM);
        goto end;
    }
    fill_rand_int((int*)inbuf, buf_size/4);
    fill_rand_int(mask, mask_size/4);

    CREATEBUF(cl_mask, CL_MEM_READ_ONLY, mask_size);
    CREATEBUF(cl_inbuf, CL_MEM_READ_ONLY, buf_size);
    CREATEBUF(cl_outbuf, CL_MEM_READ_WRITE, buf_size);

    kernel_len = strlen(ocl_bench_source);
    program = clCreateProgramWithSource(ext_opencl_env->context, 1, &ocl_bench_source,
                                        &kernel_len, &status);
    if (status != CL_SUCCESS || !program) {
        av_log(NULL, AV_LOG_ERROR, "OpenCL unable to create benchmark program\n");
        ret = AVERROR_EXTERNAL;
        goto end;
    }
    status = clBuildProgram(program, 1, &(ext_opencl_env->device_id), NULL, NULL, NULL);
    if (status != CL_SUCCESS) {
        av_log(NULL, AV_LOG_ERROR, "OpenCL unable to build benchmark program\n");
        ret = AVERROR_EXTERNAL;
        goto end;
    }
    kernel = clCreateKernel(program, "unsharp_bench", &status);
    if (status != CL_SUCCESS) {
        av_log(NULL, AV_LOG_ERROR, "OpenCL unable to create benchmark kernel\n");
        ret = AVERROR_EXTERNAL;
        goto end;
    }

    OCLCHECK(clEnqueueWriteBuffer, ext_opencl_env->command_queue, cl_inbuf, CL_TRUE, 0,
             buf_size, inbuf, 0, NULL, NULL);
    OCLCHECK(clEnqueueWriteBuffer, ext_opencl_env->command_queue, cl_mask, CL_TRUE, 0,
             mask_size, mask, 0, NULL, NULL);
    OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_mem), &cl_inbuf);
    OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_mem), &cl_outbuf);
    OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_mem), &cl_mask);
    OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_int), &width);
    OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_int), &height);

    start = av_gettime_relative();
    for (i = 0; i < OPENCL_NB_ITER; i++)
        OCLCHECK(clEnqueueNDRangeKernel, ext_opencl_env->command_queue, kernel, 2, NULL,
                 global_work_size_2d, local_work_size_2d, 0, NULL, NULL);
    clFinish(ext_opencl_env->command_queue);
    ret = (av_gettime_relative() - start)/OPENCL_NB_ITER;
end:
    if (kernel)
        clReleaseKernel(kernel);
    if (program)
        clReleaseProgram(program);
    if (cl_inbuf)
        clReleaseMemObject(cl_inbuf);
    if (cl_outbuf)
        clReleaseMemObject(cl_outbuf);
    if (cl_mask)
        clReleaseMemObject(cl_mask);
    av_free(inbuf);
    av_free(mask);
    return ret;
}

static int compare_ocl_device_desc(const void *a, const void *b)
{
    return ((OpenCLDeviceBenchmark*)a)->runtime - ((OpenCLDeviceBenchmark*)b)->runtime;
}

int opt_opencl_bench(void *optctx, const char *opt, const char *arg)
{
    int i, j, nb_devices = 0, count = 0;
    int64_t score = 0;
    AVOpenCLDeviceList *device_list;
    AVOpenCLDeviceNode *device_node = NULL;
    OpenCLDeviceBenchmark *devices = NULL;
    cl_platform_id platform;

    av_opencl_get_device_list(&device_list);
    for (i = 0; i < device_list->platform_num; i++)
        nb_devices += device_list->platform_node[i]->device_num;
    if (!nb_devices) {
        av_log(NULL, AV_LOG_ERROR, "No OpenCL device detected!\n");
        return AVERROR(EINVAL);
    }
    if (!(devices = av_malloc_array(nb_devices, sizeof(OpenCLDeviceBenchmark)))) {
        av_log(NULL, AV_LOG_ERROR, "Could not allocate buffer\n");
        return AVERROR(ENOMEM);
    }

    for (i = 0; i < device_list->platform_num; i++) {
        for (j = 0; j < device_list->platform_node[i]->device_num; j++) {
            device_node = device_list->platform_node[i]->device_node[j];
            platform = device_list->platform_node[i]->platform_id;
            score = av_opencl_benchmark(device_node, platform, run_opencl_bench);
            if (score > 0) {
                devices[count].platform_idx = i;
                devices[count].device_idx = j;
                devices[count].runtime = score;
                strcpy(devices[count].device_name, device_node->device_name);
                count++;
            }
        }
    }
    qsort(devices, count, sizeof(OpenCLDeviceBenchmark), compare_ocl_device_desc);
    fprintf(stderr, "platform_idx\tdevice_idx\tdevice_name\truntime\n");
    for (i = 0; i < count; i++)
        fprintf(stdout, "%d\t%d\t%s\t%"PRId64"\n",
                devices[i].platform_idx, devices[i].device_idx,
                devices[i].device_name, devices[i].runtime);

    av_opencl_free_device_list(&device_list);
    av_free(devices);
    return 0;
}

int opt_opencl(void *optctx, const char *opt, const char *arg)
{
    char *key, *value;
    const char *opts = arg;
    int ret = 0;
    while (*opts) {
        ret = av_opt_get_key_value(&opts, "=", ":", 0, &key, &value);
        if (ret < 0)
            return ret;
        ret = av_opencl_set_option(key, value);
        if (ret < 0)
            return ret;
        if (*opts)
            opts++;
    }
    return ret;
}
Commit	Line	Data
	1	/*
	2	* Copyright (C) 2013 Lenny Wang
	3	*
	4	* This file is part of FFmpeg.
	5	*
	6	* FFmpeg is free software; you can redistribute it and/or
	7	* modify it under the terms of the GNU Lesser General Public
	8	* License as published by the Free Software Foundation; either
	9	* version 2.1 of the License, or (at your option) any later version.
	10	*
	11	* FFmpeg is distributed in the hope that it will be useful,
	12	* but WITHOUT ANY WARRANTY; without even the implied warranty of
	13	* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
	14	* Lesser General Public License for more details.
	15	*
	16	* You should have received a copy of the GNU Lesser General Public
	17	* License along with FFmpeg; if not, write to the Free Software
	18	* Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
	19	*/
	20
	21	#include "libavutil/opt.h"
	22	#include "libavutil/time.h"
	23	#include "libavutil/log.h"
	24	#include "libavutil/opencl.h"
	25	#include "cmdutils.h"
	26
	27	typedef struct {
	28	int platform_idx;
	29	int device_idx;
	30	char device_name[64];
	31	int64_t runtime;
	32	} OpenCLDeviceBenchmark;
	33
	34	const char *ocl_bench_source = AV_OPENCL_KERNEL(
	35	inline unsigned char clip_uint8(int a)
	36	{
	37	if (a & (~0xFF))
	38	return (-a)>>31;
	39	else
	40	return a;
	41	}
	42
	43	kernel void unsharp_bench(
	44	global unsigned char *src,
	45	global unsigned char *dst,
	46	global int *mask,
	47	int width,
	48	int height)
	49	{
	50	int i, j, local_idx, lc_idx, sum = 0;
	51	int2 thread_idx, block_idx, global_idx, lm_idx;
	52	thread_idx.x = get_local_id(0);
	53	thread_idx.y = get_local_id(1);
	54	block_idx.x = get_group_id(0);
	55	block_idx.y = get_group_id(1);
	56	global_idx.x = get_global_id(0);
	57	global_idx.y = get_global_id(1);
	58	local uchar data[32][32];
	59	local int lc[128];
	60
	61	for (i = 0; i <= 1; i++) {
	62	lm_idx.y = -8 + (block_idx.y + i) * 16 + thread_idx.y;
	63	lm_idx.y = lm_idx.y < 0 ? 0 : lm_idx.y;
	64	lm_idx.y = lm_idx.y >= height ? height - 1: lm_idx.y;
	65	for (j = 0; j <= 1; j++) {
	66	lm_idx.x = -8 + (block_idx.x + j) * 16 + thread_idx.x;
	67	lm_idx.x = lm_idx.x < 0 ? 0 : lm_idx.x;
	68	lm_idx.x = lm_idx.x >= width ? width - 1: lm_idx.x;
	69	data[i16 + thread_idx.y][j16 + thread_idx.x] = src[lm_idx.y*width + lm_idx.x];
	70	}
	71	}
	72	local_idx = thread_idx.y*16 + thread_idx.x;
	73	if (local_idx < 128)
	74	lc[local_idx] = mask[local_idx];
	75	barrier(CLK_LOCAL_MEM_FENCE);
	76
	77	\n#pragma unroll\n
	78	for (i = -4; i <= 4; i++) {
	79	lm_idx.y = 8 + i + thread_idx.y;
	80	\n#pragma unroll\n
	81	for (j = -4; j <= 4; j++) {
	82	lm_idx.x = 8 + j + thread_idx.x;
	83	lc_idx = (i + 4)*8 + j + 4;
	84	sum += (int)data[lm_idx.y][lm_idx.x] * lc[lc_idx];
	85	}
	86	}
	87	int temp = (int)data[thread_idx.y + 8][thread_idx.x + 8];
	88	int res = temp + (((temp - (int)((sum + 1<<15) >> 16))) >> 16);
	89	if (global_idx.x < width && global_idx.y < height)
	90	dst[global_idx.x + global_idx.y*width] = clip_uint8(res);
	91	}
	92	);
	93
	94	#define OCLCHECK(method, ... ) \
	95	do { \
	96	status = method(__VA_ARGS__); \
	97	if (status != CL_SUCCESS) { \
	98	av_log(NULL, AV_LOG_ERROR, # method " error '%s'\n", \
	99	av_opencl_errstr(status)); \
	100	ret = AVERROR_EXTERNAL; \
	101	goto end; \
	102	} \
	103	} while (0)
	104
	105	#define CREATEBUF(out, flags, size) \
	106	do { \
	107	out = clCreateBuffer(ext_opencl_env->context, flags, size, NULL, &status); \
	108	if (status != CL_SUCCESS) { \
	109	av_log(NULL, AV_LOG_ERROR, "Could not create OpenCL buffer\n"); \
	110	ret = AVERROR_EXTERNAL; \
	111	goto end; \
	112	} \
	113	} while (0)
	114
	115	static void fill_rand_int(int *data, int n)
	116	{
	117	int i;
	118	srand(av_gettime());
	119	for (i = 0; i < n; i++)
	120	data[i] = rand();
	121	}
	122
	123	#define OPENCL_NB_ITER 5
	124	static int64_t run_opencl_bench(AVOpenCLExternalEnv *ext_opencl_env)
	125	{
	126	int i, arg = 0, width = 1920, height = 1088;
	127	int64_t start, ret = 0;
	128	cl_int status;
	129	size_t kernel_len;
	130	char *inbuf;
	131	int *mask;
	132	int buf_size = width * height * sizeof(char);
	133	int mask_size = sizeof(uint32_t) * 128;
	134
	135	cl_mem cl_mask, cl_inbuf, cl_outbuf;
	136	cl_kernel kernel = NULL;
	137	cl_program program = NULL;
	138	size_t local_work_size_2d[2] = {16, 16};
	139	size_t global_work_size_2d[2] = {(size_t)width, (size_t)height};
	140
	141	if (!(inbuf = av_malloc(buf_size)) \|\| !(mask = av_malloc(mask_size))) {
	142	av_log(NULL, AV_LOG_ERROR, "Out of memory\n");
	143	ret = AVERROR(ENOMEM);
	144	goto end;
	145	}
	146	fill_rand_int((int*)inbuf, buf_size/4);
	147	fill_rand_int(mask, mask_size/4);
	148
	149	CREATEBUF(cl_mask, CL_MEM_READ_ONLY, mask_size);
	150	CREATEBUF(cl_inbuf, CL_MEM_READ_ONLY, buf_size);
	151	CREATEBUF(cl_outbuf, CL_MEM_READ_WRITE, buf_size);
	152
	153	kernel_len = strlen(ocl_bench_source);
	154	program = clCreateProgramWithSource(ext_opencl_env->context, 1, &ocl_bench_source,
	155	&kernel_len, &status);
	156	if (status != CL_SUCCESS \|\| !program) {
	157	av_log(NULL, AV_LOG_ERROR, "OpenCL unable to create benchmark program\n");
	158	ret = AVERROR_EXTERNAL;
	159	goto end;
	160	}
	161	status = clBuildProgram(program, 1, &(ext_opencl_env->device_id), NULL, NULL, NULL);
	162	if (status != CL_SUCCESS) {
	163	av_log(NULL, AV_LOG_ERROR, "OpenCL unable to build benchmark program\n");
	164	ret = AVERROR_EXTERNAL;
	165	goto end;
	166	}
	167	kernel = clCreateKernel(program, "unsharp_bench", &status);
	168	if (status != CL_SUCCESS) {
	169	av_log(NULL, AV_LOG_ERROR, "OpenCL unable to create benchmark kernel\n");
	170	ret = AVERROR_EXTERNAL;
	171	goto end;
	172	}
	173
	174	OCLCHECK(clEnqueueWriteBuffer, ext_opencl_env->command_queue, cl_inbuf, CL_TRUE, 0,
	175	buf_size, inbuf, 0, NULL, NULL);
	176	OCLCHECK(clEnqueueWriteBuffer, ext_opencl_env->command_queue, cl_mask, CL_TRUE, 0,
	177	mask_size, mask, 0, NULL, NULL);
	178	OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_mem), &cl_inbuf);
	179	OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_mem), &cl_outbuf);
	180	OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_mem), &cl_mask);
	181	OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_int), &width);
	182	OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_int), &height);
	183
	184	start = av_gettime_relative();
	185	for (i = 0; i < OPENCL_NB_ITER; i++)
	186	OCLCHECK(clEnqueueNDRangeKernel, ext_opencl_env->command_queue, kernel, 2, NULL,
	187	global_work_size_2d, local_work_size_2d, 0, NULL, NULL);
	188	clFinish(ext_opencl_env->command_queue);
	189	ret = (av_gettime_relative() - start)/OPENCL_NB_ITER;
	190	end:
	191	if (kernel)
	192	clReleaseKernel(kernel);
	193	if (program)
	194	clReleaseProgram(program);
	195	if (cl_inbuf)
	196	clReleaseMemObject(cl_inbuf);
	197	if (cl_outbuf)
	198	clReleaseMemObject(cl_outbuf);
	199	if (cl_mask)
	200	clReleaseMemObject(cl_mask);
	201	av_free(inbuf);
	202	av_free(mask);
	203	return ret;
	204	}
	205
	206	static int compare_ocl_device_desc(const void a, const void b)
	207	{
	208	return ((OpenCLDeviceBenchmark)a)->runtime - ((OpenCLDeviceBenchmark)b)->runtime;
	209	}
	210
	211	int opt_opencl_bench(void optctx, const char opt, const char *arg)
	212	{
	213	int i, j, nb_devices = 0, count = 0;
	214	int64_t score = 0;
	215	AVOpenCLDeviceList *device_list;
	216	AVOpenCLDeviceNode *device_node = NULL;
	217	OpenCLDeviceBenchmark *devices = NULL;
	218	cl_platform_id platform;
	219
	220	av_opencl_get_device_list(&device_list);
	221	for (i = 0; i < device_list->platform_num; i++)
	222	nb_devices += device_list->platform_node[i]->device_num;
	223	if (!nb_devices) {
	224	av_log(NULL, AV_LOG_ERROR, "No OpenCL device detected!\n");
	225	return AVERROR(EINVAL);
	226	}
	227	if (!(devices = av_malloc_array(nb_devices, sizeof(OpenCLDeviceBenchmark)))) {
	228	av_log(NULL, AV_LOG_ERROR, "Could not allocate buffer\n");
	229	return AVERROR(ENOMEM);
	230	}
	231
	232	for (i = 0; i < device_list->platform_num; i++) {
	233	for (j = 0; j < device_list->platform_node[i]->device_num; j++) {
	234	device_node = device_list->platform_node[i]->device_node[j];
	235	platform = device_list->platform_node[i]->platform_id;
	236	score = av_opencl_benchmark(device_node, platform, run_opencl_bench);
	237	if (score > 0) {
	238	devices[count].platform_idx = i;
	239	devices[count].device_idx = j;
	240	devices[count].runtime = score;
	241	strcpy(devices[count].device_name, device_node->device_name);
	242	count++;
	243	}
	244	}
	245	}
	246	qsort(devices, count, sizeof(OpenCLDeviceBenchmark), compare_ocl_device_desc);
	247	fprintf(stderr, "platform_idx\tdevice_idx\tdevice_name\truntime\n");
	248	for (i = 0; i < count; i++)
	249	fprintf(stdout, "%d\t%d\t%s\t%"PRId64"\n",
	250	devices[i].platform_idx, devices[i].device_idx,
	251	devices[i].device_name, devices[i].runtime);
	252
	253	av_opencl_free_device_list(&device_list);
	254	av_free(devices);
	255	return 0;
	256	}
	257
	258	int opt_opencl(void optctx, const char opt, const char *arg)
	259	{
	260	char key, value;
	261	const char *opts = arg;
	262	int ret = 0;
	263	while (*opts) {
	264	ret = av_opt_get_key_value(&opts, "=", ":", 0, &key, &value);
	265	if (ret < 0)
	266	return ret;
	267	ret = av_opencl_set_option(key, value);
	268	if (ret < 0)
	269	return ret;
	270	if (*opts)
	271	opts++;
	272	}
	273	return ret;
	274	}