Computer Vision2014.06.01 06:41

1차원 데이터 열에 대한 커널의 실행은


clEnqueueTask

2혹은 3차원 데이터에 대한 커널의 실행은

clEnqueueNDRangeKernel

에 의해서 실행 된다.

cl_int clEnqueueNDRangeKernel (cl_command_queue command_queue,
 cl_kernel kernel,
 cl_uint work_dim,
 const size_t *global_work_offset,
 const size_t *global_work_size,
 const size_t *local_work_size,
 cl_uint num_events_in_wait_list,
 const cl_event *event_wait_list,
 cl_event *event)


command_queue : 커멘드 큐…
kernel : Kernel
work_dim   : 0 보다는 커야 하지만 3보다는 작거나 같아야 한다.
global work offset : 현재 버전에서는 그냥 0 이 되어야 한다.
global_work_size :  포인터로 정의되고 
                                work_dim의 영향을 받는다. 따라서
                                global_work_size[0] * global_work_size[1] ….  * global_work_size[work_dim-1]
                                이 전체 work_itme의 개수이다.
                                global work_size는 sizeof(size_t)를 초과할 수 없다.
                                무슨 말이냐 하면, 만약 특정 디바이스에서 size를 나타내는 최대 크기가 32이면 
                                즉 2^32이면 이보다는 작아야 한다는 것이다.
                                sizeof(size_t)는 결국 메모리 공간의 포인터의 크기를 의미하는데 32보다 크면 
                               포인터로 지정할 수 있는 공간보다 크다는
                               의미이므로 오류가 되는 것이다.
                                이런 경우에  CL_OUT_OF_RESOURCES   오류가 발생한다.

local_work_size :   work_dim의 배열을 가리키는 unsigned 값이다.
                              work_item 숫자이고 이것은 work-group을 만든다. 그리고 커널이 실행되는 횟수이다.  
                              global_work_size와 마찬가지로
                              local_work_size[0] * local_work_size[1] ….  * local_work_size[work_dim-1] 
                              으로 해서 work_item의 숫자가 결졍된다.  그리고 이 숫자는 
                              CL_DEVICE_MAX_WORK_GROUP_SIZE보다 작아야 한다. 
                              local_work_size[0] <= CL_DEVICE_MAX_WORK_GROUP_SIZE[0]
                              local_work_size[1] <= CL_DEVICE_MAX_WORK_GROUP_SIZE[1]
                                   ….
                              local_work_size[2] <= CL_DEVICE_MAX_WORK_GROUP_SIZE[2]
                    
                              local_work_size는 global_work_item을 workgroup instance로 나누어서 할당하는것을 
                              결정하는데 사용한다.
                              만약 local_work_size가 명시되어 있다면, global_work_size[0] , 
                             global_work_size[1] … global_work_size[work_dim-1]은
                              local_work_size[0] , local_work_size[1] … local_work_size[2]의 값으로 나누어
                              떨어져야 한다.
                                   
                              work-group size는 커널이 프로그램 소스에서  
                              __attribute__((reqd_work_group_size(X,Y,Z))) 에 의해서 사용할 수 있다.
                              이 경우 local_work_size는 이 reqd_work_group_size에 명시된 숫자와 같아야 하낟.

                              귀찮으면 OpenCL이 직접 실행시에 지정할 수 있다. 이럴 경우에 NULL로 명시한다.

(주) global work size는 결국 작업해야 할 전체 work item의 개수를 의미한다.
      이를 local work size로 나누어서 분할하여 작업하게 되는 것이다.
      
각각의 차원을 chunk로 나눌 것이다. 이때 나누어지는 chunk의 개수는 2,3,4….CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS 까지 가능하다.
이 값은 아래와 같은 방법으로 확인이 가능하다.

  /* the cl_device_info are preprocessor directives defined in cl.h */
  
