OpenCL
 | |
| 原作者 | 蘋果公司 |
|---|---|
| 開發者 | 科納斯組織 |
| 首次發布 | 2009年8月28日 |
| 當前版本 | 3.0.12 (2022年9月15日) |
| 操作系統 | 跨平台 |
| 類型 | API |
| 許可協議 | 免版稅 |
| 網站 | www www |
OpenCL（Open Computing Language，開放計算語言）是一個為異構平台編寫程序的框架，此異構平台可由CPU、GPU、DSP、FPGA或其他類型的處理器與硬體加速器所組成。OpenCL由一門用於編寫kernels（在OpenCL設備上運行的函數）的語言（基於C99）和一組用於定義並控制平台的API組成。OpenCL提供了基於任務分割和數據分割的並行計算機制。
OpenCL類似於另外兩個開放的工業標準OpenGL和OpenAL，這兩個標準分別用於三維圖形和計算機音頻方面。OpenCL擴充了GPU圖形生成之外的能力。OpenCL由非盈利性技術組織Khronos Group掌管。
歷史
OpenCL最初由蘋果公司開發，擁有其商標權，並在與AMD，IBM，Intel和NVIDIA技術團隊的合作之下初步完善。隨後，蘋果將這一草案提交至Khronos Group。
2008年6月16日，Khronos的通用計算工作小組成立[2]。5個月後的2008年11月18日，該工作組完成了OpenCL 1.0規範的技術細節[3]。該技術規範在由Khronos成員進行審查之後，於2008年12月8日公開發表[4]。2010年6月14日，OpenCL 1.1發布[5]。
範例
快速傅立葉變換
// create a compute context with GPU device
context = clCreateContextFromType(NULL, CL_DEVICE_TYPE_GPU, NULL, NULL, NULL);
// create a command queue
queue = clCreateCommandQueue(context, NULL, 0, NULL);
// allocate the buffer memory objects
memobjs[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 program
program = clCreateProgramWithSource(context, 1, &fft1D_1024_kernel_src, NULL, NULL);
// build the compute program executable
clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
// create the compute kernel
kernel = clCreateKernel(program, "fft1D_1024", NULL);
// set the args values
clSetKernelArg(kernel, 0, sizeof(cl_mem),(void *)&memobjs[0]);
clSetKernelArg(kernel, 1, sizeof(cl_mem),(void *)&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 kernel
global_work_size[0] = num_entries;
local_work_size[0] = 64;
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
真正的運算：（基於Fitting FFT onto the G80 Architecture（頁面存檔備份，存於網際網路檔案館））[7]
// 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 memory
in = in + blockIdx; out = out + blockIdx;
globalLoads(data, in, 64); // coalesced global reads
fftRadix16Pass(data); // in-place radix-16 pass
twiddleFactorMul(data, tid, 1024, 0);
// local shuffle using local memory
localShuffle(data, sMemx, sMemy, tid, (((tid & 15)* 65) +(tid >> 4)));
fftRadix16Pass(data); // in-place radix-16 pass
twiddleFactorMul(data, tid, 64, 4); // twiddle factor multiplication
localShuffle(data, sMemx, sMemy, tid, (((tid >> 4)* 64) +(tid & 15)));
// four radix-4 function calls
fftRadix4Pass(data); // radix-4 function number 1
fftRadix4Pass(data + 4); // radix-4 function number 2
fftRadix4Pass(data + 8); // radix-4 function number 3
fftRadix4Pass(data + 12); // radix-4 function number 4
// coalesced global writes
globalStores(data, out, 64);
}
Apple的網站上可以發現傅立葉變換的例子[8]
平行合併排序法
使用 Python 3.x 搭配 PyOpenCL 與 NumPy
import io
import random
import numpy as np
import pyopencl as cl
def dump_step(data, chunk_size):
"""顯示排序過程"""
msg = io.StringIO('')
div = io.StringIO('')
for idx, item in enumerate(data):
if idx % chunk_size == 0:
if idx > 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) {
if (r_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] < 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 Buffer
data_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 < 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 = data
data = buff
buff = 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 || 83
-- -- -- -- -- -- -- -- -- -- -- -- -- -- --
41 85 || 40 64 || 29 90 || 38 41 || 17 64 || 20 41 || 16 65 || 83
-------- -------- -------- -------- -------- -------- -------- --
40 41 64 85 || 29 38 41 90 || 17 20 41 64 || 16 65 83
-------------------- -------------------- -------------------- --------------
29 38 40 41 41 64 85 90 || 16 17 20 41 64 65 83
-------------------------------------------- --------------------------------------
16 17 20 29 38 40 41 41 41 64 64 65 83 85 90
--------------------------------------------------------------------------------------
參考文獻
- ^ Khronos OpenCL Registry. Khronos Group. 2022-09-15 [2022-09-15]. （原始內容存檔於2019-09-14） （英語）.
- ^ Khronos Launches Heterogeneous Computing Initiative (新聞稿). Khronos Group. 2008-06-16 [2008-06-18]. （原始內容存檔於2008-06-20）.
- ^ OpenCL gets touted in Texas. MacWorld. 2008-11-20 [2009-06-12]. （原始內容存檔於2009-02-18）.
- ^ The Khronos Group Releases OpenCL 1.0 Specification (新聞稿). Khronos Group. 2008-12-08 [2009-06-12]. （原始內容存檔於2010-07-13）.
- ^ Khronos Drives Momentum of Parallel Computing Standard with Release of OpenCL 1.1 Specification (新聞稿). Khronos Group. 2010-06-14 [2010-10-13]. （原始內容存檔於2010-09-23）.
- ^ OpenCL (PDF). SIGGRAPH2008. 2008-08-14 [2008-08-14]. （原始內容 (PDF)存檔於2012-03-19）.
- ^ Fitting FFT onto G80 Architecture (PDF). Vasily Volkov and Brian Kazian, UC Berkeley CS258 project report. May 2008 [2008-11-14]. （原始內容存檔 (PDF)於2012-03-19）.
- ^ . OpenCL on FFT. Apple. 16 Nov 2009 [2009-12-07]. （原始內容存檔於2009-11-30）.
外部連結
- 支援OpenCL的產品（頁面存檔備份，存於網際網路檔案館）
- 開源GPU社區（簡體中文）
參見
| ||||||||||||