Firefly rk3288 OpenCL

bunchen

Firefly rk3288采用Mali-T764的GPU，该GPU支持OpeCL 1.1。下面一步一步介绍android下OpenCL开发。

1.首先我们需要OpenCL的头文件和库。头文件可以在khronos的网站上下载：https://www.khronos.org/registry/cl/ 由于Mali-T764支持OpenCL 1.1所以我们下载1.1版的头文件并放对位置。新建include目录，把opencl.h放在include目录下，把cl_d3d10.h 、 cl_ext.h 、 cl_gl_ext.h 、 cl_gl.h 、 cl.h 、 cl.hpp 、 cl_platform.h放到include/CL目录下。在OpenCL的库文件在firefly rk3288源代码的device/rockchip/common/gpu/libMali-T760/libGLES_mali.so。

2.然后可以在这基础下开发了。下面给出一个打印OpenCL信息的Demo。由于libGLES_mali.so并没有放入的android的/system/lib路径下，所以我们要在Demo的Android.mk文件中要对libGLES_mali.so做预编译，Android.mk文件如下：

1. LOCAL_PATH := $(call my-dir)
   
2. 
   
3. 
   
4. include $(CLEAR_VARS)
   
5. LOCAL_PREBUILT_LIBS := \
   
6.     ../opencl/lib/libGLES_mali.so
   
7. include $(BUILD_MULTI_PREBUILT)
   
8. 
   
9. include $(CLEAR_VARS)
   
10. 
    
11. LOCAL_C_INCLUDES := \
    
12.         $(LOCAL_PATH)/../opencl/include
    
13. 
    
14. LOCAL_SRC_FILES := \
    
15.         opencl-info.c
    
16.         
    
17. LOCAL_SHARED_LIBRARIES := \
    
18.         libGLES_mali
    
19. 
    
20. LOCAL_MODULE := opencl-info
    
21. 
    
22. include $(BUILD_EXECUTABLE)

复制代码

在opencl-info.c中我们打印一些OpenCL的信息：

1. #include <stdio.h>
   
2. #include <stdlib.h>
   
3. #include <opencl.h>
   
4. 
   
5. void printPlatformInfo(cl_int ret ,cl_uint i, const char* part , char* buf)
   
6. {
   
7.     if(ret==CL_SUCCESS){
   
8.             printf("\t platform index=%d %s : %s\n",i,part,buf);
   
9.     }else if(ret==CL_INVALID_PLATFORM){
   
10.             printf("\t platform index=%d %s : invalid platform.\n",i,part);
    
11.         }else if(ret==CL_INVALID_VALUE){
    
12.             printf("\t platform index=%d %s : invalid value.\n",i,part);
    
13.         }else if(ret==CL_OUT_OF_HOST_MEMORY){
    
14.             printf("\t platform index=%d %s : out of host memory.\n",i,part);
    
15.         }else{
    
16.             printf("\t platform index=%d %s : i don't know why.\n",i,part);
    
17.         }
    
18. 
    
19. }
    
20. 
    
21. void printDevice(cl_device_id device)
    
22. {
    
23.     char buf[128];
    
24.     size_t size = 0 ;
    
25.     cl_device_type type;
    
26.     cl_int ret = CL_SUCCESS;
    
27.     printf("\t   ######################\n");
    
28.     ret = clGetDeviceInfo(device,CL_DEVICE_NAME,128,buf,&size);
    
29.     if(ret==CL_SUCCESS){
    
30.         printf("\t   device name=%s\n",buf);
    
31.     }else{
    
32.         printf("\t   get device name fail !\n");
    
33.     }
    
34.    
    
35.     ret = clGetDeviceInfo(device,CL_DEVICE_TYPE,sizeof(cl_device_type),&type,&size);
    
36.     if(ret==CL_SUCCESS){
    
37.         printf("\t   device type=%lu\n",type);
    
38.     }else{
    
39.         printf("\t   get device type fail !\n");
    
40.     }
    
41.    
    
42.    
    
43.     ret = clGetDeviceInfo(device,CL_DEVICE_VENDOR,128,buf,&size);
    
44.     if(ret==CL_SUCCESS){
    
45.         printf("\t   device vendor=%s\n",buf);
    
46.     }else{
    
47.         printf("\t   get device vendor fail !\n");
    
48.     }
    
49.     ret = clGetDeviceInfo(device,CL_DRIVER_VERSION,128,buf,&size);
    
50.     if(ret==CL_SUCCESS){
    
51.         printf("\t   device version=%s\n",buf);
    
52.     }else{
    
53.         printf("\t   get device version fail !\n");
    
54.     }
    
55.     ret = clGetDeviceInfo(device,CL_DEVICE_PROFILE,128,buf,&size);
    
56.     if(ret==CL_SUCCESS){
    
57.         printf("\t   device profile=%s\n",buf);
    
58.     }else{
    
59.         printf("\t   get device profile fail !\n");
    
60.     }
    
61.     printf("\t   ######################\n");
    
62. }
    
