1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
|
#include "utest_helper.hpp"
#include <string.h>
void builtin_kernel_block_motion_estimate_intel(void)
{
char* built_in_kernel_names;
size_t built_in_kernels_size;
cl_int err = CL_SUCCESS;
size_t ret_sz;
OCL_CALL (clGetDeviceInfo, device, CL_DEVICE_BUILT_IN_KERNELS, 0, 0, &built_in_kernels_size);
built_in_kernel_names = (char* )malloc(built_in_kernels_size * sizeof(char) );
OCL_CALL(clGetDeviceInfo, device, CL_DEVICE_BUILT_IN_KERNELS, built_in_kernels_size, (void*)built_in_kernel_names, &ret_sz);
OCL_ASSERT(ret_sz == built_in_kernels_size);
if (strstr(built_in_kernel_names, "block_motion_estimate_intel") == NULL)
{
free(built_in_kernel_names);
return;
}
cl_program built_in_prog = clCreateProgramWithBuiltInKernels(ctx, 1, &device, built_in_kernel_names, &err);
OCL_ASSERT(built_in_prog != NULL);
kernel = clCreateKernel(built_in_prog, "block_motion_estimate_intel", &err);
OCL_ASSERT(kernel != NULL);
cl_motion_estimation_desc_intel vmedesc = {CL_ME_MB_TYPE_16x16_INTEL, //0x0
CL_ME_SUBPIXEL_MODE_INTEGER_INTEL, //0x0
CL_ME_SAD_ADJUST_MODE_NONE_INTEL, //0x0
CL_ME_SEARCH_PATH_RADIUS_16_12_INTEL //0x5
};
cl_accelerator_intel accel = clCreateAcceleratorINTEL(ctx, CL_ACCELERATOR_TYPE_MOTION_ESTIMATION_INTEL,sizeof(cl_motion_estimation_desc_intel), &vmedesc, &err);
OCL_ASSERT(accel != NULL);
const size_t w = 71; //80
const size_t h = 41; //48
cl_image_format format;
cl_image_desc desc;
memset(&desc, 0x0, sizeof(cl_image_desc));
memset(&format, 0x0, sizeof(cl_image_format));
uint8_t* image_data1 = (uint8_t *)malloc(w * h); //src
uint8_t* image_data2 = (uint8_t *)malloc(w * h); //ref
for (size_t j = 0; j < h; j++) {
for (size_t i = 0; i < w; i++) {
if (i >= 32 && i <= 47 && j >= 16 && j <= 31)
image_data2[w * j + i] = image_data1[w * j + i] = 100;
else
image_data2[w * j + i] = image_data1[w * j + i] = 0;
}
}
format.image_channel_order = CL_R;
format.image_channel_data_type = CL_UNORM_INT8;
desc.image_type = CL_MEM_OBJECT_IMAGE2D;
desc.image_width = w;
desc.image_height = h;
desc.image_row_pitch = 0;
OCL_CREATE_IMAGE(buf[0], CL_MEM_COPY_HOST_PTR, &format, &desc, image_data1); //src
OCL_CREATE_IMAGE(buf[1], CL_MEM_COPY_HOST_PTR, &format, &desc, image_data2); //ref
const size_t mv = (80/16) * (48/16);
OCL_CREATE_BUFFER(buf[2], 0, mv * sizeof(int) * 4, NULL);
OCL_SET_ARG(0, sizeof(cl_accelerator_intel), &accel);
OCL_SET_ARG(1, sizeof(cl_mem), &buf[0]);
OCL_SET_ARG(2, sizeof(cl_mem), &buf[1]);
OCL_SET_ARG(3, sizeof(cl_mem), NULL);
OCL_SET_ARG(4, sizeof(cl_mem), &buf[2]);
OCL_SET_ARG(5, sizeof(cl_mem), NULL);
globals[0] = w;
globals[1] = h;
OCL_CALL(clEnqueueNDRangeKernel, queue, kernel, 2, NULL, globals, NULL, 0, NULL, NULL);
OCL_MAP_BUFFER(2);
short expected[] = {-64, -48,
-64, -48,
-64, -48,
-64, -48,
-64, -48,
-64, -48,
-64, -48,
0, 0,
0, -48,
-64, -48,
-64, -48,
-64, -48,
-64, -48,
0, -48,
-64, -48};
short* res = (short*)buf_data[2];
for (uint32_t j = 0; j < mv; ++j) {
OCL_ASSERT(res[j * 2 + 0] == expected[j * 2 + 0]);
OCL_ASSERT(res[j * 2 + 1] == expected[j * 2 + 1]);
}
OCL_UNMAP_BUFFER(2);
clReleaseAcceleratorINTEL(accel);
clReleaseKernel(kernel);
clReleaseProgram(built_in_prog);
free(built_in_kernel_names);
free(image_data1);
free(image_data2);
}
MAKE_UTEST_FROM_FUNCTION(builtin_kernel_block_motion_estimate_intel);
|
