macOS的OpenCL高性能計算

隨著深度學習、區塊鏈的發展，人類對計算量的需求越來越高，在傳統的計算模式下，壓榨GPU的計算能力一直是重點。
NV系列的顯卡在這方面走的比較快，CUDA框架已經普及到了高性能計算的各個方面，比如Google的TensorFlow深度學習框架，預設內置了支持CUDA的GPU計算。
AMD(ATI)及其它顯卡在這方面似乎一直不夠給力，在CUDA退出後倉促應對，使用了開放式的OPENCL架構，其中對CUDA應當說有不少的模仿。開放架構本來是一件好事，但OPENCL的發展一直不盡人意。而且為了相容更多的顯卡，程式中通用層導致的效率損失一直比較大。而實際上，現在的高性能顯卡其實也就剩下了NV/AMD兩家的競爭，這樣基本沒什麼意義的性能損失不能不說讓人糾結。所以在個人工作站和個人裝機市場，通常的選擇都是NV系列的顯卡。
mac電腦在這方面是比較尷尬的，當前的高端系列是MacPro垃圾桶。至少新款的一體機MacPro量產之前，垃圾桶仍然是mac家性能的扛鼎產品。然而其內置的顯卡就是AMD，只能使用OPENCL通用計算框架了。

下麵是蘋果官方給出的一個OPENCL的入門例子，結構很清晰，展示了使用顯卡進行高性能計算的一般結構，我在註釋中增加了中文的說明，相信可以讓你更容易的上手OPENCL顯卡計算。

//
// File:       hello.c
//
// Abstract:   A simple "Hello World" compute example showing basic usage of OpenCL which
//             calculates the mathematical square (X[i] = pow(X[i],2)) for a buffer of
//             floating point values.
//             
//
// Version:    <1.0>
//
// Disclaimer: IMPORTANT:  This Apple software is supplied to you by Apple Inc. ("Apple")
//             in consideration of your agreement to the following terms, and your use,
//             installation, modification or redistribution of this Apple software
//             constitutes acceptance of these terms.  If you do not agree with these
//             terms, please do not use, install, modify or redistribute this Apple
//             software.
//
//             In consideration of your agreement to abide by the following terms, and
//             subject to these terms, Apple grants you a personal, non - exclusive
//             license, under Apple's copyrights in this original Apple software ( the
//             "Apple Software" ), to use, reproduce, modify and redistribute the Apple
//             Software, with or without modifications, in source and / or binary forms;
//             provided that if you redistribute the Apple Software in its entirety and
//             without modifications, you must retain this notice and the following text
//             and disclaimers in all such redistributions of the Apple Software. Neither
//             the name, trademarks, service marks or logos of Apple Inc. may be used to
//             endorse or promote products derived from the Apple Software without specific
//             prior written permission from Apple.  Except as expressly stated in this
//             notice, no other rights or licenses, express or implied, are granted by
//             Apple herein, including but not limited to any patent rights that may be
//             infringed by your derivative works or by other works in which the Apple
//             Software may be incorporated.
//
//             The Apple Software is provided by Apple on an "AS IS" basis.  APPLE MAKES NO
//             WARRANTIES, EXPRESS OR IMPLIED, INCLUDING WITHOUT LIMITATION THE IMPLIED
//             WARRANTIES OF NON - INFRINGEMENT, MERCHANTABILITY AND FITNESS FOR A
//             PARTICULAR PURPOSE, REGARDING THE APPLE SOFTWARE OR ITS USE AND OPERATION
//             ALONE OR IN COMBINATION WITH YOUR PRODUCTS.
//
//             IN NO EVENT SHALL APPLE BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL OR
//             CONSEQUENTIAL DAMAGES ( INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
//             SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
//             INTERRUPTION ) ARISING IN ANY WAY OUT OF THE USE, REPRODUCTION, MODIFICATION
//             AND / OR DISTRIBUTION OF THE APPLE SOFTWARE, HOWEVER CAUSED AND WHETHER
//             UNDER THEORY OF CONTRACT, TORT ( INCLUDING NEGLIGENCE ), STRICT LIABILITY OR
//             OTHERWISE, EVEN IF APPLE HAS BEEN ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
//
// Copyright ( C ) 2008 Apple Inc. All Rights Reserved.
//
 