63. 
    
64. void printDevices(cl_platform_id platform , cl_device_type device_type)
    
65. {
    
66.     cl_device_id devices[8];
    
67.     cl_uint i ;
    
68.     cl_uint num = 0 ;
    
69.     cl_int ret = clGetDeviceIDs(platform,device_type,8,devices,&num);
    
70.     if(ret==CL_SUCCESS){
    
71.         printf("\t  getDeviceIDs success ! num=%d\n",num);
    
72.         
    
73.         for(i=0;i<num;i++){
    
74.             printDevice(devices[i]);
    
75.         }
    
76.     }else{
    
77.         printf("\t  getDeviceIDs fail !\n");
    
78.     }
    
79.    
    
80. }
    
81. 
    
82. void printOpenCL()
    
83. {
    
84.         cl_platform_id platforms[8];
    
85.         cl_uint num = 0 ;
    
86.         cl_uint i=0;
    
87.         cl_int err = clGetPlatformIDs(8,platforms,&num);
    
88.         if(err==CL_SUCCESS){
    
89.                 printf("printOpenCL got %d platforms : \n",num);
    
90.                 for(i=0;i<num;i++){
    
91.                     char buf[128];
    
92.                     size_t size = 0 ;
    
93.                     cl_int ret = 0 ;
    
94.                     
    
95.                     ret = clGetPlatformInfo(platforms[i],CL_PLATFORM_PROFILE,128,buf,&size);
    
96.                     printPlatformInfo(ret,i,"profile",buf);
    
97.                
    
98.                 ret = clGetPlatformInfo(platforms[i],CL_PLATFORM_VERSION,128,buf,&size);
    
99.                     printPlatformInfo(ret,i,"version",buf);
    
100.                     
     
101.                 ret = clGetPlatformInfo(platforms[i],CL_PLATFORM_NAME,128,buf,&size);
     
102.                 printPlatformInfo(ret,i,"name",buf);
     
103.                     
     
104.                     ret = clGetPlatformInfo(platforms[i],CL_PLATFORM_VENDOR,128,buf,&size);
     
105.                 printPlatformInfo(ret,i,"vendor",buf);
     
106.                
     
107.                 ret = clGetPlatformInfo(platforms[i],CL_PLATFORM_EXTENSIONS,128,buf,&size);
     
108.                 printPlatformInfo(ret,i,"extensions",buf);
     
109.                
     
110.                 printf("Device cpu:\n");
     
111.                 printDevices(platforms[i],CL_DEVICE_TYPE_CPU);
     
112.                
     
113.                 printf("Device gpu:\n");
     
114.                 printDevices(platforms[i],CL_DEVICE_TYPE_GPU);
     
115.                
     
116.                 printf("Device accelerator:\n");
     
117.                 printDevices(platforms[i],CL_DEVICE_TYPE_ACCELERATOR);
     
118.                
     
119.                 printf("Device default:\n");
     
120.                 printDevices(platforms[i],CL_DEVICE_TYPE_DEFAULT);
     
121.                
     
122.                 printf("Device all :\n");
     
123.                 printDevices(platforms[i],CL_DEVICE_TYPE_ALL);
     
124.                
     
125.                 }
     
126.         }else if(err==CL_INVALID_VALUE){
     
127.                 printf("printOpenCL invalid value.\n");
     
128.         }else if(err==CL_OUT_OF_HOST_MEMORY){
     
129.                 printf("printOpenCL out of host memory.\n");
     
130.         }else{
     
131.                 printf("printOpenCL i don't know why.\n");
     
132.         }
     
133. 
     
134. }
     
135. 
     
136. 
     
137. int main(){
     
138.         printOpenCL();
     
139.         return 0;
     
140. }

复制代码

这里打印的信息请参考opencl文档 : https://www.khronos.org/registry/cl/specs/opencl-1.1.pdf
把编译好的opencl-info push到/system/bin目录下，把libGLES_mali.so push到/system/lib目录下，就可以在adb shell中运行opencl-info了。


从打印的信息可以看到，rk3288支持OpenCL 1.1，支持的设备是gpu的Mali-T764。

bunchen

这里贴出一个简单的向量加法的例子，和容易出错的地方。先上代码：

1. #include <stdio.h>
   
