OpenCL內存模型
OpenCL的內存模型定義了各種各樣內存類型，各種內存模型之間有層級關系。各種內存之間的數據傳輸必須是顯式進行的，比如從host memory到device memory，從global memory到local memory等等。
WorkGroup被映射到硬件的CU上執行（在AMD 5xxx系列顯卡上，CU就是simd，一個simd中有16個pe），OpenCL並不提供各個workgroup之間的一致性，如果我們需要在各個workgroup之間共享數據或者通信之類的，要自己通過軟件實現。
Kernel函數的寫法
每個線程（workitem）都有一個kenerl函數的實例。下面我們看下kernel的寫法：
1: __kernel void vecadd(__global const float* A, __global const float* B, __global float* C)
2: {
3: int id = get_global_id(0);
4: C[id] = A[id] + B[id];
5: }
每個Kernel函數都必須以__kernel開始，而且必須返回void。每個輸入參數都必須聲明使用的內存類型。通過一些API，比如get_global_id之類的得到線程id。
內存對象地址空間標識符有以下幾種：
__global – memory allocated from global address space
__constant – a special type of read-only memory
__local – memory shared by a work-group
__private – private per work-item memory
__read_only/__write_only – used for images
Kernel函數參數如果是內存對象，那么一定是__global,__local或者constant。
運行Kernel
首先要設置線程索引空間的維數以及workgroup大小等。
我們通過函數clEnqueueNDRangeKerne把Kernel放在一個隊列里，但不保證它馬上執行，OpenCL driver會管理隊列，調度Kernel的執行。注意：每個線程執行的代碼都是相同的，但是它們執行數據卻是不同的。
該函數把要執行的Kernel函數放在指定的命令隊列中，globald大小（線程索引空間）必須指定，local大小（work group）可以指定，也可以為空。如果為空，則系統會自動根據硬件選擇合適的大小。event_wait_list用來選定一些events，只有這些events執行完后，該kernel才可能被執行，也就是通過事件機制來實現不同kernel函數之間的同步。
當Kernel函數執行完畢后，我們要把數據從device memory中拷貝到host memory中去。
釋放資源：
大多數的OpenCL資源都是指針，不使用的時候需要釋放掉。當然，程序關閉的時候這些對象也會被自動釋放掉。
釋放資源的函數是：clRelase{Resource} ，比如: clReleaseProgram(), clReleaseMemObject()等。
錯誤捕捉：
如果OpenCL函數執行失敗，會返回一個錯誤碼，一般是個負值，返回0則表示執行成功。我們可以根據該錯誤碼知道什么地方出錯了，需要修改。錯誤碼在cl.h中定義，下面是幾個錯誤碼的例子.
CL_DEVICE_NOT_FOUND -1
CL_DEVICE_NOT_AVAILABLE -2
CL_COMPILER_NOT_AVAILABLE -3
CL_MEM_OBJECT_ALLOCATION_FAILURE -4
…
下面是一個OpenCL機制的示意圖
程序模型
數據並行：work item和內存對象元素之間是一一映射關系；workgroup可以顯示指定，也可以隱式指定。
任務並行：kernel的執行獨立於線程索引空間；用其他方法表示並行，比如把不同的任務放入隊列，用設備指定的特殊的向量類型等等。
同步：workgroup內work item之間的同步；命令隊列中不同命令之間的同步。
完整代碼如下：
1: #include "stdafx.h"
2: #include <CL/cl.h>
3: #include <stdio.h>
4: #include <stdlib.h>
5: #include <time.h>
6: #include <iostream>
7: #include <fstream>
8:
9: using namespace std;
10: #define NWITEMS 262144
11:
12: #pragma comment (lib,"OpenCL.lib")
13:
14: //把文本文件讀入一個string中
15: int convertToString(const char *filename, std::string& s)
16: {
17: size_t size;
18: char* str;
19:
20: std::fstream f(filename, (std::fstream::in | std::fstream::binary));
21:
22: if(f.is_open())
23: {
24: size_t fileSize;
25: f.seekg(0, std::fstream::end);
26: size = fileSize = (size_t)f.tellg();
27: f.seekg(0, std::fstream::beg);
28:
29: str = new char[size+1];
30: if(!str)
31: {
32: f.close();
33: return NULL;
34: }
35:
36: f.read(str, fileSize);
37: f.close();
38: str[size] = '\0';
39:
40: s = str;
41: delete[] str;
42: return 0;
43: }
44: printf("Error: Failed to open file %s\n", filename);
45: return 1;
46: }
47:
48: int main(int argc, char* argv[])
49: {
50: //在host內存中創建三個緩沖區
51: float *buf1 = 0;
52: float *buf2 = 0;
53: float *buf = 0;
54:
55: buf1 =(float *)malloc(NWITEMS * sizeof(float));
56: buf2 =(float *)malloc(NWITEMS * sizeof(float));
57: buf =(float *)malloc(NWITEMS * sizeof(float));
58:
59: //初始化buf1和buf2的內容
60: int i;
61: srand( (unsigned)time( NULL ) );
62: for(i = 0; i < NWITEMS; i++)
63: buf1[i] = rand()%65535;
64:
65: srand( (unsigned)time( NULL ) +1000);
66: for(i = 0; i < NWITEMS; i++)
67: buf2[i] = rand()%65535;
68:
69: for(i = 0; i < NWITEMS; i++)
70: buf[i] = buf1[i] + buf2[i];
71:
72: cl_uint status;
73: cl_platform_id platform;
74:
75: //創建平台對象
76: status = clGetPlatformIDs( 1, &platform, NULL );
77:
78: cl_device_id device;
79:
80: //創建GPU設備
81: clGetDeviceIDs( platform, CL_DEVICE_TYPE_GPU,
82: 1,
83: &device,
84: NULL);
85: //創建context
86: cl_context context = clCreateContext( NULL,
87: 1,
88: &device,
89: NULL, NULL, NULL);
90: //創建命令隊列
91: cl_command_queue queue = clCreateCommandQueue( context,
92: device,
93: CL_QUEUE_PROFILING_ENABLE, NULL );
94: //創建三個OpenCL內存對象，並把buf1的內容通過隱式拷貝的方式
95: //拷貝到clbuf1,buf2的內容通過顯示拷貝的方式拷貝到clbuf2
96: cl_mem clbuf1 = clCreateBuffer(context,
97: CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
98: NWITEMS*sizeof(cl_float),buf1,
99: NULL );
100:
101: cl_mem clbuf2 = clCreateBuffer(context,
102: CL_MEM_READ_ONLY ,
103: NWITEMS*sizeof(cl_float),NULL,
104: NULL );
105:
106: status = clEnqueueWriteBuffer(queue, clbuf2, 1,
107: 0, NWITEMS*sizeof(cl_float), buf2, 0, 0, 0);
108:
109: cl_mem buffer = clCreateBuffer( context,
110: CL_MEM_WRITE_ONLY,
111: NWITEMS * sizeof(cl_float),
112: NULL, NULL );
113:
114: const char * filename = "add.cl";
115: std::string sourceStr;
116: status = convertToString(filename, sourceStr);
117: const char * source = sourceStr.c_str();
118: size_t sourceSize[] = { strlen(source) };
119:
120: //創建程序對象
121: cl_program program = clCreateProgramWithSource(
122: context,
123: 1,
124: &source,
125: sourceSize,
126: NULL);
127: //編譯程序對象
128: status = clBuildProgram( program, 1, &device, NULL, NULL, NULL );
129: if(status != 0)
130: {
131: printf("clBuild failed:%d\n", status);
132: char tbuf[0x10000];
133: clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0x10000, tbuf, NULL);
134: printf("\n%s\n", tbuf);
135: return -1;
136: }
137:
138: //創建Kernel對象
139: cl_kernel kernel = clCreateKernel( program, "vecadd", NULL );
140: //設置Kernel參數
141: cl_int clnum = NWITEMS;
142: clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*) &clbuf1);
143: clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*) &clbuf2);
144: clSetKernelArg(kernel, 2, sizeof(cl_mem), (void*) &buffer);
145:
146: //執行kernel
147: cl_event ev;
148: size_t global_work_size = NWITEMS;
149: clEnqueueNDRangeKernel( queue,
150: kernel,
151: 1,
152: NULL,
153: &global_work_size,
154: NULL, 0, NULL, &ev);
155: clFinish( queue );
156:
157: //數據拷回host內存
158: cl_float *ptr;
159: ptr = (cl_float *) clEnqueueMapBuffer( queue,
160: buffer,
161: CL_TRUE,
162: CL_MAP_READ,
163: 0,
164: NWITEMS * sizeof(cl_float),
165: 0, NULL, NULL, NULL );
166: //結果驗證，和cpu計算的結果比較
167: if(!memcmp(buf, ptr, NWITEMS))
168: printf("Verify passed\n");
169: else printf("verify failed");
170:
171: if(buf)
172: free(buf);
173: if(buf1)
174: free(buf1);
175: if(buf2)
176: free(buf2);
177:
178: //刪除OpenCL資源對象
179: clReleaseMemObject(clbuf1);
180: clReleaseMemObject(clbuf2);
181: clReleaseMemObject(buffer);
182: clReleaseProgram(program);
183: clReleaseCommandQueue(queue);
184: clReleaseContext(context);
185: return 0;
186: }
187:
也可以在http://code.google.com/p/imagefilter-opencl/downloads/detail?name=amdunicourseCode1.zip&can=2&q=#makechanges上下載完整版本。