OpenCL::Query OpenCL kernel (4)

이 프로그램은 OpenCL 코드를 불러들여서 컴파일하여서 kernel을 만드는데 까지 진행한다.

실행은 아직이다.

우선 크게 보면 main program에서 platform id를 확인하고 

platform id로 context를 만든다.

만들어진 context에 OpenCL 프로그램을 불러들여서 build를 하고 kernel로 준비해 준다.

그리고 해당 kernel의 정보를 받아와서 프린트 해 주는 데까지 진행한다.

OpenCL은 Runtime compile을 한다. 따라서 미리미리 컴파일을 해 고 진행하여야 한다.

이 예제에서는 Program과 Kernel을 만드는데 까지 진행한 예제이다.

(1) Platform ID의 개수를 확인한다.

(2) 확인된 Platform ID의 개수 만큼 공간을 확보한다.

(3) 확보된 메모리 공간에 Platform ID를 받아온다.

(4) Platform ID개수만큼 

     1) Device ID중 GPU 인것을 받아온다.

     2) GPU인 것이 없으면 CPU 인 것을 받아온다.

     3) Device ID에 Context를 생성한다.

     4) OpenCL 프로그램 소스를 로 드

     5) 로드된 프로그램소스를 빌드하여서 프로그램 화 한다.

         이 단계는 CreateProgramWithSource -> BuildProgram 을 거치는 2단계로 구성된다.

     6) 빌드 (컴파일)에서 문제가 생기면 그 정보(LOG)를 받아서 프린트 해준다.

 (5) clCreateKernelsInProgram 프로그램을 가지고 커널을 만든다.

      마찬가지로 2단계로 한다. 

      우선 커널의 개수를 확인하고

      확인된 커널의 개수만큼 메모리를 할당받고

      해당 할당받은 메모리에다가 커널을 빌드한다.

 (6) 확인된 커널의 개수만큼 

     1) 커널의 이름을 받아오고

     2) 인자의 개수를 받아온다.

     3) 받아온 커널의 이름과 인자의 개수를 프린트한다.

복잡하지만 줄여보면 다음과 같다.

기본 프레임은

  (1) 필요한 요소의 개수를 물어보고

  (2) 그만큼 메모리 공간을 할당 받고

  (3) 할당받은 메모리 공간에 요소를 받아온다.

이다. 

이것을 

 Platform ID -> Device ID -> context -> Kernel 

순으로 진행한다.

  

#include <stdio.h>
#include <stdlib.h>
#include <sys/types.h>
#include <alloca.h>

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

/*
 프로그램을 읽어 들인다.
 OpenCL은 On-The-Fly Compile을 수행한다.
 */
void loadProgramSource(const char** files,
                       size_t length,
                       char** buffer,
                       size_t* sizes) {
    
    /* Read each source file (*.cl) and store the contents into a temporary datastore */
    for(size_t i=0; i < length; i++) {
        FILE* file = fopen(files[i], "r");
        if(file == NULL) {
            perror("Couldn't read the program file");
            exit(1);
        }
        
        /* 파일의 끝까지 진행하여서 크기를 확인하고 다시 처음으로 돌아온다.*/
        fseek(file, 0, SEEK_END);
        sizes[i] = ftell(file);
        rewind(file); // reset the file pointer so that 'fread' reads from the front
        
        /* 크기 +1 만큼 할당한다음 */
        buffer[i] = (char*)malloc(sizes[i]+1);
        buffer[i][sizes[i]] =  '\0';
        
        /* 한방에 읽어 들이다. */
        fread(buffer[i], sizeof(char), sizes[i], file);
        fclose(file);
    }
}