2. #include <stdlib.h>
   
3. #include <opencl.h>
   
4. 
   
5. void printDeviceWorkInfo(cl_device_id device)
   
6. {
   
7.     cl_uint nMaxComputeUnits = 0;
   
8.     cl_uint nMaxWorkItemDims = 0;
   
9.     cl_uint i = 0;
   
10.     size_t* nMaxWorkItemSizes = NULL;
    
11.     size_t nMaxWorkGroupSize = 0;
    
12.     size_t size = 0 ;
    
13.     cl_int err ;
    
14.     err = clGetDeviceInfo(device,CL_DEVICE_MAX_COMPUTE_UNITS,sizeof(cl_uint),&nMaxComputeUnits,&size);
    
15.     if(err==CL_SUCCESS){
    
16.         printf("nMaxComputeUnits=%d\n",nMaxComputeUnits);
    
17.     }
    
18.    
    
19.     err = clGetDeviceInfo(device,CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS,sizeof(cl_uint),&nMaxWorkItemDims,&size);
    
20.     if(err==CL_SUCCESS){
    
21.         printf("nMaxWorkItemDims=%d\n",nMaxWorkItemDims);
    
22.         nMaxWorkItemSizes = (size_t*)malloc(sizeof(size_t)*nMaxWorkItemDims);
    
23.         err = clGetDeviceInfo(device,CL_DEVICE_MAX_WORK_ITEM_SIZES,sizeof(size_t)*nMaxWorkItemDims,nMaxWorkItemSizes,&size);
    
24.         if(err==CL_SUCCESS){
    
25.             for(i=0;i<nMaxWorkItemDims;i++){
    
26.                 printf("nMaxWorkItemSizes[%d]=%d\n",i,nMaxWorkItemSizes[i]);
    
27.             }
    
28.         }
    
29.         free(nMaxWorkItemSizes);
    
30.     }
    
31.    
    
32.     err = clGetDeviceInfo(device,CL_DEVICE_MAX_WORK_GROUP_SIZE,sizeof(size_t),&nMaxWorkGroupSize,&size);
    
33.     if(err==CL_SUCCESS){
    
34.         printf("nMaxWorkGroupSize=%d\n",nMaxWorkGroupSize);
    
35.     }
    
36. }
    
37. 
    
38. const char* program_src = ""
    
39. "__kernel void vector_add_gpu (__global const float* src_a,\n"
    
40. "   __global const float* src_b,\n"
    
41. "   __global float* res,\n"
    
42. "   const int num)\n"
    
43. "{\n"
    
44. "   int idx = get_global_id(0);\n"
    
45. "   if(idx<num){"
    
46. "       res[idx]=src_a[idx]+src_b[idx];\n"
    
47. "   }\n"
    
48. "}\n"
    
49. ;
    
50. 
    
51. static const cl_int vect_len = 10000000;
    
52. 
    
53. static float* vect_a = NULL ;
    
54. static float* vect_b = NULL ;
    
55. static float* vect_c = NULL ;
    
56. 
    
57. void initVects()
    
58. {
    
59.     cl_int i;
    
60.     vect_a = (float*)malloc(sizeof(float)*vect_len);
    
61.     vect_b = (float*)malloc(sizeof(float)*vect_len);
    
62.     vect_c = (float*)malloc(sizeof(float)*vect_len);
    
63.     for(i=0;i<vect_len;i++){
    
64.         vect_a[i]=(float)rand()/RAND_MAX;
    
65.         vect_b[i]=(float)rand()/RAND_MAX;
    
66.         vect_c[i]=0.0f;
    
67.     }
    
68. }
    
69. 
    
70. void printVects()
    
71. {
    
72.     cl_int i;
    
73.     if(vect_a && vect_b && vect_c){
    
74.         printf("######################\n");
    
75.         for(i=0;i<4;i++){
    
76.             printf("%08d : %f,%f,%f\n",i,vect_a[i],vect_b[i],vect_c[i]);
    
77.         }
    
78.         printf("    ...    \n");
    
79.         for(i=vect_len-4;i<vect_len;i++){
    
80.             printf("%08d : %f,%f,%f\n",i,vect_a[i],vect_b[i],vect_c[i]);
    
81.         }
    
82.         printf("######################\n");
    
83.     }
    
84. }
    
85. 
    
86. void releaseVects()
    
87. {
    
88.     if(vect_a){
    
89.         free(vect_a);
    
90.         vect_a=NULL;
    
91.     }
    
92.     if(vect_b){
    
93.         free(vect_b);
    
94.         vect_b=NULL;
    
95.     }
    
96.     if(vect_c){
    
97.         free(vect_c);
    
98.         vect_c=NULL;
    
99.     }
    
100. }
     