switch (param_name) {
        …
    case CL_DEVICE_VENDOR_ID : 
    
case CL_DEVICE_MAX_COMPUTE_UNITS : 
    case CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS : {
            cl_uint* ret = (cl_uint*) alloca(
sizeof(cl_uint) * paramSize);
            error = clGetDeviceInfo( id, param_name, paramSize, ret, 
NULL );
            
if (error != CL_SUCCESS ) {
                perror(
"Unable to obtain device info for param\n");
                
return;
            }
            
switch (param_name) {
                
case CL_DEVICE_VENDOR_ID: printf("\tVENDOR ID: 0x%x\n", *ret); break;
                
case CL_DEVICE_MAX_COMPUTE_UNITS: printf("\tMaximum number of parallel compute units: %d\n", *ret); break;
                
case CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS: 
                                   printf("\tMaximum dimensions for global/local work-item IDs: %d\n", *ret); break;
            }
         }
break;


processing group을 병렬로 실행할 수 있고, 이 것을 work group으로 부른다. work group은 work item이라 부르는 개별적인 병렬 실행 단위를 가진다. 저자는 이 것을 executable thread라고 개념적으로 보고 있다.



위의 예에서 보면 2차원으로 구성된 공간을 9개의 work group으로 나눈다. 각각의  work group은 4x4의 work item으로 구성된다.
work item을 할당하는 데에는 2가지 방법이 있다.
하나의 work item을 하나의 work group에 할당하는 방법과
n x n 개의 work item을 (이 경우 4x4가 된다.) work group에 할당하는 방법이 있다.

우선 첫번째 방식으로 진행한다고 가정하고 진행한다.

…. 코드…

프로그램을 실행시키면 오류가 발생한다.

Number of OpenCL platforms found: 1
Number of detected OpenCL devices: 3
Running on CPU ........
=> Kernel name: copy2Dfloat4 with arity: 2
=> About to create command queue and enqueue this kernel...
Unable to enqueue task to command-queue: No such file or directory 
...

어떤 원인으로 발생된지 모르므로 이를 다시 확인하기 위해서 코드를 추가하였다.


……..
void Check_clEnqueueNDRangeKernel_Error(cl_int error)
{
    
    
switch( error )
    {
        
case CL_INVALID_PROGRAM_EXECUTABLE : printf("CL_INVALID_PROGRAM_EXECUTABLE\n"); break;
        
case CL_INVALID_COMMAND_QUEUE      : printf("CL_INVALID_COMMAND_QUEUE\n"); break;
        
case CL_INVALID_KERNEL            : printf("CL_INVALID_KERNEL\n");break;
        
case CL_INVALID_CONTEXT            : printf("CL_INVALID_CONTEXT\n");break;
        
case CL_INVALID_KERNEL_ARGS        : printf("CL_INVALID_KERNEL_ARGS\n"); break;
        
case CL_INVALID_WORK_DIMENSION    : printf("CL_INVALID_WORK_DIMENSION\n"); break;
        
case CL_INVALID_WORK_GROUP_SIZE    : printf("CL_INVALID_WORK_GROUP_SIZE\n") ; break;
        
case CL_INVALID_WORK_ITEM_SIZE    : printf("CL_INVALID_WORK_ITEM_SIZE\n"); break ;
        
case CL_INVALID_GLOBAL_OFFSET      : printf("CL_INVALID_GLOBAL_OFFSET\n"); break;
        
case CL_OUT_OF_RESOURCES          : printf("CL_OUT_OF_RESOURCES\n"); break;
        
case CL_MEM_OBJECT_ALLOCATION_FAILURE : printf("CL_MEM_OBJECT_ALLOCATION_FAILURE\n");break;
        
case CL_INVALID_EVENT_WAIT_LIST    : printf("CL_INVALID_EVENT_WAIT_LIST\n");break;
        
case CL_OUT_OF_HOST_MEMORY        : printf("CL_OUT_OF_HOST_MEMORY\n");break;       
    }   
}
….


추가된 코드에 의해서 

Number of OpenCL platforms found: 1
Number of detected OpenCL devices: 3
Running on CPU ........
=> Kernel name: copy2Dfloat4 with arity: 2
=> About to create command queue and enqueue this kernel...
Unable to enqueue task to command-queue: No such file or directory
CL_INVALID_WORK_GROUP_SIZE
Program ended with exit code: 1


과 같은 것을 얻을 수 있다.

확인하기 위해서 코드를 보면


            /* Enqueue the kernel to the command queue */
                
size_t globalThreads[2];
                globalThreads[
0]=1024;
                globalThreads[
1]=1024;
                
size_t localThreads[2];
                localThreads[
0] = 64;
                localThreads[
1] = 2;

                
cl_event evt;
            error = 
clEnqueueNDRangeKernel(    cQ,  // command queue
                                               kernels[j], // kernel
                                               
2,   // work_dim
                                               
0,   // global work_offset
                                               globalThreads, // global work size
                                               localThreads,  // local work_size
                                               
0,              // event related parameter
                                               
NULL, &evt);
                
clWaitForEvents(1, &evt);
            
if (error != CL_SUCCESS) { 
                
perror("Unable to enqueue task to command-queue");
                    
Check_clEnqueueNDRangeKernel_Error(error);
                
exit(1);
            }




으로 되어 있다.


계속 같은 에러가 나타나서 다음과 같이 디바이스 정보를 얻어내는 함수를 넣고 실행을 해 본다.


void displayDeviceDetails(cl_device_id id,
                          
cl_device_info param_name,
                          
const char* paramNameAsStr) {
    
cl_int error = 0;
    
size_t paramSize = 0;
    
    error = 
clGetDeviceInfo( id, param_name, 0NULL, &paramSize );
    
if (error != CL_SUCCESS ) {
        
perror("Unable to obtain device info for param\n");
        
return;
    }
    
    
/* the cl_device_info are preprocessor directives defined in cl.h */
    
switch (param_name) {
        
case CL_DEVICE_TYPE: {
            
cl_device_type* devType = (cl_device_type*) alloca(sizeof(cl_device_type) * paramSize);
            error = 
clGetDeviceInfo( id, param_name, paramSize, devType, NULL );
            
if (error != CL_SUCCESS ) {
                
perror("Unable to obtain device info for param\n");
                
return;
            }
            
switch (*devType) {
                
case CL_DEVICE_TYPE_CPU : printf("CPU detected\n");break;
                
case CL_DEVICE_TYPE_GPU : printf("GPU detected\n");break;
                
case CL_DEVICE_TYPE_ACCELERATOR : printf("Accelerator detected\n");break;
                
case CL_DEVICE_TYPE_DEFAULT : printf("default detected\n");break;
            }
        }
break;
        
case CL_DEVICE_VENDOR_ID :
        
case CL_DEVICE_MAX_COMPUTE_UNITS :
        
case CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS : {
            
cl_uint* ret = (cl_uint*) alloca(sizeof(cl_uint) * paramSize);
            error = 
clGetDeviceInfo( id, param_name, paramSize, ret, NULL );
            
if (error != CL_SUCCESS ) {
                
perror("Unable to obtain device info for param\n");
                
return;
            }
            
switch (param_name) {
                
case CL_DEVICE_VENDOR_IDprintf("\tVENDOR ID: 0x%x\n", *ret); break;
                
case CL_DEVICE_MAX_COMPUTE_UNITSprintf("\tMaximum number of parallel compute units: %d\n", *ret); break;
                
case CL_DEVICE_MAX_WORK_ITEM_DIMENSIONSprintf("\tMaximum dimensions for global/local work-item IDs: %d\n", *ret); break;
            }
        }
break;
        
case CL_DEVICE_MAX_WORK_ITEM_SIZES : {
            
cl_uint maxWIDimensions;
            
size_t* ret = (size_t*) alloca(sizeof(size_t) * paramSize);
            error = 
clGetDeviceInfo( id, param_name, paramSize, ret, NULL );
            
            error = 
clGetDeviceInfo( id, CL_DEVICE_MAX_WORK_ITEM_DIMENSIONSsizeof(cl_uint), &maxWIDimensions, NULL );
            
if (error != CL_SUCCESS ) {
                
perror("Unable to obtain device info for param\n");
                
return;
            }
            
printf("\tMaximum number of work-items in each dimension: ( ");
            
for(cl_int i =0; i < maxWIDimensions; ++i ) {
                
printf("%d ", ret[i]);
            }
            
printf(" )\n");
        }
break;
        
case CL_DEVICE_MAX_WORK_GROUP_SIZE : {
            
size_t* ret = (size_t*) alloca(sizeof(size_t) * paramSize);
            error = 
clGetDeviceInfo( id, param_name, paramSize, ret, NULL );
            
if (error != CL_SUCCESS ) {
                
perror("Unable to obtain device info for param\n");
                
return;
            }
            
printf("\tMaximum number of work-items in a work-group: %d\n", *ret);
        }
break;
        
case CL_DEVICE_NAME :
        
case CL_DEVICE_VENDOR : {
            
char data[48];
            error = 
clGetDeviceInfo( id, param_name, paramSize, data, NULL );
            
if (error != CL_SUCCESS ) {
                
perror("Unable to obtain device name/vendor info for param\n");
                
return;
            }
            
switch (param_name) {
                
case CL_DEVICE_NAME : printf("\tDevice name is %s\n", data);break;
                
case CL_DEVICE_VENDOR : printf("\tDevice vendor is %s\n", data);break;
            }
        } 
break;
        
case CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE: {
            
cl_uint* size = (cl_uint*) alloca(sizeof(cl_uint) * paramSize);
            error = 
clGetDeviceInfo( id, param_name, paramSize, size, NULL );
            
if (error != CL_SUCCESS ) {
                
perror("Unable to obtain device name/vendor info for param\n");
                
return;
            }
            
printf("\tDevice global cacheline size: %d bytes\n", (*size)); break;
        } 
break;
        
case CL_DEVICE_GLOBAL_MEM_SIZE:
        
case CL_DEVICE_MAX_MEM_ALLOC_SIZE: {
            
cl_ulong* size = (cl_ulong*) alloca(sizeof(cl_ulong) * paramSize);
            error = 
clGetDeviceInfo( id, param_name, paramSize, size, NULL );
            
if (error != CL_SUCCESS ) {
                
perror("Unable to obtain device name/vendor info for param\n");
                
return;
            }
            
switch (param_name) {
                
case CL_DEVICE_GLOBAL_MEM_SIZEprintf("\tDevice global mem: %ld mega-bytes\n", (*size)>>20); break;
                
case CL_DEVICE_MAX_MEM_ALLOC_SIZEprintf("\tDevice max memory allocation: %ld mega-bytes\n", (*size)>>20); break;
            }
        } 
break;
            
    } 
//end of switch
    
}



void displayDeviceInfo(cl_platform_id platform_id,
                       
cl_device_type dev_type) {
    
/* OpenCL 1.1 device types */
    
cl_int error = 0;
    
cl_uint numOfDevices = 0;
    
    
/* Determine how many devices are connected to your platform */
    error = 
clGetDeviceIDs(platform_id, dev_type, 0NULL, &numOfDevices);
    
if (error != CL_SUCCESS ) {
        
perror("Unable to obtain any OpenCL compliant device info");
        
exit(1);
    }
    
cl_device_id* devices = (cl_device_id*) alloca(sizeof(cl_device_id) * numOfDevices);
    
    
/* Load the information about your devices into the variable 'devices' */
    error = 
clGetDeviceIDs(platform_id, dev_type, numOfDevices, devices, NULL);
    
if (error != CL_SUCCESS ) {
        
perror("Unable to obtain any OpenCL compliant device info");
        
exit(1);
    }
    
printf("Number of detected OpenCL devices: %d\n", numOfDevices);
    
/* We attempt to retrieve some information about the devices. */
    
for(int i = 0; i < numOfDevices; ++ i ) {
        
displayDeviceDetails( devices[i], CL_DEVICE_TYPE"CL_DEVICE_TYPE" );
        
displayDeviceDetails( devices[i], CL_DEVICE_NAME"CL_DEVICE_NAME" );
        
displayDeviceDetails( devices[i], CL_DEVICE_VENDOR"CL_DEVICE_VENDOR" );
        
displayDeviceDetails( devices[i], CL_DEVICE_VENDOR_ID"CL_DEVICE_VENDOR_ID" );
        
displayDeviceDetails( devices[i], CL_DEVICE_MAX_MEM_ALLOC_SIZE"CL_DEVICE_MAX_MEM_ALLOC_SIZE" );
        
displayDeviceDetails( devices[i], CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE"CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE" );
        
displayDeviceDetails( devices[i], CL_DEVICE_GLOBAL_MEM_SIZE"CL_DEVICE_GLOBAL_MEM_SIZE" );
        
displayDeviceDetails( devices[i], CL_DEVICE_MAX_COMPUTE_UNITS"CL_DEVICE_MAX_COMPUTE_UNITS" );
        
displayDeviceDetails( devices[i], CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS"CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS" );
        
displayDeviceDetails( devices[i], CL_DEVICE_MAX_WORK_ITEM_SIZES"CL_DEVICE_MAX_WORK_ITEM_SIZES" );
        
displayDeviceDetails( devices[i], CL_DEVICE_MAX_WORK_GROUP_SIZE"CL_DEVICE_MAX_WORK_GROUP_SIZE" );
    }
}
..

결과는 아래와 같다.

Number of OpenCL platforms found: 1
Number of detected OpenCL devices: 3
Running on CPU ........
Number of detected OpenCL devices: 3
CPU detected
Device name is Intel(R) Core(TM) i7-3615QM CPU @ 2.30GHz
Device vendor is Intel
VENDOR ID: 0xffffffff
Device max memory allocation: 4096 mega-bytes
Device global cacheline size: 6291456 bytes
Device global mem: 16384 mega-bytes
Maximum number of parallel compute units: 8
Maximum dimensions for global/local work-item IDs: 3
Maximum number of work-items in each dimension: ( 1024 1 1  )
Maximum number of work-items in a work-group: 1024

GPU detected
Device name is GeForce GT 650M
Device vendor is NVIDIA
VENDOR ID: 0x1022700
Device max memory allocation: 256 mega-bytes
Device global cacheline size: 0 bytes
Device global mem: 1024 mega-bytes
Maximum number of parallel compute units: 2
Maximum dimensions for global/local work-item IDs: 3
Maximum number of work-items in each dimension: ( 1024 1024 64  )
Maximum number of work-items in a work-group: 1024

GPU detected
Device name is HD Graphics 4000
Device vendor is Intel
VENDOR ID: 0x1024400
Device max memory allocation: 256 mega-bytes
Device global cacheline size: 0 bytes
Device global mem: 1024 mega-bytes
Maximum number of parallel compute units: 16
Maximum dimensions for global/local work-item IDs: 3
Maximum number of work-items in each dimension: ( 512 512 512  )
Maximum number of work-items in a work-group: 512

numOfKernels : 1
=> Kernel name: copy2Dfloat4 with arity: 2
=> About to create command queue and enqueue this kernel



처음에는 잘 몰랐는데 보면 work item의 차원 수와 각 차원의 갯수를 알수 있다.

우선 CPU에서 실행하게 되므로 {1024 , 1 , 1 }이 할당 될 수 있다.
앞서의 코드에서


                size_t globalThreads[2];
                globalThreads[
0]=1024;
                globalThreads[
1]=1024;
                
size_t localThreads[2];
                localThreads[
0] = 64;
                localThreads[
1] = 2;


이 되어 있으므로 localThread를 2를 할당해서 그런 것이다.
수정을 해서 진행하면
넘어간다.



…    

                cl_event evt;
                error = 
clEnqueueNDRangeKernel(cQ, 
                                               kernels[j],
                                               
2,
                                               
0,
                                               globalThreads,
                                               localThreads,
                                               
0

                                               NULL, &evt); 
                clWaitForEvents(1, &evt);

에서 걸리는데


clWaitForEvents 을 호출하면 대기하는 동안 EXC_BAD_ACCESS 가 발생한다.
즉 잘못된 메모리를 참조 하는 것으로 발생하는 것이다.

이것은 event list의 오류로 보기 보다는 kernel에서 메모리 참조 오류로 보아야 하는게 아닌가 생각이 든다.

처음부터 찬찬히 보면은


#define DATA_SIZE (128*128)      // for test runs, 데이터의 크기는 128 * 128개의 크기이다.
….

main()
{
...
   cl_float* h_in = (float*) mallocsizeof(cl_float4) * DATA_SIZE); // input to device
   
cl_float* h_out = (float*) mallocsizeof(cl_float4) * DATA_SIZE); // output from device
   
