安卓下使用OpenCL进行PowerVR GPU编程
http://blog.****.net/sinat_27685435/article/details/50884175
参考文献
[1]http://blog.****.net/wcj0626/article/details/38061223?c=ac6f159fad7177ce723a5358e13a336c
[2]http://bbs.****.net/topics/390738559
在android 设备上运行opencl ，步骤如下：
第一步、确认
这是最关键一步，首先是检查你的设备是否支持opencl （如果不支持，是无法运行的）；
第二步、找到opencl库
在你的设备中找到支持opencl 的库（这个库通常位于/vendor/lib/libOpenCL.so ）；
第三步、建立动态库
把第二步中找到的库，pull出来，作为作为编译程序的动态库
第四步、引用 opencl 头文件
在android 工程中正常编写opencl 程序，引用opencl 头文件。
备注：android 编译的时候直接链接上面的库，在 Java端编译的时候用System.load()来载入该 opencl 库。
第五步、按照普通NDK程序运行
下面针对NokiaN1平板上的GPU编程进行简介。NokiaN1平板的GPU型号为PowerVRG6430，支持OpenCL1.2。
步骤一、连接手机到电脑，在终端运行：adb pull/system/vendor/lib/libPVROCL.so,此时会把手机里的libOpenCL.so拷贝到电脑当前文件夹下，把该文件拷贝到：android-ndk-r9d/platforms/android-14/arch-x86/usr/lib/（请自行选择拷贝到哪个版本下）
步骤二、下载OpenCL头文件
https://github.com/KhronosGroup/OpenCL-Headers/
步骤三、新建安卓工程，NDK开发。将头文件放到jni/CL文件夹下。
java代码：
1. public class MainActivity extends ActionBarActivity {
2.
3. static
4. {
5. System.loadLibrary("ocl");
6. }
7. public native String testopencl();
8. public native String getPlatformName();
9. public native String getDeviceName();
10. @Override
11. protected void onCreate(Bundle savedInstanceState) {
12. super.onCreate(savedInstanceState);
13. setContentView(R.layout.activity_main);
14.
15. TextView testView2=(TextView)findViewById(R.id.textView2);
16. TextView testView4=(TextView)findViewById(R.id.textView4);
17. TextView testView6=(TextView)findViewById(R.id.textView6);
18. testView6.setText(testopencl());
19. testView2.setText(getPlatformName());
20. testView4.setText(getDeviceName());
21.
22. }
23.
24.
25.
26.}
NDK代码：
1. #include <jni.h>
2. #include <CL/cl.h>
3. #include<malloc.h>
4. #include<stdio.h>
5. #include<stdlib.h>
6. #include"com_example_ocl_MainActivity.h"
7. #define LEN(arr) sizeof(arr) / sizeof(arr[0])
8. #define N 1024
9. #define NUM_THREAD 128
10.
11. cl_uint num_device;
12.cl_uint num_platform;
13. cl_platform_id *platform;
14.cl_device_id *devices;
15. cl_int err;
16.cl_context context;
17. cl_command_queue cmdQueue;
18. cl_mem buffer,sum_buffer;
19. cl_program program ;
20. cl_kernel kernel;
21. const char* src[] = {
22. " __kernel void redution( \n"
23. " __global int *data, \n"
24. " __global int *output, \n"
25. " __local int *data_local \n"
26. " ) \n"
27. " { \n"
28. " int gid=get_group_id(0); \n"
29. " int tid=get_global_id(0); \n"
30. " int size=get_local_size(0); \n"
31. " int id=get_local_id(0); \n"
32. " data_local[id]=data[tid]; \n"
33. " barrier(CLK_LOCAL_MEM_FENCE); \n"
34. " for(int i=size/2;i>0;i>>=1){ \n"
35. " if(id<i){ \n"
36. " data_local[id]+=data_local[id+i]; \n"
37. " } \n"
38. " barrier(CLK_LOCAL_MEM_FENCE); \n"
39. " } \n"
40. " if(id==0){ \n"
41. " output[gid]=data_local[0]; \n"
42. " } \n"
43. " } \n"
44.
45. };
46.int num_block;
47.
48.
49. int test()
50.{
51. int* in,*out;
52. num_block=N/NUM_THREAD;
53. in=(int*)malloc(sizeof(int)*N);
54. out=(int*)malloc(sizeof(int)*num_block);
55. for(int i=0;i<N;i++){
56. in[i]=1;
57. }
58. Init_OpenCL();
59. Context_cmd();
60. Create_Buffer(in);
61. Create_program();
62. Set_arg();
63. Execution();
64. CopyOutResult(out);
65. int sum=0;
66. for(int i=0;i<num_block;i++){
67. sum+=out[i];
68. }
69. return sum;
70.}
71.
72. JNIEXPORT jstring JNICALL Java_com_example_ocl_MainActivity_testopencl (JNIEnv * env, jobject thisobject)
73. {
74. char result[10];
75. sprintf(result,"%d\n",test());
76. return env->NewStringUTF(result);
77. }
78. JNIEXPORT jstring JNICALL Java_com_example_ocl_MainActivity_getPlatformName(JNIEnv *env , jobject thisobject)
79. {
80. char buffer[1024];
81. clGetPlatformInfo(platform[0],CL_PLATFORM_NAME,sizeof(buffer),buffer,NULL);
82. return env->NewStringUTF(buffer);
83. }
84. JNIEXPORT jstring JNICALL Java_com_example_ocl_MainActivity_getDeviceName(JNIEnv *env , jobject thisobject)
85. {
86.
87. char buffer[1024];
88. clGetDeviceInfo(devices[0],CL_DEVICE_NAME,sizeof(buffer),buffer,NULL);
89. return env->NewStringUTF(buffer);
90. }
91.
92. void Init_OpenCL()
93. {
94. size_t nameLen1;
95. char platformName[1024];
96.
97. err = clGetPlatformIDs(0, 0, &num_platform);
98. platform=(cl_platform_id*)malloc(sizeof(cl_platform_id)*num_platform);
99. err = clGetPlatformIDs(num_platform, platform, NULL);
100.
101. err=clGetDeviceIDs(platform[0],CL_DEVICE_TYPE_GPU,0,NULL,&num_device);
102. devices=(cl_device_id*)malloc(sizeof(cl_device_id)*num_device);
103. err=clGetDeviceIDs(platform[0],CL_DEVICE_TYPE_GPU,num_device,devices,NULL);
104.
105. }
106.
107. void Context_cmd()
108. {
109. context=clCreateContext(NULL,num_device,devices,NULL,NULL,&err);
110. cmdQueue=clCreateCommandQueue(context,devices[0],0,&err);
111. }
112.
113. void Create_Buffer(int *data)
114. {
115.
116. buffer=clCreateBuffer(context,CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,sizeof(int)*N,data,&err);
117. sum_buffer=clCreateBuffer(context,CL_MEM_WRITE_ONLY,sizeof(int)*num_block,0,&err);
118. }
119.
120. void Create_program()
121. {
122. program=clCreateProgramWithSource(context, LEN(src), src, NULL, NULL);
123. err=clBuildProgram(program,num_device,devices,NULL,NULL,NULL);
124. kernel = clCreateKernel(program, "redution", NULL);
125. }
126.
127. void Set_arg()
128.{
129. err=clSetKernelArg(kernel,0,sizeof(cl_mem),&buffer);
130. err=clSetKernelArg(kernel,1,sizeof(cl_mem),&sum_buffer);
131. err=clSetKernelArg(kernel,2,sizeof(int)*NUM_THREAD,NULL);
132.}
133.
134.void Execution()
135. {
136. const size_t globalWorkSize[1]={N};
137. const size_t localWorkSize[1]={NUM_THREAD};
138. err=clEnqueueNDRangeKernel(cmdQueue,kernel,1,NULL,globalWorkSize,localWorkSize,0,NULL,NULL);
139. clFinish(cmdQueue);
140.}
141.
142.void CopyOutResult(int*out)
143. {
144. err=clEnqueueReadBuffer(cmdQueue,sum_buffer,CL_TRUE,0,sizeof(int)*num_block,out,0,NULL,NULL);
145. }
Android.mk
1. LOCAL_PATH := $(call my-dir)
2.
3. include $(CLEAR_VARS)
4. LOCAL_MODULE := ocl
5. LOCAL_SRC_FILES := ocl.cpp
6. LOCAL_LDFLAGS += -llog -lOpenCL
7. include $(BUILD_SHARED_LIBRARY)
Application.mk
1. APP_CPPFLAGS := -frtti -fexceptions
2. APP_STL := stlport_static
3. APP_ABI := x86
4. APP_PLATFORM := android-14
执行结果
工程文件下载地址：
http://download.****.net/detail/sinat_27685435/9460498