101. 
     
102. size_t shrRoundUp(size_t f , size_t s)
     
103. {
     
104.     return (s+f-1)/f*f;
     
105. }
     
106. 
     
107. 
     
108. 
     
109. void test()
     
110. {
     
111.     cl_int error = 0 ;
     
112.     cl_platform_id platform;
     
113.     cl_context context;
     
114.     cl_command_queue queue;
     
115.     cl_device_id device;
     
116.     cl_mem inbuf_a ;
     
117.     cl_mem inbuf_b ;
     
118.     cl_mem outbuf_r ;
     
119.     const cl_int size = vect_len;
     
120.     cl_int i ;
     
121.     const size_t mem_size = sizeof(float)*size;
     
122.     size_t program_len = strlen(program_src);
     
123.     char build_log[1024];
     
124.     size_t log_size;
     
125.     size_t local_ws;
     
126.     size_t global_ws;
     
127.     cl_kernel vector_add_kernel;
     
128.    
     
129.     error = clGetPlatformIDs(1,&platform,NULL);
     
130.     if(error != CL_SUCCESS){
     
131.         printf("get platform id fail !\n");
     
132.         exit(1);
     
133.     }
     
134.    
     
135.     error = clGetDeviceIDs(platform,CL_DEVICE_TYPE_GPU,1,&device,NULL);
     
136.     if(error != CL_SUCCESS){
     
137.         printf("get gpu device fail !\n");
     
138.         exit(1);
     
139.     }
     
140.    
     
141.     printDeviceWorkInfo(device);
     
142.    
     
143.     cl_context_properties properties[]={
     
144.         CL_CONTEXT_PLATFORM,
     
145.         (cl_context_properties)platform,
     
146.         0
     
147.     };
     
148.    
     
149.     // 这里要配置properties
     
150.     context = clCreateContext(properties,1,&device,NULL,NULL,&error);
     
151.     if(error != CL_SUCCESS){
     
152.         printf("create context fail !\n");
     
153.         exit(1);
     
154.     }
     
155.    
     
156.     queue = clCreateCommandQueue(context,device,CL_QUEUE_PROFILING_ENABLE,&error);
     
157.     if(error != CL_SUCCESS){
     
158.         printf("create command queue fail !\n");
     
159.         exit(1);
     
160.     }
     
161.    
     
162.     initVects();
     
163.     printVects();
     
164.    
     
165.     inbuf_a = clCreateBuffer(context,CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,mem_size,vect_a,&error);
     
166.     if(error!=CL_SUCCESS){
     
167.         printf("create buffer inbuf_a fail !\n");
     
168.         exit(1);
     
169.     }
     
170.     inbuf_b = clCreateBuffer(context,CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,mem_size,vect_b,&error);
     
171.     if(error!=CL_SUCCESS){
     
172.         printf("create buffer inbuf_b fail !\n");
     
173.         exit(1);
     
174.     }
     
175.     outbuf_r = clCreateBuffer(context,CL_MEM_WRITE_ONLY,mem_size,NULL,&error);
     
176.     if(error!=CL_SUCCESS){
     
177.         printf("create buffer outbuf_r fail !\n");
     
178.         exit(1);
     
179.     }
     
180.    
     
181.     cl_program program = clCreateProgramWithSource(context,1,&program_src,&program_len,&error);
     
182.     if(error!=CL_SUCCESS){
     
183.         printf("create program fail !\n");
     
184.         exit(1);
     
185.     }
     
186.     error = clBuildProgram(program,1,&device,NULL,NULL,NULL);
     
187.     if(error!=CL_SUCCESS){
     
188.         printf("build program fail !\n");
     
189.         clGetProgramBuildInfo(program,device,CL_PROGRAM_BUILD_LOG,1024,build_log,&log_size);
     
190.         printf("build_log : %s\n",build_log);
     
191.         exit(1);
     
192.     }
     
193.    
     
194.     vector_add_kernel = clCreateKernel(program,"vector_add_gpu",&error);
     
195.     if(error!=CL_SUCCESS){
     
196.         printf("create kernel fail !\n");
     
197.         exit(1);
     
198.     }
     
199.    
     
200.     error = clSetKernelArg(vector_add_kernel,0,sizeof(cl_mem),&inbuf_a);
     
201.     error |= clSetKernelArg(vector_add_kernel,1,sizeof(cl_mem),&inbuf_b);
     
202.     error |= clSetKernelArg(vector_add_kernel,2,sizeof(cl_mem),&outbuf_r);
     
203.     error |= clSetKernelArg(vector_add_kernel,3,sizeof(cl_int),&size);
     
204.     if(error!=CL_SUCCESS){
     
205.         printf("set kernel arg fail !\n");
     
206.         exit(1);
     
207.     }
     
208.    
     
209.     local_ws = 256;  //我们使用一维的clEnqueueNDRangeKernel，这里local_ws选择nMaxWorkItemSizes[0]=256
     
210.     global_ws = shrRoundUp(local_ws,size); //这里是线程总数，应该是local_ws的倍数。
     
211.     printf("local_ws=%d,global_ws=%d\n",local_ws,global_ws);
     
212.    
     
213.     error = clEnqueueNDRangeKernel(queue,vector_add_kernel,1,NULL,&global_ws,&local_ws,0,NULL,NULL);
     
214.     if(error!=CL_SUCCESS){
     
215.         printf("enqueue kernel fail !\n");
     
216.         exit(1);
     
217.     }
     
218.    
     
219.     clEnqueueReadBuffer(queue,outbuf_r,CL_TRUE,0,mem_size,vect_c,0,NULL,NULL);
     
220.     printVects();
     
221.    
     
222.     clReleaseKernel(vector_add_kernel);
     
223.     clReleaseProgram(program);
     
224.     clReleaseCommandQueue(queue);
     
225.     clReleaseContext(context);
     
226.     clReleaseMemObject(inbuf_a);
     
227.     clReleaseMemObject(inbuf_b);
     
228.     clReleaseMemObject(outbuf_r);
     
229.     releaseVects();
     
230. }
     