forint i = 0; i < DATA_SIZE; ++i) {
        h_in[i] = (
float)i;
   }
  h_in에 (128*128)개의 공간을 할당하고 그 메모리를 초기화 한다.
...


위에서 코드를 보면 cl_float와 float 그리고 cl_float4가 혼용되어 있다.
이를 아래와 같이 수정해둔다.

...
    cl_float* h_in = (float*) mallocsizeof(cl_float) * DATA_SIZE); // input to device
    
cl_float* h_out = (float*) mallocsizeof(cl_float) * DATA_SIZE); // output from device
   
forint i = 0; i < DATA_SIZE; ++i) {
        h_in[i] = (
cl_float)i;
   }




   // Search for a CPU/GPU device through the installed platforms
   
// Build a OpenCL program and do not run it.
   
for(cl_uint i = 0; i < numOfPlatforms; i++ ) {
…..
    /* For each device, create a buffer and partition that data among the devices for compute! */
    
cl_mem memInObj = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR
                                    
sizeof(cl_float4) * (DATA_SIZE), h_in, &error);
    
if(error != CL_SUCCESS) {
        
perror("Can't create an input buffer object");
        
exit(1);
    }


        int offset = 0;

        
for(int i = 0; i < numOfDevices; ++i, ++offset ) {

            /* Loop thru each kernel and execute on device */
            
// 커널의 갯수 만큼 반복한다.
        
for(cl_uint j = 0; j < numOfKernels; j++) {

            cl_mem memOutObj = clCreateBuffer(context, CL_MEM_WRITE_ONLY ,
                                              
sizeof(cl_float4) * (DATA_SIZE), NULL, &error);
                
            
if(error != CL_SUCCESS) {
                
perror("Can't create an output buffer object");
                
exit(1);
            }

            /* Let OpenCL know that the kernel is suppose to receive two arguments */
            error = 
clSetKernelArg(kernels[j], 0sizeof(cl_mem), &memInObj);
            
if (error != CL_SUCCESS) { 
                
perror("Unable to set buffer object in kernel");
                
exit(1);
            }
                
            error = 
clSetKernelArg(kernels[j], 1sizeof(cl_mem), &memOutObj);
            
if (error != CL_SUCCESS) { 
                
perror("Unable to set buffer object in kernel");
                
exit(1);
            }


                // 커널을 커맨드 큐에 넣는다.
            
/* Enqueue the kernel to the command queue */
                
size_t globalThreads[3]= { 1024 , 2 , 1 };
                size_t localThreads[3] = { 128 , 1 , 1 } ;
                cl_event  evt ;//= new cl_event[1];

            error = 
clEnqueueNDRangeKernel(cQ, 
                                               kernels[j],
                                               
2,
                                               
0,
                                               globalThreads,
                                               localThreads,
                                               
0
                                               
NULL, &evt);
                
                
clWaitForEvents(1, (const cl_event*)&evt);


            /* Enqueue the read-back from device to host */
            error = 
clEnqueueReadBuffer(cQ, memOutObj,
                                        
CL_TRUE,               // blocking read
                                        
0,         // write from the last offset
                                        (
DATA_SIZE)*sizeof(cl_float4),           // how much to copy
                                        h_out, 
0NULLNULL);

                /* Check the returned data */
            
if ( valuesOK(h_in, h_out, DATA_SIZE) ) {
                
printf("Check passed!\n");
            } 
else printf("Check failed!\n");

            
/* Release the resources */
            
clReleaseCommandQueue(cQ);
            
clReleaseMemObject(memOutObj);
     }


이제 실행 결과를 보면 아래와 같다.

=> Task has been enqueued successfully!
Checking data of size: 16384
to:0.000000, from:0.000000
to:1.000000, from:1.000000
to:2.000000, from:2.000000
to:3.000000, from:3.000000
to:1.000000, from:0.000000

..
위와 같이 제대로 데이터가 카피가 안된다.

OpenCL 함수는 아래와 같이 코딩되어 있다.

#define WIDTH 128
#define DATA_TYPE float4

/* 
  The following macros are convenience 'functions' 
  for striding across a 2-D array of coordinates (x,y)
  by a factor which happens to be the width of the block
  i.e. WIDTH
*/

#define A(x,y) A[(x)* WIDTH + (y)]
#define C(x,y) C[(x)* WIDTH + (y)]
__kernel void copy2Dfloat4(__global DATA_TYPE *A, __global DATA_TYPE *C)
{
    
int x = get_global_id(0);
    
int y = get_global_id(1);
    
// its like a vector load/store of 4 elements
    C(x,y) = A(x,y);

}


global id를 두가지를 받아 와서 동작하돌고 만들어져 있다.


Kernel과 work item에 대해서 관계를 정리하면 아래와 같다.

커널은 하나의 워크 아이템에 대해서 한번만 실행이 된다.
각각의 워크 아이템은 하나의 private 메모리 공간을 가지고 있다.

워크 아이템을 묶어서 work group으로 나타내어 진다.
워크 그룹은 local memory를 가진다.

전체 워크 아이템의 갯수는 global work size로 명시된다.
global과 constant 메모리는 모든 워크 아이템과 워크 그룹 사이에서 공유된다





위의 그림에서
하나의 사각형은 워크 아이템을 나타낸다.
그리고 그룹핑된 박스는 워크 그룹을 나타낸다.

OpenCL은 Dimension이라는 개념을 기반으로 움직인다. 이것은 work item을  인덱스를 이용해서 선언할 수 있다는 의미이다.
위의 그림에서 work group의 크기는  Sx = 4 , Sy = 4 가 된다.
얼마나 큰 dimension을 사용 할 지는 전적으로 프로그래머에게 달려 있다.
그러나 전체적으로 물리적인 한계 때문에 그룹당 워크 아이템이라던가 global index등의 한계는 존재하게 된다.

커널의 내에서는 핸재의 커널이 할당된 work item의 위치를 물어 볼 수 있다.

get_global_id(dim) 

으로 알아 낼 수 있다.

위의 함수는 
 get_local_size(dim)*get_group_id(dim) + get_local_id(dim).
과 동일한 결과를 가지고 온디ㅏ.

 get_local_size(dim) : group size 
 get_group_id(dim) : group position in dim
 get_local_id(dim)  : he position of a work item relative to the group

 그림으로 나타내면 아래와 같다.




연산을 어떻게 할 것인가를 나타내기 위해서는 
현재의 위치를 알고 현재의 위치에 맞는 크기를 가지고 와야 한다.

global size는 global thread로 코드에 표시가 되는데

1024개의 글로벌 Size가 있다면

128 * 128 / 1024 = 16 이 되어야 한다.

local 은 global을 local단위로 잘라내는 역활을 한다.

            /* Enqueue the kernel to the command queue */
                
size_t globalThreads[3]= { 256 , 64 , 1 };
//                globalThreads[0]=1024;
//                globalThreads[1]=1024;
                
                
size_t localThreads[3] = { 128 , 1 , 1 } ;
//                localThreads[0] = 128;
//                localThreads[1] = 128;


위와 같이 되어 있다면

256 x 64의 global size를 가지게 된다.
이것을 128 x 1으로 나누어서 한번에 계산하게 되는 것이다.
 이때 Kernel은 현재의 위치를 알기 위해서는
global_id(0)와 global_id(1)을 확인해야 한다.

OpenCL에서

#define A(x,y) A[(x)* WIDTH + (y)]

#define C(x,y) C[(x)* WIDTH + (y)]

..
    int x = get_global_id(0);
    
int y = get_global_id(1);
..

으로 되어 있다.

따라서 global thread가 0/1/2가 있는데 0은 OpenCL에서 Index(0)를 1은 Index(1)을 반영한다.(고 가정할 수 있다.)
그렇게 되면 위에서 처럼 256 x 64 이면 16384가 되므로 (1024,1,1)을 초과하여서 안되고
일단 1024,1,1로만 한정하여서 다시 해본다.

                size_t globalThreads[3]= { 1024 , 1 , 1 };

로 지정하면 global thread의 개수가 1024개까지 실행이 된다고 볼 수 있다.

여기서 local thread는 global thread를 나누어서 실행하는 것을 의미한다.
여기서는 128,1,1로 되어 있다면 
1024를 8개 Work Group으로 나누어서 실행한다고 볼 수 있다.
즉 한번에 128개를 동시에 실행하는 것이다.

여기까지는 문제가 없다.

OpenCL Code를 보면

..
#define A(x,y) A[(x)* WIDTH + (y)]
#define C(x,y) C[(x)* WIDTH + (y)]
__kernel void copy2Dfloat4(__global DATA_TYPE *A, __global DATA_TYPE *C)
{
    
int x = get_global_id(0);
    
int y = get_global_id(1);
    
//int y = 0;
    
// its like a vector load/store of 4 elements
    C(x,y) = A(x,y);

}


위와 같은데 
위치를 계산하는 x와 y를 보면

x * width + y의 값을 가진다.
그런데 x가 index 0를 y가 index 1을 가지고 오고 있다.
x축과 y축으로 본다면 index 0와 index 1이 바뀐 것으로 볼 수 있다.

따라서 다음과 같이 한다.

..
    int x = get_global_id(1);
    
int y = get_global_id(0);

으로 하고 컴파일하면


from:1019.000000, to:1019.000000
from:1020.000000, to:1020.000000
from:1021.000000, to:1021.000000
from:1022.000000, to:1022.000000
from:1023.000000, to:1023.000000
Check passed!
Program ended with exit code: 0

..
으로 테스트가 통과하고 있다.
결국 OpenCL에서 Index의 선택이 잘 못되어 있었다.


그런데
local thread를
                size_t localThreads[3] = { 512 1 , 1 } ; —> FAIL
                size_t localThreads[3] = { 256 1 , 1 } ; —> FAIL
                size_t localThreads[3] = { 128 , 1 , 1 } ; —> WORK
로 된다.

Local Thread의 갯수가 128일때 동작을 한다.

이런 것을 일일이 찾아서 만드는 것이 귀찮으면, 그리고 플랫폼에 따른 커널의 일관성을 유지하기가 쉽지 않을 것 같으므로
Local Thread를 NULL로 해두면 OpenCL Compiler가 알아서 할당하게 된다.
아래 코드로 해도 정상적으로 실행이 된다.

..
            error = clEnqueueNDRangeKernel(cQ, 
                                               kernels[j],
                                               
2,
                                               
0,
                                               globalThreads,
                                               
NULL,
                                               
0
                                               
NULL, &evt);



테스트를 위해서 다음과 같이 크기를 변경하였다.

#define DATA_SIZE (128*128)      // for test runs,

그랬더니 다음과 같은 오류가 발생하였다.

...
=> Task has been enqueued successfully!
Checking data of size: 16384
from:4096.000000, to:0.000000
from:4097.000000, to:0.000000
from:4098.000000, to:0.000000
from:4099.000000, to:0.000000
Check failed!
Program ended with exit code: 0



4096에서 오류가 발생한 것이다.

복사가 안된 것이다.


                size_t globalThreads[3]= { 512 , 1 , 1 };
..

로 바꾼뒤에 다시 실행하면
=> Task has been enqueued successfully!
Checking data of size: 16384
from:2048.000000, to:0.000000
from:2049.000000, to:0.000000
from:2050.000000, to:0.000000
from:2051.000000, to:0.000000
Check failed!

..
으로 2048에서 오류가 발생한다.
마찬가지로

...
                size_t globalThreads[3]= { 256 , 1 , 1 };
...

=> Task has been enqueued successfully!
Checking data of size: 16384
from:1024.000000, to:0.000000
from:1025.000000, to:0.000000
from:1026.000000, to:0.000000
from:1027.000000, to:0.000000



1024일때 4096
512일 경우 2048
256 일 경우 1024
이후에 오류가 발생한다.
즉 Global Size * 4에서 부터 오류가 발생하는 것이다.

체크를 하기 위햇 더 프린트를 해 보면 아래와 같다.

Checking data of size: 16384
from:1024.000000, to:0.000000
from:1025.000000, to:0.000000
from:1026.000000, to:0.000000
from:1027.000000, to:0.000000
from:0.000000, to:0.000000
from:0.000000, to:0.000000
Check failed!



값이 1027 이후에는 아예 없어졌다.
현상은 두가지 이다. 하나는 1023까지는 정확하게 가는데
1024부터는 안가는 현상이고
두번째는 1027 이후 1028부터는 아예 데이터가 없다는 점이다.

원래 부터 없었는지 확인을 해보기 위해서 아래와 같은 코드를 넣어 보자

                printf("Before start check data \n");
                
printf("  from:%f, to:%f\n", h_in[1024+0] ,h_out[1024+0]);
                
printf("  from:%f, to:%f\n", h_in[1024+1] ,h_out[1024+1]);
                
printf("  from:%f, to:%f\n", h_in[1024+2] ,h_out[1024+2]);
                
printf("  from:%f, to:%f\n", h_in[1024+3] ,h_out[1024+3]);
                
printf("  from:%f, to:%f\n", h_in[1024+4] ,h_out[1024+4]);
..

결과는 아래와 같다.

Before start check data 
  from:1024.000000, to:0.000000
  from:1025.000000, to:0.000000
  from:1026.000000, to:0.000000
  from:1027.000000, to:0.000000
  from:0.000000, to:0.000000


우선 1028부터는 아예 값이 안들어간 문제부터 확인해보자

    forint i = 0; i < DATA_SIZE; ++i) {
          …      
    }

