OpenCL 포팅

1. C로 구현.

int j = 0;

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

{

j =  j * 3;

a[i] = j ;

}

이 코드는 a[i]에 3^i을 넣는 코드이다.

물론 j =  j * 3; 대신 a[i] = 3^i;를 써도 되지만 j * 3이 3^i 보다 light(빠름)다면 이렇게 쓸 것이다.

2. OpenMP로 구현.

우선 OpenMP로 포팅해서 data dependency를 없애주자.

int j = 0;

#pragma omp parallel

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

{

j =  j * 3;

a[i] = j ;

}

3. OpenMP Data Dependency 제거

이렇게 하고 인텔 composer같은 툴을 돌리면 

j끼리 디펜던시가 있다고 알려준다.

바꾸자.

#pragma omp parallel

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

{

int j =  3 ^ i;

a[i] = j ;

}

이제 j는 for안의 local variable이다.

디펜던시제거 하였다.

이렇게 하면 4코어이면 빠르면 3.7배 정도 성능 향상이 있을 것이다.

4. OpenCL 구현 - Workitem 1개

이제 Workitem 1개의 OpenCL로 구현해보자.

(편의상 pseudo 코드)

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

{

write_buffer(a[], clmem); //인풋 넣기

executeKernel(1); //Workitem 1개로 실행

read_buffer(clmem, a[]); //결과값 빼오기.

}

__kernel execute(a[], i) //i값도 인풋으로 넣자.

{

int j =  3 ^ i;

a[i] = j ;

}

여기서 write_buffer와 read_buffer를 한번의 iteration마다 전체 배열a를 넣다 빼준다.

계산은 i번째 하나만 해준다.

그걸 for로 256번 돈다.

결과는 되게 느리다.

이렇게 하는 이유는

이건 되게 간단하지만 예이지만, 실전에서 커널 코드가 복잡해지면

코드의 correctness를 확신하기 힘들다.

Step by Step으로 가자.

OpenCL은 더럽다. 

Host구현도 더럽고, Kernel 문법도 더럽고, 디버깅도 더럽다.

5. OpenCL 구현 - Workitem Max

앞의 구현된 코드의 결과가 맞다면 이제 Workitem 갯수를 늘려서 GPU의 많은 코어를 활용해보자.

//for( int i=0 ; i<256 ; i ++)

{

write_buffer(a[], clmem); //인풋 넣기

executeKernel(256); //Workitem 256개로 실행

read_buffer(clmem, a[]); //결과값 빼오기.

}

__kernel execute(a[]) 

{

int i =  get_global_id(0); // i 대신  get_global_id(0)를 사용하자.

int j =  3 ^ i;

a[i] = j ;

}

끝.