231. 
     
232. int main(){
     
233.         test();
     
234.         return 0;
     
235. }

复制代码

这里说一下容易出错的地方：
（1）clCreateContext，网上有些例子把第一个参数置成0，经试验这样不行，要设置properties。
（2）clEnqueueNDRangeKernel的global_work_size和local_work_size。我们使用一维的clEnqueueNDRangeKernel，这里local_work_size选择nMaxWorkItemSizes[0]=256，global_work_size是线程总数，应该是local_ws的倍数。
一维情况下：


二维情况下：



（3）kernel代码中
size_t get_global_id (uint dimindx)
返回这个线程的global_id,参数是维度索引，我们一维的情况下参数是0。


最后运行结果：


在加一个nv12转rgb的例子：

1. #include <stdio.h>
   
2. #include <stdlib.h>
   
3. #include <opencl.h>
   
4. #include <sys/time.h>
   
5. 
   
6. void printDeviceWorkInfo(cl_device_id device)
   
7. {
   
8.     cl_uint nMaxComputeUnits = 0;
   
9.     cl_uint nMaxWorkItemDims = 0;
   
10.     cl_uint i = 0;
    
11.     size_t* nMaxWorkItemSizes = NULL;
    
12.     size_t nMaxWorkGroupSize = 0;
    
13.     size_t size = 0 ;
    
14.     cl_int err ;
    
15.     err = clGetDeviceInfo(device,CL_DEVICE_MAX_COMPUTE_UNITS,sizeof(cl_uint),&nMaxComputeUnits,&size);
    
16.     if(err==CL_SUCCESS){
    
17.         printf("nMaxComputeUnits=%d\n",nMaxComputeUnits);
    
18.     }
    
19.    
    
20.     err = clGetDeviceInfo(device,CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS,sizeof(cl_uint),&nMaxWorkItemDims,&size);
    
21.     if(err==CL_SUCCESS){
    
22.         printf("nMaxWorkItemDims=%d\n",nMaxWorkItemDims);
    
23.         nMaxWorkItemSizes = (size_t*)malloc(sizeof(size_t)*nMaxWorkItemDims);
    
24.         err = clGetDeviceInfo(device,CL_DEVICE_MAX_WORK_ITEM_SIZES,sizeof(size_t)*nMaxWorkItemDims,nMaxWorkItemSizes,&size);
    
25.         if(err==CL_SUCCESS){
    
26.             for(i=0;i<nMaxWorkItemDims;i++){
    
27.                 printf("nMaxWorkItemSizes[%d]=%d\n",i,nMaxWorkItemSizes[i]);
    
28.             }
    
29.         }
    
30.         free(nMaxWorkItemSizes);
    
31.     }
    
32.    
    
33.     err = clGetDeviceInfo(device,CL_DEVICE_MAX_WORK_GROUP_SIZE,sizeof(size_t),&nMaxWorkGroupSize,&size);
    
34.     if(err==CL_SUCCESS){
    
35.         printf("nMaxWorkGroupSize=%d\n",nMaxWorkGroupSize);
    
36.     }
    
37. }
    
