安卓下使用OpenCL进行PowerVR GPU编程

参考文献

[1]http://blog.csdn.net/wcj0626/article/details/38061223?c=ac6f159fad7177ce723a5358e13a336c

[2]http://bbs.csdn.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.csdn.net/detail/sinat_27685435/9460498