#include <math.h>


//在使用 OpenCL API 之前,和絕大部份所有其它的 API 一樣,都需要 include 相關的 header 檔案。
//由於在 MacOS X 10.6 下 OpenCL 的 header 檔案命名方式和在其它作業系統下不同,
//因此,通常要使用一個 #ifdef 來進行區分。如下所示:

//#ifdef __APPLE__
//#include <OpenCL/opencl.h>
//#else
//#include <CL/cl.h>
//#endif

#include "CL\opencl.h"
#include "utility.h"

static const size_t vectorSize = 4096; //must be evenly disible by workSize
static const size_t workSize = 256;


int main(void)
{
    cl_int error;

    //Setup Platform
    cl_uint testPlatformCount;

    //要先取得系統上所有的 OpenCL platform。在 MacOS X 10.6 下,
    //目前只有一個由 Apple 提供的 OpenCL platform,
    //但是在其它系統上,可能會有不同廠商提供的多個不同的 OpenCL platform,
    //因此需要先取得 platform 的數目:
    //大部份的 OpenCL API 會傳回錯誤值。如果傳回值是 CL_SUCCESS 則表示執行成功,
    //否則會傳回某個錯誤值,表示失敗的原因。

    error = clGetPlatformIDs(0, NULL, &testPlatformCount);
    
    
    //Get Platform ID
    cl_platform_id testPlatform;
    //接著,再取得 platform 的 ID,這在建立 OpenCL context 時會用到:
    error = clGetPlatformIDs(1, &testPlatform, NULL);
    assert(error==CL_SUCCESS);
    print_platform_info(&testPlatform);


    //Setup Device
    cl_uint test_deviceCount;
    //取得系統上所有的 OpenCL Device
    error = clGetDeviceIDs(testPlatform, CL_DEVICE_TYPE_ALL, 0, NULL, &test_deviceCount);
    

    //Get Device ID
    cl_device_id test_device;
    

    //CL_DEVICE_TYPE_CPU:使用 CPU 裝置
    //CL_DEVICE_TYPE_GPU:使用顯示晶片裝置
    //CL_DEVICE_TYPE_ACCELERATOR:特定的 OpenCL 加速裝置,例如 CELL
    //CL_DEVICE_TYPE_DEFAULT:系統預設的 OpenCL 裝置
    //CL_DEVICE_TYPE_ALL:所有系統中的 OpenCL 裝置

    
    error = clGetDeviceIDs(testPlatform, CL_DEVICE_TYPE_ALL, 1, &test_device, NULL);
    assert(error==CL_SUCCESS);
    print_device_info(&test_device);
    

    //Create Context
    //要建立一個 OpenCL context。如下:
 
    cl_context testcontext = clCreateContext(0, 1, &test_device, NULL, NULL, &error);
    assert(error==CL_SUCCESS);

    
    //執行這個程式,如果建立 OpenCL context 成功的話,應該會顯示出找到的 OpenCL 裝置的名稱,例如
    //clGetDeviceInfo(devices[0], CL_DEVICE_NAME, cb, &devname[0], 0);
    
    //-------------------------------------------------------------------------------------------------------
    //Create Command queue
    //大部份 OpenCL 的操作,都要透過 command queue。Command queue 可以接收對一個 OpenCL 裝置的各種操作,
    //並按照順序執行(OpenCL 也容許把一個 command queue 指定成不照順序執行,即 out-of-order execution,
    //但是這裡先不討論這個使用方式)。所以,下一步是建立一個 command queue:

    cl_command_queue testqueue = clCreateCommandQueue(testcontext, test_device, 0, &error);
    assert(error==CL_SUCCESS);

    //上面的程式中,是把裝置列表中的第一個裝置(即 test_device)建立 command queue。
    //如果想要同時使用多個 OpenCL 裝置,則每個裝置都要有自己的 command queue。


    //配置記憶體並複製資料

    //要使用 OpenCL 裝置進行運算時,通常會需要在 OpenCL 裝置上配置記憶體,並把資料從主記憶體中複製到裝置上。
    //有些 OpenCL 裝置可以直接從主記憶體存取資料,但是速度通常會比較慢,因為 OpenCL 裝置(例如顯示卡)通常會有專用的高速記憶體。
    //以下的程式配置三塊記憶體:

    cl_mem kernelIn = clCreateBuffer(testcontext, CL_MEM_READ_ONLY, sizeof(cl_float) * vectorSize, NULL, &error);
    assert(error==CL_SUCCESS);
    cl_mem kernelIn2 = clCreateBuffer(testcontext, CL_MEM_READ_ONLY, sizeof(cl_float) * vectorSize, NULL, &error);
    assert(error==CL_SUCCESS);

    cl_mem kernelOut = clCreateBuffer(testcontext, CL_MEM_WRITE_ONLY, sizeof(cl_float) * vectorSize, NULL, &error);
    assert(error==CL_SUCCESS);

    /*
    clCreateBuffer 函式可以用來配置記憶體。
    第二個參數可以指定記憶體的使用方式,包括:
        CL_MEM_READ_ONLY:表示 OpenCL kernel 只會對這塊記憶體進行讀取的動作
        CL_MEM_WRITE_ONLY:表示 OpenCL kernel 只會對這塊記憶體進行寫入的動作
        CL_MEM_READ_WRITE:表示 OpenCL kernel 會對這塊記憶體進行讀取和寫入的動作
        CL_MEM_USE_HOST_PTR:表示希望 OpenCL 裝置直接使用指定的主記憶體位址。要注意的是,如果 OpenCL 裝置無法直接存取主記憶體,
                        它可能會將指定的主記憶體位址的資料複製到 OpenCL 裝置上。
        CL_MEM_ALLOC_HOST_PTR:表示希望配置的記憶體是在主記憶體中,而不是在 OpenCL 裝置上。不能和 CL_MEM_USE_HOST_PTR 同時使用。
        CL_MEM_COPY_HOST_PTR:將指定的主記憶體位址的資料,複製到配置好的記憶體中。不能和 CL_MEM_USE_HOST_PTR 同時使用。
    第三個參數是指定要配置的記憶體大小,以 bytes 為單位。在上面的程式中,指定的大小是 sizeof(cl_float) * DATA_SIZE。
    第四個參數是指定主記憶體的位置。因為對 cl_a 和 cl_b 來說,在第二個參數中,指定了 CL_MEM_COPY_HOST_PTR,因此要指定想要複製的資料的位址。cl_res 則不需要指定。
    第五個參數是指定錯誤碼的傳回位址。在這裡並沒有使用到。
    
    如果 clCreateBuffer 因為某些原因無法配置記憶體(例如 OpenCL 裝置上的記憶體不夠),則會傳回 0。要釋放配置的記憶體,可以使用 clReleaseMemObject 函式。
    */

    //Inputs and Outputs to Kernel, X and Y are inputs, Z is output
    void *X, *Y, *Z;
    
    //Allocates memory with value from 0 to 1000
    float LO= 0;   float HI=1000;
    allocate_generate(&X, &Y, &Z, LO, HI, vectorSize);
    //Create Buffers for input and output

    //Write data to device
    //
    error = clEnqueueWriteBuffer(testqueue, kernelIn, CL_FALSE, 0, sizeof(cl_float) * vectorSize, X, 0, NULL, NULL);
    error = clEnqueueWriteBuffer(testqueue, kernelIn2, CL_FALSE, 0, sizeof(cl_float) * vectorSize, Y, 0, NULL, NULL);
    clFinish(testqueue);
    assert(error==CL_SUCCESS);

    //-------------------------------------------------------------------------------------------------------
    //執行 OpenCL kernel

    // create the kernel
    const char *kernel_name = "SimpleKernel";

    size_t lengths;
    unsigned char* test_binaries = get_binary("SimpleKernel.aocx", &lengths);
    cl_int kernel_status;
    
    // Create the Program from the AOCX file.
    /*
    本例中的 kernel 程式是作為靜態字串讀入的(單獨的文字檔也一樣),
    所以使用的是 clCreateProgramWithSource,如果你不想讓 kernel 程式讓其他人看見,可以先生成二進位檔案,
    再通過 clCreateProgramWithBinary 函數動態讀入二進位檔案,做一定的保密。詳細請參閱 OpenCL Specifications。    
    */
    
    cl_program program = clCreateProgramWithBinary(testcontext, 1, &test_device, &lengths, (const unsigned char**)&test_binaries, &kernel_status, &error);

    assert(error==CL_SUCCESS);
      
    // build the program
    //     Compile the Kernel...
    //使用 clBuildProgram 編譯 kernel:
    error = clBuildProgram(program, 1, &test_device, NULL, NULL, NULL);
    assert(error==CL_SUCCESS);


    // create the kernel
    //////////////       Find Kernel in Program
    //一個 OpenCL kernel 程式裡面可以有很多個函式。因此,還要取得程式中函式的進入點:
    cl_kernel kernel = clCreateKernel(program, kernel_name, &error);
    assert(error==CL_SUCCESS);

    //////////////     Set Arguments to the Kernels
    /*
    先設定好函式的參數。adder 函式有三個參數要設定:
    設定參數是使用 clSetKernelArg 函式。它的參數很簡單:
        第一個參數是要設定的 kernel object,
        第二個是參數的編號(從 0 開始),
        第三個參數是要設定的參數的大小,
        第四個參數則是實際上要設定的參數內部。
        以這裡的 adder 函式來說,三個參數都是指向 memory object 的指標。
    設定好參數後,就可以開始執行了。如下:
    */


    error = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&kernelIn);
    assert(error==CL_SUCCESS);
    error = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*)&kernelIn2);
    assert(error==CL_SUCCESS);
    error = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void*)&kernelOut);
    assert(error==CL_SUCCESS);

    printf("\nLaunching the kernel...\n");


    // launch kernel
    /*
    clEnqueueNDRangeKernel 會把執行一個 kernel 的動作加到 command queue 裡面。
    第三個參數(1)是指定 work item 數目的維度,在這裡就是一維。
    第五個參數是指定 work item 的總數目,也就是 DATA_SIZE。後面的參數現在暫時先不用管。如果成功加入的話,會傳回 CL_SUCCESS。否則會傳回錯誤值。
    
    在執行 kernel 被加到 command queue 之後,就可能會開始執行(如果 command queue 現在沒有別的工作的話)。
    但是 clEnqueueNDRangeKernel 是非同步的,也就是說,它並不會等待 OpenCL 裝置執行完畢才傳回。這樣可以讓 CPU 在 OpenCL 裝置在進行運算的同時,
    進行其它的動作。    
    */
    error = clEnqueueNDRangeKernel(testqueue, kernel, 1, NULL, &vectorSize, &workSize, 0, NULL, NULL);
    clFinish(testqueue);
    assert(error==CL_SUCCESS);

    // read the output
    /*
    由於執行的結果是在 OpenCL 裝置的記憶體中,所以要取得結果,需要把它的內容複製到 CPU 能存取的主記憶體中。這可以透過下面的程式完成:
    clEnqueueReadBuffer 函式會把「將記憶體資料從 OpenCL 裝置複製到主記憶體」的動作加到 command queue 中。
    第三個參數表示是否要等待複製的動作完成才傳回,CL_TRUE 表示要等待。第五個參數是要複製的資料大小,
    第六個參數則是目標的位址。
    由於這裡指定要等待複製動作完成,所以當函式傳回時,資料已經完全複製完成了。
    */
    
    error = clEnqueueReadBuffer(testqueue, kernelOut, CL_TRUE, 0, sizeof(cl_float) * vectorSize, Z, 0, NULL, NULL);
    assert(error==CL_SUCCESS);
    
    
    //最後是進行驗證,確定資料正確:
    void * CalcZ = malloc(sizeof(float)*vectorSize);

    for (int i=0; i<vectorSize; i++)
    {
        //////////////  Equivalent Code runnign on CPUs
        ((float*) CalcZ)[i]=sin(((float*) X)[i] + ((float*) Y)[i]); 
                
    }
    
    //-------------------------------------------------------------------------------------------------------


    // Clean up Stuff
    if(kernel) 
        clReleaseKernel(kernel);  
    if(program) 
        clReleaseProgram(program);
    if(testqueue) 
        clReleaseCommandQueue(testqueue);
    if(testcontext) 
        clReleaseContext(testcontext);
    if(kernelIn) 
        clReleaseMemObject(kernelIn);
    if(kernelOut) 
        clReleaseMemObject(kernelOut);
    if(X) 
        free (X);
    if(Y) 
        free (Y);
    if(Z) 
        free (Z);
    if(CalcZ) 
        free (CalcZ);

    return 1;
}
 

 

 

 


//編譯 OpenCL kernel 程式               
//現在執行 OpenCL kernel 的準備工作已經大致完成了。所以,現在剩下的工作,就是把 OpenCL kernel 程式編釋並執行。

__kernel  void SimpleKernel(global const float * restrict in, global const float * restrict in2, __global float * restrict out)
{
    // Get Index of the work item
    int index = get_global_id(0);

    //Perform the Math Operation
    out[index] = sin (in[index] + in2[index]);
}
 

 

 

其他相關資料
 

OpenCL 教學(一)

http://www.kimicat.com/opencl-1/opencl-jiao-xue-yi

 

 

 

立你斯 發表在 痞客邦 留言(0) 人氣()