위의 코드를 보면

데이터 크기까지만 카피 하도록 되어 있다.
그런데 데이터 크기의 타입이 float4이므로 실제로는 그 4배가 되어야 한다.

    forint i = 0; i < DATA_SIZE*4; ++i) {
        h_in[i] = (
float)(i);
    }



위와 같이 수정하여서 실행하면 다음과 같은 결과를 얻는다.


Checking data of size: 16384
from:1024.000000, to:0.000000
from:1025.000000, to:0.000000
from:1026.000000, to:0.000000
from:1027.000000, to:0.000000
from:1028.000000, to:0.000000
from:1029.000000, to:0.000000
Check failed!
Program ended with exit code: 0

...
이제 문제중 두번째 문제는 해결 되었다.
첫번째 문제인 1024 이후에는 복사를 하지 않는 문제를 해결 해 보자

아까 말한대로 global thread의 4배까지만 복사해주는 문제가 있다.
이는 Enqueue하는 것이 global thread까지만 한다고 가정 할 수 있지 않을까 ?

그런데 여기서 전체적으로 작업해야 할 크기가 1024이고 float4이므로 실제로는 4096까지 간다는 것은 이해 할 수 있다.
한번 이것을 늘려 보자

그러면 

                size_t globalThreads[3]= { (128 * 128 / 4 ) , 1 , 1 };