38. 
    
39. const char* program_src = ""
    
40. "__kernel void nv12_to_rgb (__global const unsigned char* nv12,\n"
    
41. "   __global unsigned char* rgb,\n"
    
42. "   const int width,\n"
    
43. "   const int height)\n"
    
44. "{\n"
    
45. "   int idi = get_global_id(0);\n"
    
46. "   int idj = get_global_id(1);\n"
    
47. "   int k = 0 ;\n"
    
48. "   int y = 0 ;\n"
    
49. "   int u = 0 ;\n"
    
50. "   int v = 0 ;\n"
    
51. "   int t = 0 ;\n"
    
52. "   if(idi<width && idj<height){\n"
    
53. "       y=nv12[idj*width+idi];\n"
    
54. "       k=width*height+((idj/2)*(width/2)+idi/2)*2;\n"
    
55. "       u=nv12[k];\n"
    
56. "       v=nv12[k+1];\n"
    
57. "       k=(idj*width+idi)*3;\n"
    
58. "       t=(int)(y+1.370705*v-175.4502);\n"
    
59. "       rgb[k]=t>255?255:t<0?0:t;\n"
    
60. "       t=(int)(y-0.698001*v-0.337633*u+132.56124);\n"
    
61. "       rgb[k+1]=t>255?255:t<0?0:t;\n"
    
62. "       t=(int)(y+1.732446*u-221.7531);"
    
63. "       rgb[k+2]=t>255?255:t<0?0:t;\n"
    
64. "   }\n"
    
65. "}\n"
    
66. ;
    
67. 
    
68. size_t shrRoundUp(size_t f , size_t s)
    
69. {
    
70.     return (s+f-1)/f*f;
    
71. }
    
72. 
    
73. void nv12_to_rgb(uint8_t* rgb , uint8_t* nv12 , int width , int height)
    