int main(int argc, char** argv) {
    
    /* OpenCL 1.1 data structures */
    cl_platform_id* platforms;
    cl_program program;
    cl_device_id device;
    cl_context context;
    
    /* OpenCL 1.1 scalar data types */
    cl_uint numOfPlatforms;
    cl_int  error;
    
    /*
     Get the number of platforms
     Remember that for each vendor's SDK installed on the computer,
     the number of available platform also increased.
     */

    /* Step 1 : Platform ID의 개수를 확인한다. */
    error = clGetPlatformIDs(0, NULL, &numOfPlatforms);
    if(error != CL_SUCCESS) {
        perror("Unable to find any OpenCL platforms");
        exit(1);
    }
    
    /* Step 2 : 확인된 Platform ID의 개수만큼 메모리 공간을 확보한다. */
    platforms = (cl_platform_id*) alloca(sizeof(cl_platform_id) * numOfPlatforms);
    printf("Number of OpenCL platforms found: %d\n", numOfPlatforms);
    
    /* Step 3 : 확보된 메모리 공간에 Platform ID 데이터를 받아온다. */
    error = clGetPlatformIDs(numOfPlatforms, platforms, NULL);
    if(error != CL_SUCCESS) {
        perror("Unable to find any OpenCL platforms");
        exit(1);
    }
    
    // GPU를 찾는다.
    // Search for a CPU/GPU device through the installed platforms
    // Build a OpenCL program and do not run it.

    // Platform의 개수 만큼 진행한다.
    for(cl_uint i = 0; i < numOfPlatforms; i++ ) {

        // Get the GPU device

        // Step 1 : Device ID 중에서 GPU 인 것을  확인한다.
        error = clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_GPU, 1, &device, NULL);
        if(error != CL_SUCCESS) {
            // Otherwise, get the CPU
            // Step 1-1 : Device ID 중 GPU가 없으면 CPU인 것을 확인한다.
            error = clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_CPU, 1, &device, NULL);
        }
        if(error != CL_SUCCESS) {
            perror("Can't locate any OpenCL compliant device");
            exit(1);
        }
        
        /* Context를 생성한다. */
        // Step 2 : Device ID에 Context를 생성한다.
        context = clCreateContext(NULL, 1, &device, NULL, NULL, &error);
        if(error != CL_SUCCESS) {
            perror("Can't create a valid OpenCL context");
            exit(1);
        }
        
        /* Load the two source files into temporary datastores */
        const char *file_names[] = {"/Users/freegear/Desktop/OpenCL/cl_test1/KernelQuery/KernelQuery/simple.cl", "/Users/freegear/Desktop/OpenCL/cl_test1/KernelQuery/KernelQuery/simple_2.cl"};
        const int NUMBER_OF_FILES = 2;
        char* buffer[NUMBER_OF_FILES];
        size_t sizes[NUMBER_OF_FILES];
        loadProgramSource(file_names, NUMBER_OF_FILES, buffer, sizes);
        
        /* OpenCL program object를 생성한다.*/

        // Step 3 : Context에 프로그램 source를 컴파일하여 program을 생성한다.
        program = clCreateProgramWithSource(context, NUMBER_OF_FILES, (const char**)buffer, sizes, &error);
        if(error != CL_SUCCESS) {
            perror("Can't create the OpenCL program object");
            exit(1);
        }
        
        /* Build OpenCL program object and dump the error message, if any */
        char *program_log;
        const char options[] = "-cl-finite-math-only -cl-no-signed-zeros";
        size_t log_size;
        //error = clBuildProgram(program, 1, &device, argv[1], NULL, NULL);
        // Uncomment the line below, comment the line above; re-build the program to use build options statically
        // Step 4: 생성된 프로그램을 빌드한다.
        error = clBuildProgram(program, 1, &device, options, NULL, NULL);
        
        if(error != CL_SUCCESS) {
            // If there's an error whilst building the program, dump the log
            clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
            program_log = (char*) malloc(log_size+1);
            program_log[log_size] = '\0';
            
            // Step 4-1: 빌드 정보에 대한 Log를 확인한다.
            clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG,
                                  log_size+1, program_log, NULL);
            
            printf("\n=== ERROR ===\n\n%s\n=============\n", program_log);
            free(program_log);
            exit(1);
        }
        
        /* Query the program as to how many kernels were detected */
        cl_uint numOfKernels;

        // Step 5 : Kernel의 개수를 학인한다.
        error = clCreateKernelsInProgram(program, 0, NULL, &numOfKernels);
        if (error != CL_SUCCESS) {
            perror("Unable to retrieve kernel count from program");
            exit(1);
        }
        

        // Step 6 : 확인된 Kernel의 개수만큼 메모리를 할당한다.
        cl_kernel* kernels = (cl_kernel*) alloca(sizeof(cl_kernel) * numOfKernels);
        error = clCreateKernelsInProgram(program, numOfKernels, kernels, NULL);
        
        // Step 7 : 확인된 Kernel의 개수만큼 반복한다.
        for(cl_uint i = 0; i < numOfKernels; i++) {
            char kernelName[32];
            cl_uint argCnt;
            // Step 7-1 : 확인된 Kernel의 이름을 받아온다.
            clGetKernelInfo(kernels[i], CL_KERNEL_FUNCTION_NAME, sizeof(kernelName), kernelName, NULL);
            // Step 7-2 : 확인된 Kernel 인자의 개수를 받아온다.
            clGetKernelInfo(kernels[i], CL_KERNEL_NUM_ARGS, sizeof(argCnt), &argCnt, NULL);
            printf("Kernel name: %s with arity: %d\n", kernelName, argCnt);
        }
        

     
        /* Clean up */

        // Step 8 :  청소하기
        for(cl_uint i = 0; i < numOfKernels; i++) { clReleaseKernel(kernels[i]); }
        for(i=0; i< NUMBER_OF_FILES; i++) { free(buffer[i]); }
        clReleaseProgram(program);
        clReleaseContext(context);
    }
}

실행 결과는 아래와 같다.

Number of OpenCL platforms found: 1

Kernel name: simpleAdd with arity: 3

Kernel name: simpleAdd_2 with arity: 3

Program ended with exit code: 0

찾아낸 플랫폼의 개수는 1개이고 여기서 포함되는 커널의 개수는 2개이다.

각각 

   simpleAdd

   simpleAdd_2

가 된다.

OpenCL source 코드는 아래와 같이 되어 있고 여기서는 아직 중요한 포인트가 아니므로 그냥 넘어간다.

 simple.cl

__kernel void simpleAdd(__global float *a,

                   __global float *b,

                   __global float *c) {

   

   *c = *a + *b;

}

simple_2.cl

__kernel void simpleAdd_2(__global float *a,

                   __global float *b,

                   __global float *c) {

   

   *c = *a + *b;

}