////////////////////////////////////////////////////////////////////////////////
 
#include <fcntl.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <math.h>
#include <unistd.h>
#include <sys/types.h>
#include <sys/stat.h>
#include <OpenCL/opencl.h>
 
////////////////////////////////////////////////////////////////////////////////
 
// Use a static data size for simplicity
//
#define DATA_SIZE (1024)
 
////////////////////////////////////////////////////////////////////////////////
 
// Simple compute kernel which computes the square of an input array 
// 這是OPENCL用於計算的內核部分源碼，跟C相同的語法格式，通過編譯後將發佈到GPU設備
//（或者將來專用的計算設備）上面去執行。因為顯卡通常有幾十、上百個內核，所以這部分
// 需要設計成可併發的程式邏輯。
// 
const char *KernelSource = "\n" \
"__kernel void square(                                                       \n" \
"   __global float* input,                                              \n" \
"   __global float* output,                                             \n" \
"   const unsigned int count)                                           \n" \
"{                                                                      \n" \
// 併發邏輯主要是在下麵這一行體現的，i的初始值獲取當前內核的id（整數）,根據id計算自己的那一小塊任務
"   int i = get_global_id(0);                                           \n" \
"   if(i < count)                                                       \n" \
"       output[i] = input[i] * input[i];                                \n" \
"}                                                                      \n" \
"\n";
 
////////////////////////////////////////////////////////////////////////////////
 
