OpenCL

對於異質計算系統(如GPU或GPU)程式設計開放標準

OpenCLOpen Computing Language,开放计算语言)是一个为异构平台编写程序的框架,此异构平台可由CPUGPUDSPFPGA或其他类型的处理器與硬體加速器所组成。OpenCL由一门用于编写kernels(在OpenCL设备上运行的函数)的语言(基于C99)和一组用于定义并控制平台的API组成。OpenCL提供了基于任务分割和数据分割的并行计算机制。

OpenCL
原作者苹果公司
開發者科纳斯组织
首次发布2009年8月28日,​14年前​(2009-08-28
当前版本3.0.12 (2022年9月15日,​19個月前​(2022-09-15)
操作系统跨平台
类型API
许可协议免版税
网站www.khronos.org/opencl
www.khronos.org/webcl

OpenCL类似于另外两个开放的工业标准OpenGLOpenAL,这两个标准分别用于三维图形和计算机音频方面。OpenCL擴充了GPU圖形生成之外的能力。OpenCL由非盈利性技术组织Khronos Group掌管。

历史编辑

OpenCL最初由苹果公司开发,拥有其商标权,并在与AMDIBMIntelNVIDIA技术团队的合作之下初步完善。随后,苹果将这一草案提交至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]

範例编辑

快速傅立葉變換编辑

一個快速傅立葉變換的式子:[6]

  // 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 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 > 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 --------------------------------------------------------------------------------------

參考文獻编辑

  1. ^ Khronos OpenCL Registry. Khronos Group. 2022-09-15 [2022-09-15]. (原始内容存档于2019-09-14) (英语). 
  2. ^ Khronos Launches Heterogeneous Computing Initiative (新闻稿). Khronos Group. 2008-06-16 [2008-06-18]. (原始内容存档于2008-06-20). 
  3. ^ OpenCL gets touted in Texas. MacWorld. 2008-11-20 [2009-06-12]. (原始内容存档于2009-02-18). 
  4. ^ The Khronos Group Releases OpenCL 1.0 Specification (新闻稿). Khronos Group. 2008-12-08 [2009-06-12]. (原始内容存档于2010-07-13). 
  5. ^ 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). 
  6. ^ OpenCL (PDF). SIGGRAPH2008. 2008-08-14 [2008-08-14]. (原始内容 (PDF)存档于2012-03-19). 
  7. ^ 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). 
  8. ^ .OpenCL on FFT. Apple. 16 Nov 2009 [2009-12-07]. (原始内容存档于2009-11-30). 

外部連結编辑

参见编辑