으로 하면 

Running GPU ........
numOfKernels : 1
=> Kernel name: copy2Dfloat4 with arity: 2
=> About to create command queue and enqueue this kernel...
[0000] Before start check data 
  from:1024.000000, to:1024.000000
  from:1025.000000, to:1025.000000
  from:1026.000000, to:1026.000000
  from:1027.000000, to:1027.000000
  from:1028.000000, to:1028.000000
=> Task has been enqueued successfully!
Checking data of size: 16384
Check passed!


정확하게 종료 하는 것을 알 수 있다.

따라서 global thread  값은 실제 수행해야 하는 work item의 갯수를 의미한다.

이를 local thread로 나누어서 수행하는 것이다.

이제 일반화 시키기 위해서 다음과 같이 한다.

..
                size_t globalThreads[3]= { (DATA_SIZE / 4 ) , 1 , 1 };
...

그래도 결과는 성공적으로 끝난다.
이번에는 2차원으로 해보자

..
               size_t globalThreads[3]= { (DATA_SIZE / 4 ) >> 1  , 1 << 1 , 1 };
..

from:12288.000000, to:0.000000
from:12289.000000, to:0.000000
from:12290.000000, to:0.000000
from:12291.000000, to:0.000000
from:12292.000000, to:0.000000
from:12293.000000, to:0.000000
Check failed!

위에서 오류가 발생하였음을 알 수 있다.


                size_t globalThreads[3]= { (DATA_SIZE / 4 ) >> 1  , 1 << 1 , 1 };

에 대한 OpenCL에서 Code를 반영하지 않아서 생긴 문제이다.

...
#define WIDTH (16384/8)



으로 하면 정상적으로 실행이 된다.

                size_t globalThreads[3]= { (DATA_SIZE / 4 ) >> 2  , 1 << 2 , 1 };

#define WIDTH (16384/16)

위와 같이 되어도 정상적으로 동작을 한다.

정리하면

size_t globalThreads[3]= { (DATA_SIZE )   , 1   1 };
#define WIDTH (16384)

size_t globalThreads[3]= { (DATA_SIZE /2 )   , 1   1 };
#define WIDTH (16384/2)

size_t globalThreads[3]= { (DATA_SIZE / 4 )  , 1 1 };
#define WIDTH (16384/4)

size_t globalThreads[3]= { (DATA_SIZE / 8 )  , 2 , 1 };
#define WIDTH (16384/8)

size_t globalThreads[3]= { (DATA_SIZE / 16 )   , 4 , 1 };
#define WIDTH (16384/16)

가 정상적으로 동작하는 케이스이다.

2차원의 값이 증가하는 경우 숫자가 증가하므로 1차원의 숫자가 /2가 된다.





개념은 위와 같은 모델이 된다.
잘 생각해보면 당연한 이야기 이다.





'Computer Vision' 카테고리의 다른 글

증강 현실 테스트  (0) 2014.10.04
OpenCL : Using work item to partition data (5)  (1) 2014.06.01
OpenCL::Query OpenCL kernel (4)  (0) 2014.05.02
OpenCL Test Program (3)  (0) 2014.03.31
OpenCL Test Program (2)  (0) 2014.03.18
OpenCL test program (1)  (0) 2014.03.15
Posted by GUNDAM_IM
Computer Vision2014.03.31 19:28

디바이스 별로 고유한 확장 기능을 부여할 수 있다.

이런 기능이 추가되어 있는지 확인을 한다.


추가되는 함수는 아래에 표시되어 있다.


void displayDeviceInfo( cl_platform_id id ,

                        cl_device_type dev_type)

......

        displayDeviceDetails( devices[i], CL_DEVICE_EXTENSIONS"CL_DEVICE_EXTENSIONS" );

.....



void displayDeviceDetails( cl_device_id id,

                            cl_device_info param_name,

                            const char * paramNameAsStr)

...

    

    switch(param_name)

    {

       ......

            break;

        case CL_DEVICE_EXTENSIONS : {

            // beware of buffer overflow; alternatively use the OpenCL C++ bindings

            char* extension_info[4096];

            error = clGetDeviceInfo( id, CL_DEVICE_EXTENSIONSsizeof(extension_info), 

extension_info, NULL);

            printf("\tSupported extensions: %s\n", extension_info);

        }break;

  }

....


확장에 대한 정보를 확인하기 위해서 CL_DEVICE_EXTENSONS를 확인하면 된다.

세부 정보에서 해당 정보를 프린트하도록 한다.


완성된 코드와 실행 결과는 아래와 같다.




//

//  main.cpp

//  TestOpenCL

//

//  Created by freegear on 2014. 2. 8..

//  Copyright (c) 2014 freegear. All rights reserved.

//


#include <iostream>


#include <stdio.h>

#include <stdlib.h>



#ifdef __APPLE__

#   include <OpenCL/opencl.h>

#else

#   include <CL/cl.h>

#endif


//Function proto type

void displayDeviceDetails(cl_device_id id ,

                          cl_device_info param_name,

                          const char * paramNameAsStr);


void displayDeviceInfo( cl_platform_id id ,

                        cl_device_type dev_type)

{

    /* OpenCL 1.1 device type */

    

    cl_int error = 0 ;

    cl_uint numOfDevices = 0 ;

    

    /* Determine how many devices are connected to your platform */

    error = clGetDeviceIDs(id, dev_type, 0, NULL, &numOfDevices);

    

    if( error != CL_SUCCESS){

        perror("Unable to obtain any OpenCL compliant device info");

        exit(1);

    }

    

    cl_device_id * devices = (cl_device_id*)alloca(sizeof(cl_device_id) * numOfDevices);

    

    /* Load the information about your devices into the variable 'devices' */

    

    error = clGetDeviceIDs(id, dev_type, numOfDevices, devices, NULL);

    

    if (error != CL_SUCCESS)

    {

        perror("Unable to obtain any OpenCL compliant device info");

        exit(1);

    }

    

    printf("Numnber of detected OpenCL devices %d\n", numOfDevices);

    

    /* we attempt to retrieve some information about the devices */

    for( int i = 0 ; i < numOfDevices ; ++i)

    {

        displayDeviceDetails(devices[i], CL_DEVICE_TYPE, "CL_DEVICE_TYPE");

        displayDeviceDetails(devices[i], CL_DEVICE_VENDOR_ID, "CL_DEVICE_VENDOR_ID");

        displayDeviceDetails(devices[i], CL_DEVICE_MAX_COMPUTE_UNITS, "CL_DEVICE_MAX_COMPUTE_UNITS");

        displayDeviceDetails(devices[i], CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, "CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS");

        displayDeviceDetails(devices[i], CL_DEVICE_MAX_WORK_ITEM_SIZES, "CL_DEVICE_MAX_WORK_ITEM_SIZES");

        displayDeviceDetails(devices[i], CL_DEVICE_MAX_WORK_GROUP_SIZE, "CL_DEVICE_MAX_WORK_GROUP_SIZE");

        displayDeviceDetails( devices[i], CL_DEVICE_EXTENSIONS, "CL_DEVICE_EXTENSIONS" );

    }

}