int main(int argc, char** argv)
{
    int err;                            // error code returned from api calls
      
    float data[DATA_SIZE];              // original data set given to device
    float results[DATA_SIZE];           // results returned from device
    unsigned int correct;               // number of correct results returned
 
    size_t global;                      // global domain size for our calculation
    size_t local;                       // local domain size for our calculation
 
    cl_device_id device_id;             // compute device id 
    cl_context context;                 // compute context
    cl_command_queue commands;          // compute command queue
    cl_program program;                 // compute program
    cl_kernel kernel;                   // compute kernel
    
    cl_mem input;                       // device memory used for the input array
    cl_mem output;                      // device memory used for the output array
    
    // Fill our data set with random float values
    //
    int i = 0;
    unsigned int count = DATA_SIZE;
    //隨機產生一組浮點數據，用於給GPU進行計算
    for(i = 0; i < count; i++)
        data[i] = rand() / (float)RAND_MAX;
    
    // Connect to a compute device
    //
    int gpu = 1;
    // 獲取GPU設備，OPENCL的優勢是可以使用CPU進行模擬，當然這種功能只是為了在沒有GPU設備上進行調試
    // 如果上面變數gpu=0的話，則使用CPU模擬
    err = clGetDeviceIDs(NULL, gpu ? CL_DEVICE_TYPE_GPU : CL_DEVICE_TYPE_CPU, 1, &device_id, NULL);
    if (err != CL_SUCCESS)
    {
        printf("Error: Failed to create a device group!\n");
        return EXIT_FAILURE;
    }
  
    // Create a compute context 
    // 建立一個GPU計算的上下文環境，一組上下文環境保存一組相關的狀態、記憶體等資源
    context = clCreateContext(0, 1, &device_id, NULL, NULL, &err);
    if (!context)
    {
        printf("Error: Failed to create a compute context!\n");
        return EXIT_FAILURE;
    }
 
    // Create a command commands
    //使用獲取到的GPU設備和上下文環境監理一個命令隊列，其實就是給GPU的任務隊列
    commands = clCreateCommandQueue(context, device_id, 0, &err);
    if (!commands)
    {
        printf("Error: Failed to create a command commands!\n");
        return EXIT_FAILURE;
    }
 
    // Create the compute program from the source buffer
    //將內核程式的字元串載入到上下文環境
    program = clCreateProgramWithSource(context, 1, (const char **) & KernelSource, NULL, &err);
    if (!program)
    {
        printf("Error: Failed to create compute program!\n");
        return EXIT_FAILURE;
    }
 
    // Build the program executable
    //根據所使用的設備，將程式編譯成目標機器語言代碼，跟通常的編譯類似，
    //內核程式的語法類錯誤信息都會在這裡出現，所以一般儘可能列印完整從而幫助判斷。
    err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
    if (err != CL_SUCCESS)
    {
        size_t len;
        char buffer[2048];
 
        printf("Error: Failed to build program executable!\n");
        clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);
        printf("%s\n", buffer);
        exit(1);
    }
 
    // Create the compute kernel in the program we wish to run
    //使用內核程式的函數名建立一個計算內核
    kernel = clCreateKernel(program, "square", &err);
    if (!kernel || err != CL_SUCCESS)
    {
        printf("Error: Failed to create compute kernel!\n");
        exit(1);
    }
 
    // Create the input and output arrays in device memory for our calculation
    // 建立GPU的輸入緩衝區，註意READ_ONLY是對GPU而言的，這個緩衝區是建立在顯卡顯存中的
    input = clCreateBuffer(context,  CL_MEM_READ_ONLY,  sizeof(float) * count, NULL, NULL);
    // 建立GPU的輸出緩衝區，用於輸出計算結果
    output = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(float) * count, NULL, NULL);
    if (!input || !output)
    {
        printf("Error: Failed to allocate device memory!\n");
        exit(1);
    }    
    
    // Write our data set into the input array in device memory 
    // 將CPU記憶體中的數據，寫入到GPU顯卡記憶體（內核函數的input部分）
    err = clEnqueueWriteBuffer(commands, input, CL_TRUE, 0, sizeof(float) * count, data, 0, NULL, NULL);
    if (err != CL_SUCCESS)
    {
        printf("Error: Failed to write to source array!\n");
        exit(1);
    }
 
    // Set the arguments to our compute kernel
    // 設定內核函數中的三個參數
    err = 0;
    err  = clSetKernelArg(kernel, 0, sizeof(cl_mem), &input);
    err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &output);
    err |= clSetKernelArg(kernel, 2, sizeof(unsigned int), &count);
    if (err != CL_SUCCESS)
    {
        printf("Error: Failed to set kernel arguments! %d\n", err);
        exit(1);
    }
 
    // Get the maximum work group size for executing the kernel on the device
    //獲取GPU可用的計算核心數量
    err = clGetKernelWorkGroupInfo(kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(local), &local, NULL);
    if (err != CL_SUCCESS)
    {
        printf("Error: Failed to retrieve kernel work group info! %d\n", err);
        exit(1);
    }
 
    // Execute the kernel over the entire range of our 1d input data set
    // using the maximum number of work group items for this device
    // 這是真正的計算部分，計算啟動的時候採用隊列的方式，因為一般計算任務的數量都會遠遠大於可用的內核數量，
    // 在下麵函數中，local是可用的內核數，global是要計算的數量，OPENCL會自動執行隊列，完成所有的計算
    // 所以在前面強調了，內核程式的設計要考慮、並儘力利用這種併發特征
    global = count;
    err = clEnqueueNDRangeKernel(commands, kernel, 1, NULL, &global, &local, 0, NULL, NULL);
    if (err)
    {
        printf("Error: Failed to execute kernel!\n");
        return EXIT_FAILURE;
    }
 
    // Wait for the command commands to get serviced before reading back results
    // 阻塞直到OPENCL完成所有的計算任務
    clFinish(commands);
 
    // Read back the results from the device to verify the output
    // 從GPU顯存中把計算的結果複製到CPU記憶體
    err = clEnqueueReadBuffer( commands, output, CL_TRUE, 0, sizeof(float) * count, results, 0, NULL, NULL );  
    if (err != CL_SUCCESS)
    {
        printf("Error: Failed to read output array! %d\n", err);
        exit(1);
    }
    
    // Validate our results
    // 下麵是使用CPU計算來驗證OPENCL計算結果是否正確
    correct = 0;
    for(i = 0; i < count; i++)
    {
        if(results[i] == data[i] * data[i])
            correct++;
    }
    
    // Print a brief summary detailing the results
    // 顯示驗證的結果
    printf("Computed '%d/%d' correct values!\n", correct, count);
    
    // Shutdown and cleanup
    // 清理各類對象及關閉OPENCL環境
    clReleaseMemObject(input);
    clReleaseMemObject(output);
    clReleaseProgram(program);
    clReleaseKernel(kernel);
    clReleaseCommandQueue(commands);
    clReleaseContext(context);
 
    return 0;
}

因為使用了mac的OPENCL框架，所以編譯的時候要加上對框架的引用，如下所示：

gcc -o hello hello.c -framework OpenCL