본문 바로가기
Program Language/OpenCL

OpenCL 포팅

by Leo 리오 2012. 10. 5.
반응형

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번 돈다.
결과는 되게 느리다.



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 ;

}


끝.











반응형

'Program Language > OpenCL' 카테고리의 다른 글

Intel opencl platform analyzer  (0) 2012.10.15
DeviceQuery  (0) 2012.10.08
Release  (0) 2012.09.24
clCreateBuffer clEnqueueWriteBuffer clSetKernelArg clEnqueueReadBuffer  (0) 2012.09.24
OpenCL 외부파일 쉽게 컴파일하기  (0) 2012.09.21

댓글