高通平台下安卓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代码：

public class MainActivity extends ActionBarActivity {

	static
    {
        System.loadLibrary("ocl");
    }
	public native String  testopencl();
	public native String getPlatformName();
	public native String getDeviceName();
    @Override
    protected void onCreate(Bundle savedInstanceState) {
        super.onCreate(savedInstanceState);
        setContentView(R.layout.activity_main);

      TextView testView2=(TextView)findViewById(R.id.textView2);
      TextView testView4=(TextView)findViewById(R.id.textView4);
      TextView testView6=(TextView)findViewById(R.id.textView6);
      testView6.setText(testopencl());
      testView2.setText(getPlatformName());
       testView4.setText(getDeviceName());
       
    }



}

#include <jni.h>
#include <CL/cl.h>
#include<malloc.h>
#include<stdio.h>
#include<stdlib.h>
#include"com_example_ocl_MainActivity.h"
#define LEN(arr) sizeof(arr) / sizeof(arr[0])
#define N 1024
#define NUM_THREAD 128

cl_uint num_device;
cl_uint num_platform;
cl_platform_id *platform;
cl_device_id *devices;
cl_int err;
cl_context context;
 cl_command_queue cmdQueue;
 cl_mem buffer,sum_buffer;
 cl_program program ;
 cl_kernel kernel;
 const char* src[] = {
		 "  __kernel void redution(  \n"
		 "	__global int *data,     \n"
		 "	__global int *output,   \n"
		 "	__local int *data_local   \n"
		 "	)  \n"
		" {   \n"
		 "	int gid=get_group_id(0);   \n"
		 "	int tid=get_global_id(0);    \n"
		 "	int size=get_local_size(0);   \n"
		 "	int id=get_local_id(0);     \n"
		 "	data_local[id]=data[tid];   \n"
		 "	barrier(CLK_LOCAL_MEM_FENCE);   \n"
		 "	for(int i=size/2;i>0;i>>=1){    \n"
		 "		if(id<i){   \n"
		 "			data_local[id]+=data_local[id+i];   \n"
		 "		}   \n"
		 "		barrier(CLK_LOCAL_MEM_FENCE);   \n"
		 "	}    \n"
		 "	if(id==0){    \n"
		 "		output[gid]=data_local[0];   \n"
		 "	}    \n"
		" }   \n"

 };
int num_block;


int  test()
{
	int* in,*out;
	 num_block=N/NUM_THREAD;
	in=(int*)malloc(sizeof(int)*N);
	out=(int*)malloc(sizeof(int)*num_block);
	for(int i=0;i<N;i++){
		in[i]=1;
	}
	Init_OpenCL();
	Context_cmd();
	Create_Buffer(in);
	Create_program();
	Set_arg();
	Execution();
	CopyOutResult(out);
	int sum=0;
	for(int i=0;i<num_block;i++){
		sum+=out[i];
	}
	return sum;
}

 JNIEXPORT jstring JNICALL Java_com_example_ocl_MainActivity_testopencl (JNIEnv * env, jobject thisobject)
{
    char result[10];
    sprintf(result,"%d\n",test());
	 return env->NewStringUTF(result);
}
 JNIEXPORT jstring JNICALL Java_com_example_ocl_MainActivity_getPlatformName(JNIEnv *env , jobject thisobject)
 {
	 	 char buffer[1024];
	 	 clGetPlatformInfo(platform[0],CL_PLATFORM_NAME,sizeof(buffer),buffer,NULL);
	 	return env->NewStringUTF(buffer);
 }
 JNIEXPORT jstring JNICALL Java_com_example_ocl_MainActivity_getDeviceName(JNIEnv *env , jobject thisobject)
 {

	 char buffer[1024];
	 clGetDeviceInfo(devices[0],CL_DEVICE_NAME,sizeof(buffer),buffer,NULL);
	return env->NewStringUTF(buffer);
 }

 void Init_OpenCL()
 {
	 size_t nameLen1;
	 char platformName[1024];

	 err = clGetPlatformIDs(0, 0, &num_platform);
	 platform=(cl_platform_id*)malloc(sizeof(cl_platform_id)*num_platform);
	 err = clGetPlatformIDs(num_platform, platform, NULL);

	 err=clGetDeviceIDs(platform[0],CL_DEVICE_TYPE_GPU,0,NULL,&num_device);
	 devices=(cl_device_id*)malloc(sizeof(cl_device_id)*num_device);
	 err=clGetDeviceIDs(platform[0],CL_DEVICE_TYPE_GPU,num_device,devices,NULL);

 }

 void Context_cmd()
 {
	 context=clCreateContext(NULL,num_device,devices,NULL,NULL,&err);
	 cmdQueue=clCreateCommandQueue(context,devices[0],0,&err);
 }

 void Create_Buffer(int *data)
 {

	 buffer=clCreateBuffer(context,CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,sizeof(int)*N,data,&err);
	 sum_buffer=clCreateBuffer(context,CL_MEM_WRITE_ONLY,sizeof(int)*num_block,0,&err);
 }

 void Create_program()
 {
	 program=clCreateProgramWithSource(context, LEN(src), src, NULL, NULL);
	 err=clBuildProgram(program,num_device,devices,NULL,NULL,NULL);
	 kernel = clCreateKernel(program, "redution", NULL);
 }

void Set_arg()
{
	err=clSetKernelArg(kernel,0,sizeof(cl_mem),&buffer);
	err=clSetKernelArg(kernel,1,sizeof(cl_mem),&sum_buffer);
	err=clSetKernelArg(kernel,2,sizeof(int)*NUM_THREAD,NULL);
}

void Execution()
{
	const size_t globalWorkSize[1]={N};
		const size_t localWorkSize[1]={NUM_THREAD};
	  err=clEnqueueNDRangeKernel(cmdQueue,kernel,1,NULL,globalWorkSize,localWorkSize,0,NULL,NULL);
	  clFinish(cmdQueue);
}

void CopyOutResult(int*out)
{
	err=clEnqueueReadBuffer(cmdQueue,sum_buffer,CL_TRUE,0,sizeof(int)*num_block,out,0,NULL,NULL);
}

LOCAL_PATH := $(call my-dir)

include $(CLEAR_VARS)
LOCAL_MODULE    := ocl
LOCAL_SRC_FILES := ocl.cpp
LOCAL_LDFLAGS += -llog  -lOpenCL
include $(BUILD_SHARED_LIBRARY)

搞定！！！