74. {
    
75.     cl_int error = 0 ;
    
76.     cl_platform_id platform;
    
77.     cl_context context;
    
78.     cl_command_queue queue;
    
79.     cl_device_id device;
    
80.     cl_mem inbuf_nv12 ;
    
81.     cl_mem outbuf_rgb ;
    
82.     cl_int i ;
    
83.     size_t program_len = strlen(program_src);
    
84.     char build_log[1024];
    
85.     size_t log_size;
    
86.     size_t local_ws;
    
87.     size_t global_ws;
    
88.     cl_kernel nv12_to_rgb;
    
89.    
    
90.     cl_int nv12_size = width*height*3/2;
    
91.     cl_int rgb_size = width*height*3;
    
92.    
    
93.     error = clGetPlatformIDs(1,&platform,NULL);
    
94.     if(error != CL_SUCCESS){
    
95.         printf("get platform id fail !\n");
    
96.         exit(1);
    
97.     }
    
98.    
    
99.     error = clGetDeviceIDs(platform,CL_DEVICE_TYPE_GPU,1,&device,NULL);
    
100.     if(error != CL_SUCCESS){
     
101.         printf("get gpu device fail !\n");
     
102.         exit(1);
     
103.     }
     
104.    
     
105.     printDeviceWorkInfo(device);
     
106.    
     
107.    
     
108.     cl_context_properties properties[]={
     
109.         CL_CONTEXT_PLATFORM,
     
110.         (cl_context_properties)platform,
     
111.         0
     
112.     };
     
113.    
     
114.     // 这里要配置properties
     
115.     context = clCreateContext(properties,1,&device,NULL,NULL,&error);
     
116.     if(error != CL_SUCCESS){
     
117.         printf("create context fail !\n");
     
118.         exit(1);
     
119.     }
     
120.    
     
121.     queue = clCreateCommandQueue(context,device,CL_QUEUE_PROFILING_ENABLE,&error);
     
122.     if(error != CL_SUCCESS){
     
123.         printf("create command queue fail !\n");
     
124.         exit(1);
     
125.     }
     
126.    
     
127.     cl_program program = clCreateProgramWithSource(context,1,&program_src,&program_len,&error);
     
128.     if(error!=CL_SUCCESS){
     
129.         printf("create program fail !\n");
     
130.         exit(1);
     
131.     }
     
132.     error = clBuildProgram(program,1,&device,NULL,NULL,NULL);
     
133.     if(error!=CL_SUCCESS){
     
134.         printf("build program fail !\n");
     
135.         clGetProgramBuildInfo(program,device,CL_PROGRAM_BUILD_LOG,1024,build_log,&log_size);
     
136.         printf("build_log : %s\n",build_log);
     
137.         exit(1);
     
138.     }
     
139.    
     
140.     nv12_to_rgb = clCreateKernel(program,"nv12_to_rgb",&error);
     
141.     if(error!=CL_SUCCESS){
     
142.         printf("create kernel fail !\n");
     
143.         exit(1);
     
144.     }
     
145. 
     
146.     struct timeval val;
     
147.     gettimeofday(&val,NULL);
     
148.     long t1 = val.tv_sec*1000000 + val.tv_usec;
     
149. 
     
150.     inbuf_nv12 = clCreateBuffer(context,CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,nv12_size,nv12,&error);
     
151.     if(error!=CL_SUCCESS){
     
152.         printf("create buffer inbuf_a fail !\n");
     
153.         exit(1);
     
154.     }
     
155.     outbuf_rgb = clCreateBuffer(context,CL_MEM_WRITE_ONLY,rgb_size,NULL,&error);
     
156.     if(error!=CL_SUCCESS){
     
157.         printf("create buffer outbuf_r fail !\n");
     
158.         exit(1);
     
159.     }
     
160.    
     
161.     error = clSetKernelArg(nv12_to_rgb,0,sizeof(cl_mem),&inbuf_nv12);
     
162.     error |= clSetKernelArg(nv12_to_rgb,1,sizeof(cl_mem),&outbuf_rgb);
     
163.     error |= clSetKernelArg(nv12_to_rgb,2,sizeof(cl_int),&width);
     
164.     error |= clSetKernelArg(nv12_to_rgb,3,sizeof(cl_int),&height);
     
165.    
     
166.     if(error!=CL_SUCCESS){
     
167.         printf("set kernel arg fail !\n");
     
168.         exit(1);
     
169.     }
     
170.    
     
171.     size_t lws[2] = {16,16};
     
172.     size_t gws[2];
     
173.     gws[0] = shrRoundUp(lws[0],width);
     
174.     gws[1] = shrRoundUp(lws[1],height);
     
175.     printf("lws={%d,%d},gws={%d,%d}\n",lws[0],lws[1],gws[0],gws[1]);
     
176.    
     
177.     error = clEnqueueNDRangeKernel(queue,nv12_to_rgb,2,NULL,gws,lws,0,NULL,NULL);
     
178.     if(error!=CL_SUCCESS){
     
179.         printf("enqueue kernel fail !\n");
     
180.         exit(1);
     
181.     }
     
182.    
     
183.     clEnqueueReadBuffer(queue,outbuf_rgb,CL_TRUE,0,rgb_size,rgb,0,NULL,NULL);
     
184.    
     
185. 
     
186.     clReleaseMemObject(inbuf_nv12);
     
187.     clReleaseMemObject(outbuf_rgb);   
     
188. 
     
189.    
     
190.     gettimeofday(&val,NULL);
     
191.     long t2 = val.tv_sec*1000000 + val.tv_usec;
     
192. 
     
193.     printf("nv12_to_rgb spend %ld (us)\n",t2-t1);
     
194. 
     
195.     clReleaseKernel(nv12_to_rgb);
     
196.     clReleaseProgram(program);
     
197.     clReleaseCommandQueue(queue);
     
198.     clReleaseContext(context);
     
199. }
     
200. 
     
201. int main(){
     
202.     uint8_t* nv12;
     
203.     uint8_t* rgb;
     
204.     int width=1920;
     
205.     int height=1080;
     
206.     int i,j,k;
     
207.     nv12=(uint8_t*)malloc(width*height*3/2);
     
208.     rgb=(uint8_t*)malloc(width*height*3);
     
209.    
     
210.     memset(nv12,0,width*height*3/2);
     
211.    
     
212.    
     
213.     memset(rgb,0,width*height*3);
     
214.     struct timeval val;
     
215.     gettimeofday(&val,NULL);
     
216.     long t1 = val.tv_sec*1000000 + val.tv_usec;
     
217.         nv12_to_rgb(rgb,nv12,width,height);
     
218.     gettimeofday(&val,NULL);
     
219.     long t2 = val.tv_sec*1000000 + val.tv_usec;
     
220.     printf("nv12_to_rgb spend %ld (us)\n",t2-t1);
     
221.         for(j=0;j<8;j++){
     
222.             for(i=0;i<8;i++){
     
223.                 k=(j*width+i)*3;
     
224.                 printf("(%02x,%02x,%02x) ",rgb[k],rgb[k+1],rgb[k+2]);
     
225.             }
     
226.             printf("\n");
     
227.             }
     
228.         printf("......\n");
     
229.         for(j=height-8;j<height;j++){
     
230.             for(i=width-8;i<width;i++){
     
231.                 k=(j*width+i)*3;
     
232.                 printf("(%02x,%02x,%02x) ",rgb[k],rgb[k+1],rgb[k+2]);
     
233.             }
     
234.             printf("\n");
     
235.             }
     
236. 
     
237.         
     
238.         free(nv12);
     
239.         free(rgb);
     
240.         return 0;
     
241. }