void displayDeviceDetails( cl_device_id id,

                            cl_device_info param_name,

                            const char * paramNameAsStr)

{

    cl_int error = 0 ;

    size_t paramSize = 0 ;

    

    error = clGetDeviceInfo( id ,param_name, 0 , NULL, &paramSize);

    if(error != CL_SUCCESS)

    {

        perror("Unable to obtain device info for param");

        return;

    }

    

    /* 

     The cl_device_info are preporcessor directives define in cl.h

     */

    

    switch(param_name)

    {

        case CL_DEVICE_TYPE :

        {

            cl_device_type * devType  = (cl_device_type*)alloca(sizeof(cl_device_type)*paramSize);

            error = clGetDeviceInfo(id, param_name, paramSize, devType, NULL);

            

            if(error != CL_SUCCESS)

            {

                perror("Unable to optain device info for param\n");

                return;

            }

            

            switch( *devType)

            {

                case CL_DEVICE_TYPE_CPU : printf("CPU Detected \n"); break;

                case CL_DEVICE_TYPE_GPU : printf("GPU Detected \n"); break;

                case CL_DEVICE_TYPE_ACCELERATOR : printf("Accelerator detected\n");break;

                case CL_DEVICE_TYPE_DEFAULT : printf("default detected \n"); break;

            }

            break;

        }

        case CL_DEVICE_VENDOR_ID :

        case CL_DEVICE_MAX_COMPUTE_UNITS :

        case CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS : {

            

                cl_uint* ret = (cl_uint*) alloca(sizeof(cl_uint) * paramSize);

                error = clGetDeviceInfo( id, param_name, paramSize, ret, NULL );

            

                if (error != CL_SUCCESS ) {

                    perror("Unable to obtain device info for param\n");

                    return;

                }

            

                switch (param_name) {

                    case CL_DEVICE_VENDOR_ID: printf("\tVENDOR ID: 0x%x\n", *ret); break;

                    case CL_DEVICE_MAX_COMPUTE_UNITS: printf("\tMaximum number of parallel compute units: %d\n", *ret); break;

                    case CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS: printf("\tMaximum dimensions for global/local work-item IDs: %d\n", *ret); break;

                }

            }

            break;

      

        case CL_DEVICE_MAX_WORK_ITEM_SIZES :

            {

                cl_uint maxWIDimensions;

                size_t* ret = (size_t*) alloca(sizeof(size_t) * paramSize);

                error = clGetDeviceInfo( id, param_name, paramSize, ret, NULL );

            

                error = clGetDeviceInfo( id, CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, sizeof(cl_uint), &maxWIDimensions, NULL );

                if (error != CL_SUCCESS ) {

                    perror("Unable to obtain device info for param\n");

                    return;

                }

                printf("\tMaximum number of work-items in each dimension: ( ");

                for(cl_int i =0; i < maxWIDimensions; ++i ) {

                    printf("%d ", (int)ret[i]);

                }

                printf(" )\n");

            }break;

        

        case CL_DEVICE_MAX_WORK_GROUP_SIZE :

            {

                size_t* ret = (size_t*) alloca(sizeof(size_t) * paramSize);

                error = clGetDeviceInfo( id, param_name, paramSize, ret, NULL );

                if (error != CL_SUCCESS ) {

                    perror("Unable to obtain device info for param\n");

                    return;

                }

                printf("\tMaximum number of work-items in a work-group: %d\n", (int)*ret);

            }

            break;

        case CL_DEVICE_NAME :

        case CL_DEVICE_VENDOR :

            {

                char data[48];

                error = clGetDeviceInfo( id, param_name, paramSize, data, NULL );

                if (error != CL_SUCCESS ) {

                    perror("Unable to obtain device name/vendor info for param\n");

                    return;

                }

                switch (param_name) {

                    case CL_DEVICE_NAME : printf("\tDevice name is %s\n", data);break;

                    case CL_DEVICE_VENDOR : printf("\tDevice vendor is %s\n", data);break;

                }

            }

            break;

        case CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE: {

            cl_uint* size = (cl_uint*) alloca(sizeof(cl_uint) * paramSize);

            error = clGetDeviceInfo( id, param_name, paramSize, size, NULL );

            if (error != CL_SUCCESS ) {

                perror("Unable to obtain device name/vendor info for param\n");

                return;

            }

            printf("\tDevice global cacheline size: %d bytes\n", (*size)); break;

        } break;

            

        case CL_DEVICE_GLOBAL_MEM_SIZE:

        case CL_DEVICE_MAX_MEM_ALLOC_SIZE:

            {

                cl_ulong* size = (cl_ulong*) alloca(sizeof(cl_ulong) * paramSize);

                error = clGetDeviceInfo( id, param_name, paramSize, size, NULL );

                if (error != CL_SUCCESS ) {

                    perror("Unable to obtain device name/vendor info for param\n");

                    return;

                }

                switch (param_name) {

                    case CL_DEVICE_GLOBAL_MEM_SIZE: printf("\tDevice global mem: %ld mega-bytes\n", (*size)>>20); break;

                    case CL_DEVICE_MAX_MEM_ALLOC_SIZE: printf("\tDevice max memory allocation: %ld mega-bytes\n", (*size)>>20); break;

                }

            }

            break;

        case CL_DEVICE_EXTENSIONS : {

            // beware of buffer overflow; alternatively use the OpenCL C++ bindings

            char* extension_info[4096];

            error = clGetDeviceInfo( id, CL_DEVICE_EXTENSIONS, sizeof(extension_info), extension_info, NULL);

            printf("\tSupported extensions: %s\n", extension_info);

        }break;

    }//end of switch(param_name

    

    

}// end of displayDeviceDetailes


void displayPlatformInfo(

                         cl_platform_id id ,

                         cl_platform_info param_name,

                         const char * paramNameAsStr

                         )

{

    

    cl_int error = 0 ;

    size_t paramSize = 0 ;

    

    

    error = clGetPlatformInfo(

                              id,

                              param_name,

                              0,

                              NULL,

                              &paramSize

                              );

    

    char * moreInfo = (char *) alloca(sizeof(char)*paramSize);

    

    error = clGetPlatformInfo(

                              id, /* The platform ID returned by clGetPlatformIDs or can be NULL. 

                                     If platform is NULL, the behavior is implementation-defined.*/

                              param_name, /* An enumeration constant that identifies the platform 

                                             information being queried. It can be one of the following values 

                                             as specified in the table below. */

                              paramSize,  /* Specifies the size in bytes of memory pointed to by param_value. 

                                             This size in bytes must be greater than or equal to size of return type

                                             specified in the table below. */

                              moreInfo,   /* A pointer to memory location where appropriate values for a given 

                                             param_value will be returned. Acceptable param_value values are listed in 

                                             the table below. If param_value is NULL, it is ignored.*/

                              NULL        /* Returns the actual size in bytes of data being queried by param_value. 

                                             If param_value_size_ret is NULL, it is ignored */

                              );

    

    if( error != CL_SUCCESS)

    {

        if ( error == CL_INVALID_PLATFORM)

            perror("Error CL_INVALID_PLATFORM");

        else if (error == CL_INVALID_VALUE)

            perror("Error CL_INVALID_VALUE");

        else

            perror("Unable to find any OpenCL platform information");

        return ;

    }

    

    printf("%s : %s\n" , paramNameAsStr, moreInfo);

    

} // end of displayPlatformInfo




int main(int argc, const char * argv[])

