高通平台下安卓opencl小例子

http://blog.csdn.net/wcj0626/article/details/26272019

 

先到高通的qdn下载adreno GPU SDK，里面有OpenCL的例子。
https://developer.qualcomm.com/software/adreno-gpu-sdk

例子在以下目录：
AdrenoSDK\Development\Samples\OpenCL

 

1、从高通官网下载Adreon SDK（需要注册为会员）

2、解压以后，把文件夹Development/Inc/内的CL文件夹拷贝到：Android-ndk-r9d/platforms/android-19/arch-arm/usr/include/（请自行选择拷贝到哪个版本下）

3、连接手机到电脑，在终端运行：adb pull /system/vendor/lib/libOpenCL.so,此时会把手机里的libOpenCL.so拷贝到电脑当前文件夹下，把该文件拷贝到：android-ndk-r9d/platforms/android-19/arch-arm/usr/lib/（请自行选择拷贝到哪个版本下）

4、新建安卓工程，NDK开发。不知道NDK怎么弄，请移步这里

5、Java代码：

 

[java]  view plain  copy

 

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. }  

6、NDK代码：ocl.cpp文件

 

 

[cpp]  view plain  copy

 

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. }  

7、Android.mk代码：

 

 

[plain]  view plain  copy

 

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)  

8、执行结果：

 

搞定！！！