OpenCL

OpenCL和一組用於定義並控制平臺的API組成。OpenCL提供了基於任務分割和資料分割的平行計算機制。OpenCL類似於另外兩個開放的工業標準OpenGL和OpenAL，這兩個標準分別用於三維圖形和計算機音訊方面。OpenCL擴充了GPU圖形生成之外的能力。OpenCL由非盈利性技術組織Khronos Group掌管。歷史.OpenCL最初由蘋果公司開發，擁有其商標權，並在與AMD，IBM，Intel和NVIDIA技術團隊的合作之下初步完善。隨後，蘋果將這一草案提交至Khronos Group。2008年6月16日，Khronos的通用計算工作小組成立。5個月後的2008年11月18日，該工作組完成了OpenCL 1.0規範的技術細節。該技術規範在由Khronos成員進行審查之後，於2008年12月8日公開發表。2010年6月14日，OpenCL 1.1釋出。範例.快速傅立葉變換.一個快速傅立葉變換的式子：// create a compute context with GPU devicecontext = clCreateContextFromType(NULL, CL_DEVICE_TYPE_GPU, NULL, NULL, NULL);// create a command queuequeue = clCreateCommandQueue(context, NULL, 0, NULL);// allocate the buffer memory objectsmemobjs[0] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float)*2*num_entries, srcA, NULL);memobjs[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float)*2*num_entries, NULL, NULL);// create the compute programprogram = clCreateProgramWithSource(context, 1, &amp;fft1D_1024_kernel_src, NULL, NULL);// build the compute program executableclBuildProgram(program, 0, NULL, NULL, NULL, NULL);// create the compute kernelkernel = clCreateKernel(program, "fft1D_1024", NULL);// set the args valuesclSetKernelArg(kernel, 0, sizeof(cl_mem),(void *)&amp;memobjs[0]);clSetKernelArg(kernel, 1, sizeof(cl_mem),(void *)&amp;memobjs[1]);clSetKernelArg(kernel, 2, sizeof(float)*(local_work_size[0]+1)*16, NULL);clSetKernelArg(kernel, 3, sizeof(float)*(local_work_size[0]+1)*16, NULL);// create N-D range object with work-item dimensions and execute kernelglobal_work_size[0] = num_entries;local_work_size[0] = 64;clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);真正的運算：// This kernel computes FFT of length 1024. The 1024 length FFT is decomposed into// calls to a radix 16 function, another radix 16 function and then a radix 4 function__kernel void fft1D_1024(__global float2 *in, __global float2 *out,__local float *sMemx, __local float *sMemy){int tid = get_local_id(0);int blockIdx = get_group_id(0) * 1024 + tid;float2 data[16];// starting index of data to/from global memoryin = in + blockIdx; out = out + blockIdx;globalLoads(data, in, 64); // coalesced global readsfftRadix16Pass(data); // in-place radix-16 passtwiddleFactorMul(data, tid, 1024, 0);// local shuffle using local memorylocalShuffle(data, sMemx, sMemy, tid, (((tid &amp; 15)* 65) +(tid » 4)));fftRadix16Pass(data); // in-place radix-16 passtwiddleFactorMul(data, tid, 64, 4); // twiddle factor multiplicationlocalShuffle(data, sMemx, sMemy, tid, (((tid » 4)* 64) +(tid &amp; 15)));// four radix-4 function callsfftRadix4Pass(data); // radix-4 function number 1fftRadix4Pass(data + 4); // radix-4 function number 2fftRadix4Pass(data + 8); // radix-4 function number 3fftRadix4Pass(data + 12); // radix-4 function number 4// coalesced global writesglobalStores(data, out, 64);Apple的網站上可以發現傅立葉變換的例子平行合併排序法.使用 Python 3.x 搭配 PyOpenCL 與 NumPyimport ioimport randomimport numpy as npimport pyopencl as cldef dump_step(data, chunk_size):"""顯示排序過程"""msg = io.StringIO(")div = io.StringIO(")for idx, item in enumerate(data):if idx % chunk_size == 0:if idx &gt; 0:msg.write(' ||')div.write(' ')div.write(' --')else:msg.write(' ')div.write('------')msg.write(' {:2d}'.format(item))out = msg.getvalue()if chunk_size == 1: print(' ' + '-' * (len(out) - 1))print(out)print(div.getvalue())msg.close()div.close()def cl_merge_sort_sbs(data_in):"""平行合併排序"""# OpenCL kernel 函式程式碼CL_CODE =kernel void merge(int chunk_size, int size, global long* data, global long* buff) {// 取得分組編號const int gid = get_global_id(0);// 根據分組編號計算責任範圍const int offset = gid * chunk_size;const int real_size = min(offset + chunk_size, size) - offset;global long* data_part = data + offset;global long* buff_part = buff + offset;// 設定合併前的初始狀態int r_beg = chunk_size » 1;int b_ptr = 0;int l_ptr = 0;int r_ptr = r_beg;// 進行合併while (b_ptr = real_size) {// 若右側沒有資料，取左側資料堆入緩衝區buff_part[b_ptr] = data_part[l_ptr++];} else if (l_ptr == r_beg) {// 若左側沒有資料，取右側資料堆入緩衝區buff_part[b_ptr] = data_part[r_ptr++];} else {// 若兩側都有資料，取較小資料堆入緩衝區if (data_part[l_ptr] &lt; data_part[r_ptr]) {buff_part[b_ptr] = data_part[l_ptr++];} else {buff_part[b_ptr] = data_part[r_ptr++];b_ptr++;# 配置計算資源，編譯 OpenCL 程式ctx = cl.Context(dev_type=cl.device_type.GPU)prg = cl.Program(ctx, CL_CODE).build()queue = cl.CommandQueue(ctx)mf = cl.mem_flags# 資料轉換成 numpy 形式以利轉換為 OpenCL Bufferdata_np = np.int64(data_in)buff_np = np.empty_like(data_np)# 建立緩衝區，並且複製數值到緩衝區data = cl.Buffer(ctx, mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=data_np)buff = cl.Buffer(ctx, mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=buff_np)# 設定合併前初始狀態data_len = np.int32(len(data_np))chunk_size = np.int32(1)dump_step(data_np, chunk_size)while chunk_size &lt; data_len:# 更新分組大小，每一回合變兩倍chunk_size «= 1# 換算平行作業組數group_size = ((data_len - 1) // chunk_size) + 1# 進行分組合併作業prg.merge(queue, (group_size,), (1,), chunk_size, data_len, data, buff)# 將合併結果作為下一回合的原始資料temp = datadata = buffbuff = temp# 顯示此回合狀態cl.enqueue_copy(queue, data_np, data)dump_step(data_np, chunk_size)queue.finish()data.release()buff.release()def main():n = random.randint(5, 16)data = []for i in range(n):data.append(random.randint(1, 99))cl_merge_sort_sbs(data)if __name__ == '__main__':main()執行結果：85 || 41 || 64 || 40 || 90 || 29 || 38 || 41 || 64 || 17 || 20 || 41 || 16 || 65 || 8341 85 || 40 64 || 29 90 || 38 41 || 17 64 || 20 41 || 16 65 || 8340 41 64 85 || 29 38 41 90 || 17 20 41 64 || 16 65 8329 38 40 41 41 64 85 90 || 16 17 20 41 64 65 8316 17 20 29 38 40 41 41 41 64 64 65 83 85 90