{


    

    // OpenCL 1.2 Data struction

    cl_platform_id * platforms ;

    

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

    */

    

    

    error = clGetPlatformIDs(0, /*

                                 The number of cl_platform_id entries that can be added to platforms.

                                 If platforms is not NULL, the num_entries must be greater than zero. 

                                 */

                             NULL, /* Returns a list of OpenCL platforms found. 

                                      The cl_platform_id values returned in platforms can be used to identify 

                                      a specific OpenCL platform. If platforms argument is NULL, 

                                      this argument is ignored. 

                                      The number of OpenCL platforms returned is the mininum of the value 

                                      specified by num_entries or the number of OpenCL platforms available. */

                             &numOfPlatforms

                                    /* Returns the number of OpenCL platforms available. 

                                       If num_platforms is NULL, this argument is ignored. */

                             );

    if ( error < 0)

    {

        perror("Unable to find any OpenCL platforms");

        exit(1);

    }

    

    // allocate memory for the number of installed platforms

    // alloca(....) occupies some stack space but is

    // automatically freed on return

    

    platforms = (cl_platform_id*) alloca(sizeof(cl_platform_id)*numOfPlatforms);

    

    printf("Number of OpenCL platforms found : %d\n", numOfPlatforms);

    

    

    error = clGetPlatformIDs(numOfPlatforms, /*

                                 The number of cl_platform_id entries that can be added to platforms.

                                 If platforms is not NULL, the num_entries must be greater than zero.

                                 */

                             platforms, /* Returns a list of OpenCL platforms found.

                                    The cl_platform_id values returned in platforms can be used to identify

                                    a specific OpenCL platform. If platforms argument is NULL,

                                    this argument is ignored.

                                    The number of OpenCL platforms returned is the mininum of the value

                                    specified by num_entries or the number of OpenCL platforms available. */

                             NULL

                             /* Returns the number of OpenCL platforms available.

                              If num_platforms is NULL, this argument is ignored. */

                             );

    if ( error < 0)

    {

        perror("Unable to find any OpenCL platforms");

        exit(1);

    }

    // We invoke the API 'clPlatformInfo' twice for each

    // parameter we are trying to extract

    // and we use the return value to create temporary data

    // structure (on the stack) to store

    // the returned information ot the second invocation.

    

    for( cl_uint i = 0 ; i < numOfPlatforms ; ++i)

    {

        displayPlatformInfo( platforms[i],  CL_PLATFORM_PROFILE     , "CL_PLATFORM_PROFILE");

        displayPlatformInfo( platforms[i],  CL_PLATFORM_VERSION     , "CL_PLATFORM_VERSION");

        displayPlatformInfo( platforms[i],  CL_PLATFORM_NAME        , "CL_PLATFORM_NAME");

        displayPlatformInfo( platforms[i],  CL_PLATFORM_VENDOR      , "CL_PLATFORM_VENDOR");

        displayPlatformInfo( platforms[i],  CL_PLATFORM_EXTENSIONS  , "CL_PLATFORM_EXTENSIONS");

        

        // Assume that we don't know how many devices are OpenCL compliant, we locate everything !

        displayDeviceInfo( platforms[i], CL_DEVICE_TYPE_ALL );

    }

    

    return 0;

}




실행 결과


Number of OpenCL platforms found : 1

CL_PLATFORM_PROFILE : FULL_PROFILE

CL_PLATFORM_VERSION : OpenCL 1.2 (Aug 24 2013 21:03:27)

CL_PLATFORM_NAME : Apple

CL_PLATFORM_VENDOR : Apple

CL_PLATFORM_EXTENSIONS : cl_APPLE_SetMemObjectDestructor cl_APPLE_ContextLoggingFunctions cl_APPLE_clut cl_APPLE_query_kernel_names cl_APPLE_gl_sharing cl_khr_gl_event


Numnber of detected OpenCL devices 3


CPU Detected 

VENDOR ID: 0xffffffff

Maximum number of parallel compute units: 8

Maximum dimensions for global/local work-item IDs: 3

Maximum number of work-items in each dimension: ( 1024 1 1  )

Maximum number of work-items in a work-group: 1024

Supported extensions: cl_APPLE_SetMemObjectDestructor cl_APPLE_ContextLoggingFunctions cl_APPLE_clut cl_APPLE_query_kernel_names cl_APPLE_gl_sharing cl_khr_gl_event cl_khr_fp64 cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics cl_khr_byte_addressable_store cl_khr_int64_base_atomics cl_khr_int64_extended_atomics cl_khr_3d_image_writes cl_khr_image2d_from_buffer cl_APPLE_fp64_basic_ops cl_APPLE_fixed_alpha_channel_orders cl_APPLE_biased_fixed_point_image_formats cl_APPLE_command_queue_priority


GPU Detected 

VENDOR ID: 0x1022700

Maximum number of parallel compute units: 2

Maximum dimensions for global/local work-item IDs: 3

Maximum number of work-items in each dimension: ( 1024 1024 64  )

Maximum number of work-items in a work-group: 1024

Supported extensions: cl_APPLE_SetMemObjectDestructor cl_APPLE_ContextLoggingFunctions cl_APPLE_clut cl_APPLE_query_kernel_names cl_APPLE_gl_sharing cl_khr_gl_event cl_khr_byte_addressable_store cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics cl_APPLE_fp64_basic_ops cl_khr_fp64 cl_khr_3d_image_writes cl_khr_depth_images cl_khr_gl_depth_images cl_khr_gl_msaa_sharing cl_khr_image2d_from_buffer 


GPU Detected 

VENDOR ID: 0x1024400

Maximum number of parallel compute units: 16

Maximum dimensions for global/local work-item IDs: 3

Maximum number of work-items in each dimension: ( 512 512 512  )

Maximum number of work-items in a work-group: 512

Supported extensions: cl_APPLE_SetMemObjectDestructor cl_APPLE_ContextLoggingFunctions cl_APPLE_clut cl_APPLE_query_kernel_names cl_APPLE_gl_sharing cl_khr_gl_event cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics cl_khr_byte_addressable_store cl_khr_image2d_from_buffer cl_khr_gl_depth_images cl_khr_depth_images 

Program ended with exit code: 0


위에서 나온 용어에 대해서 다음페이지에서 설명한다.

'Computer Vision' 카테고리의 다른 글

OpenCL : Using work item to partition data (5)  (1) 2014.06.01
OpenCL::Query OpenCL kernel (4)  (0) 2014.05.02
OpenCL Test Program (3)  (0) 2014.03.31
OpenCL Test Program (2)  (0) 2014.03.18
OpenCL test program (1)  (0) 2014.03.15
OpenCL Architecture  (0) 2014.03.15
Posted by GUNDAM_IM
TAG openCL
Computer Vision2014.03.15 18:03

OpenCL test program

책의 예제를 넣어서 실행


책에서 오류가 두군데가 있었다.

(1) ifdef에서 APPLE platform을 확인하는 매크로 정의 오류


== WAS == 

 #ifdef APPLE

 == IS ==

#ifdef __APPLE__


(2) OpenCL의 원칙은 OpenCL에 정보의 크기를 물어서 확인하고 그만큼 공간을 만든뒤에 그 정보를 받아들이는 것인데

     예제에서는 정보를 물어서 공간을 만드는 데 까지는 있지만, 그 이후에 정보를 다시 받아들이는 함수를 호출 하지 않았다.


....

main(...)

{

 ...

  먼저 필요한 공간의 크기를 확인

  error = clGetPlatformIDs(0,    

                             NULL,  

                             &numOfPlatforms

                           );

 if ( error < 0)

  ...


   공간을 확보 하고


   이제 확보된 공간으로 값을 받아온다. 


    error = clGetPlatformIDs(numOfPlatforms,  <== 이 함수를 호출 하는 것을 추가함.

                          platforms,  

                          NULL

                          );

  ...


}




전체 코드는 아래와 같다.


//

//  main.cpp

//  TestOpenCL

//

//  Created by freegear on 2014. 2. 8..

//  Copyright (c) 2014 freegear. All rights reserved.

//


#include <iostream>


#include <stdio.h>

#include <stdlib.h>



#ifdef __APPLE__

#include <OpenCL/opencl.h>

#else

#include <CL/cl.h>

#endif



void displayPlatformInfo(

                         cl_platform_id id ,

                         cl_platform_info param_name,

                         const char * paramNameAsStr

                         )