复制代码

这是一个二维的例子，local_work_item[0]*local_work_item[1]要小于最大的workitem数，即 local_work_item[0]*local_work_item[1]<256，所以取
local_work_item[0]=local_work_item[1]=16

duoduomu

bunchen 发表于 2015-10-9 10:08
二楼待用

不错，不过还可以参考http://malideveloper.arm.com/dow ... .0.0a36a7_Linux.tgz 这个SDK，里面都有，还有demo！

bunchen

duoduomu 发表于 2015-10-9 10:15
不错，不过还可以参考http://malideveloper.arm.com/downloads/tools/oclsdk/Mali_OpenCL_SDK_v1.1.0.0a3 ...

谢谢，之前一直想下载mali的sdk来看，奈何下载页面总是进不去，谢谢分享。

duoduomu

bunchen 发表于 2015-10-9 10:25
谢谢，之前一直想下载mali的sdk来看，奈何下载页面总是进不去，谢谢分享。

哈哈 的确你按常理下载是不行的 ！

fxlsunny

不错正打算做opencl的事，不过想在Ubuntu下做，不知3288是否可行？

ff_20150814

楼主大神你好，请问你的那个NV12转RGB耗时多少呢，GPU的频率是跑的多少M?

tk1user

太厉害了，顶起！

总有刁民想害朕

printOpenCL got 1 platforms :
         platform index=0 profile : FULL_PROFILE
         platform index=0 version : OpenCL 1.1 v1.r6p0-02rel0.0f4218be5cc66c20a4
f31b6cc856ee46
         platform index=0 name : ARM Platform
         platform index=0 vendor : ARM
file /dev/mali0 is not of a compatible version (user 9.0, kernel 8.0)
ERROR: The DDK is not compatible with any of the Mali GPUs on the system.
The DDK was built for 0x750 r0p0 status range [1..1], but none of the GPUs match
ed:
file /dev/mali0 is not of a compatible version (user 9.0, kernel 8.0)
         platform index=0 extensions : out of host memory.
Device cpu:
file /dev/mali0 is not of a compatible version (user 9.0, kernel 8.0)
ERROR: The DDK is not compatible with any of the Mali GPUs on the system.
The DDK was built for 0x750 r0p0 status range [1..1], but none of the GPUs match
ed:
file /dev/mali0 is not of a compatible version (user 9.0, kernel 8.0)
          getDeviceIDs fail !
Device gpu:
file /dev/mali0 is not of a compatible version (user 9.0, kernel 8.0)
ERROR: The DDK is not compatible with any of the Mali GPUs on the system.
The DDK was built for 0x750 r0p0 status range [1..1], but none of the GPUs match
ed:
file /dev/mali0 is not of a compatible version (user 9.0, kernel 8.0)
          getDeviceIDs fail !
Device accelerator:
file /dev/mali0 is not of a compatible version (user 9.0, kernel 8.0)
ERROR: The DDK is not compatible with any of the Mali GPUs on the system.
The DDK was built for 0x750 r0p0 status range [1..1], but none of the GPUs match
ed:
file /dev/mali0 is not of a compatible version (user 9.0, kernel 8.0)
          getDeviceIDs fail !
Device default:
file /dev/mali0 is not of a compatible version (user 9.0, kernel 8.0)
ERROR: The DDK is not compatible with any of the Mali GPUs on the system.
The DDK was built for 0x750 r0p0 status range [1..1], but none of the GPUs match
ed:
file /dev/mali0 is not of a compatible version (user 9.0, kernel 8.0)
          getDeviceIDs fail !
Device all :
file /dev/mali0 is not of a compatible version (user 9.0, kernel 8.0)
ERROR: The DDK is not compatible with any of the Mali GPUs on the system.
The DDK was built for 0x750 r0p0 status range [1..1], but none of the GPUs match
ed:
file /dev/mali0 is not of a compatible version (user 9.0, kernel 8.0)
          getDeviceIDs fail !

Eric.y

duoduomu 发表于 2015-10-9 10:15
不错，不过还可以参考http://malideveloper.arm.com/downloads/tools/oclsdk/Mali_OpenCL_SDK_v1.1.0.0a3 ...

Android-OpenCL-v1.1.zip和ARM官网的SDK一样吗，我没灯泡下载不了