{

    

    cl_int error = 0 ;

    size_t paramSize = 0 ;

    

    

    error = clGetPlatformInfo(

                              id,

                              param_name,

                              0,

                              NULL,

                              &paramSize

                              );

    

    char * moreInfo = (char *) alloca(sizeof(char)*paramSize);

    

    error = clGetPlatformInfo(

                              id, /* The platform ID returned by clGetPlatformIDs or can be NULL. 

                                     If platform is NULL, the behavior is implementation-defined.*/

                              param_name, /* An enumeration constant that identifies the platform 

                                             information being queried. It can be one of the following values 

                                             as specified in the table below. */

                              paramSize,  /* Specifies the size in bytes of memory pointed to by param_value. 

                                             This size in bytes must be greater than or equal to size of return type

                                             specified in the table below. */

                              moreInfo,   /* A pointer to memory location where appropriate values for a given 

                                             param_value will be returned. Acceptable param_value values are listed in 

                                             the table below. If param_value is NULL, it is ignored.*/

                              NULL        /* Returns the actual size in bytes of data being queried by param_value. 

                                             If param_value_size_ret is NULL, it is ignored */

                              );

    

    if( error != CL_SUCCESS)

    {

        if ( error == CL_INVALID_PLATFORM)

            perror("Error CL_INVALID_PLATFORM");

        else if (error == CL_INVALID_VALUE)

            perror("Error CL_INVALID_VALUE");

        else

            perror("Unable to find any OpenCL platform information");

        return ;

    }

    

    printf("%s : %s\n" , paramNameAsStr, moreInfo);

    

} // end of displayPlatformInfo




int main(int argc, const char * argv[])

{


    

    // OpenCL 1.2 Data struction

    cl_platform_id * platforms ;

    

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

    */

    

    

    error = clGetPlatformIDs(0, /*

                                 The number of cl_platform_id entries that can be added to platforms.

                                 If platforms is not NULL, the num_entries must be greater than zero. 

                                 */

                             NULL, /* Returns a list of OpenCL platforms found. 

                                      The cl_platform_id values returned in platforms can be used to identify 

                                      a specific OpenCL platform. If platforms argument is NULL, 

                                      this argument is ignored. 

                                      The number of OpenCL platforms returned is the mininum of the value 

                                      specified by num_entries or the number of OpenCL platforms available. */

                             &numOfPlatforms

                                    /* Returns the number of OpenCL platforms available. 

                                       If num_platforms is NULL, this argument is ignored. */

                             );

    if ( error < 0)

    {

        perror("Unable to find any OpenCL platforms");

        exit(1);

    }

    

    // allocate memory for the number of installed platforms

    // alloca(....) occupies some stack space but is

    // automatically freed on return

    

    platforms = (cl_platform_id*) alloca(sizeof(cl_platform_id)*numOfPlatforms);

    

    printf("Number of OpenCL platforms found : %d\n", numOfPlatforms);

    

    

    error = clGetPlatformIDs(numOfPlatforms, /*

                                 The number of cl_platform_id entries that can be added to platforms.

                                 If platforms is not NULL, the num_entries must be greater than zero.

                                 */

                             platforms, /* Returns a list of OpenCL platforms found.

                                    The cl_platform_id values returned in platforms can be used to identify

                                    a specific OpenCL platform. If platforms argument is NULL,

                                    this argument is ignored.

                                    The number of OpenCL platforms returned is the mininum of the value

                                    specified by num_entries or the number of OpenCL platforms available. */

                             NULL

                             /* Returns the number of OpenCL platforms available.

                              If num_platforms is NULL, this argument is ignored. */

                             );

    if ( error < 0)

    {

        perror("Unable to find any OpenCL platforms");

        exit(1);

    }

    // We invoke the API 'clPlatformInfo' twice for each

    // parameter we are trying to extract

    // and we use the return value to create temporary data

    // structure (on the stack) to store

    // the returned information ot the second invocation.

    

    for( cl_uint i = 0 ; i < numOfPlatforms ; ++i)

    {

        displayPlatformInfo( platforms[i],  CL_PLATFORM_PROFILE     , "CL_PLATFORM_PROFILE");

        displayPlatformInfo( platforms[i],  CL_PLATFORM_VERSION     , "CL_PLATFORM_VERSION");

        displayPlatformInfo( platforms[i],  CL_PLATFORM_NAME        , "CL_PLATFORM_NAME");

        displayPlatformInfo( platforms[i],  CL_PLATFORM_VENDOR      , "CL_PLATFORM_VENDOR");

        displayPlatformInfo( platforms[i],  CL_PLATFORM_EXTENSIONS  , "CL_PLATFORM_EXTENSIONS");

        

    }

    

    return 0;

}




실행 결과는 아래와 같다.


Number of OpenCL platforms found : 1

CL_PLATFORM_PROFILE : FULL_PROFILE

CL_PLATFORM_VERSION : OpenCL 1.2 (Aug 24 2013 21:03:27)

CL_PLATFORM_NAME : Apple

CL_PLATFORM_VENDOR : Apple

CL_PLATFORM_EXTENSIONS : cl_APPLE_SetMemObjectDestructor cl_APPLE_ContextLoggingFunctions cl_APPLE_clut cl_APPLE_query_kernel_names cl_APPLE_gl_sharing cl_khr_gl_event


'Computer Vision' 카테고리의 다른 글

OpenCL Test Program (3)  (0) 2014.03.31
OpenCL Test Program (2)  (0) 2014.03.18
OpenCL test program (1)  (0) 2014.03.15
OpenCL Architecture  (0) 2014.03.15
Optical Flow에 의한 영상 정보 분석  (0) 2013.05.28
OpenCV를 이용한 Face Detection  (0) 2013.05.19
Posted by GUNDAM_IM
Books2014.02.19 02:04


OpenCL에 대한 이슈가 있어서 책을 몇권 사서 보는 중입니다.

그중에서 이 책은 Kindle로도 책을 보고 싶어서 산 책입니다.


눈으로만 읽을 때는 잘 몰랐는지만, 책의 예제 코드를 인터넷에서 다운 받아서 컴파일을 하게 되면서 알게 된것이,

저자가 이 책을 쓸때 코드를 (어쩌면 한번도) 컴파일을 해보지 않았구나 하는 생각이 들었습니다.


기본적으로 컴파일이 안되는 코드들이 많았고,

그리고 코드가 컴파일이 되어도 NULL 포인터에 값을 마구 대입하는 등 코드에 이런 저런 버그가 꽤 많이 있습니다.


그래도 오기로라도 돌려볼 요령으로 버그 잡으면서 책을 따라가면서 진행하였습니다.

나중에 가만히 생각하다보니, 저자의 의도가 그러면서 배우라고 하는게 아닌가 싶을 정도로 코드에 대해서 무성의한 부분이 많이 있습니다. 


그리고,

설명의 순서를 먼저 데이터 구조/메모리 구조를 설명해주고 Host Program을  설명해주고 그 이후에 OpenCL을 설명해주면 이해가  쉬울텐데 설명 순서를 완전히 반대로 하고 있어서 앞에서 설명이 없이 무작정 OpenCL 코드만을 보면 도대체 왜 이렇게 하는지 한참 시간을 투자해야 이해할 수 있는 순서로 책이 설명이 되어 있습니다.  한참 코드를 보면서 데이터 구조를 손으로 그리면서 따라갔더니 챕터의 뒤에서 그림이 나오는 것을 보고 헐... 하면서.. 당황하였습니다.


그래서 이 책은 그닥 비추입니다


반대 이 책의 좋은 점은 이런 저런 깊은 이해나 코드를 돌려보지 않겠다는 마음을 가지고 눈팅으로만 책을 빨리 읽어 가면서 개념을 잡겠다면 의외로 설명만으로 보면 좋은 부분도 있습니다. 그리고 (당연히) 책도 빨리 읽을 수 있습니다. 책에서는 코드를 다 설명하는 것이 아니라 코드에서 설명해야될 포인트 부분만 설명하고 있고 그 외의 나머지 코드  부분은 굳이 그 설명을 이해하기 위해서 알 필요가 있는것이 아니기 때문입니다. 물론 코드를 돌려본다면 다른 문제가 되곘지만요.


책은 앞단에서는 OpenCL의 기본적인 설명을 하여주고

뒷단에서는 OpenCL을 이용한 프로그램 작성에 대해서 Histogram이나 Sobel Filter등을 OpenCL을 이용해서 만들어가는 것을 설명해주고 있습니다.


간단하게 개념만 잡겠다는 컨셉으로 빠르게 읽어간다면 그나마 볼 수 있는 책이고

반대로 이것 저것 생각해보면서 읽어가겠다면 비추입니다.


끝으로 책을 살 때에는 처음 나왔기 때문에 아마존에서 별다른 코멘트가 없었는데 

최근에 다시 들어가니 그닥 좋은 코멘트가 없더군요


Posted by GUNDAM_IM
TAG openCL