GPGPU 컴퓨팅의 개념에 대해서 설명하도록 하겠다.


더 심화된 내용은 GPU나 디바이스에 따라서 다르겠지만, 


기본적인 전체적인 개념, 한번씩 머리에 담아두고 있으면 좋을 만한 개념을 설명하도록 하겠다.






GPU 컴퓨팅에 가장 중요한 개념 두가지는


1. 막강한 컴퓨팅 능력.


2. 엄청느린 메모리 접근


3. CPU-GPU와의 통신 


이다.



1. 막강한 컴퓨팅 능력.


  GPU는 기본적으로 간단한 프로세서가 수백 수 천개가 들어있다.

그렇기 때문에 GPU는 CPU에 비해 계산을 엄청 빠르게 할 수 있다.


그러나, GPU를 이용한 컴퓨팅은 GPU의 몇가지 특성 때문에 성능이 낮아 질 수 있다.


a. SIMT 모델 (branch divergence)


많은 코어를 갖고도 저전력, 소형화(?)를 위해서 

GPU는 SIMT(singl instruction multiple threads)모델을 사용한다.

SIMT 모델의 단점은 한번에 하나의 instruction만 실행 가능 하다는 것으로,


if(a>0)  b++;

else b--;


와 같은 코드에서 덧샘 한번, 다음 뺄셈 한번이 실행된다. 

따라서 이런 분기문이 있을 때마다 2배의 실행 시간이 소요된다.


이런 분기를 줄이고 같은 warp에서는 같은 분기로 빠지게하는 최적화(optimization)가 꼭 필요하다.



2. 엄청느린 메모리 접근


이 부분 또한 GPU컴퓨팅의 성능에 큰 영향을 미친다. GPU가 global메모리(dram)에 접근하기 위해선 수백 cycle이 필요하다.

이 큰 지연(latency)를 숨기기(hiding) 위해 warp을 context switching 하는 기술을 사용한다.

또한, GPU메모리엔 hierarchy가 있으며,

cache또한 있다.


a. warp context switching


앞 서 말했듯이 큰 지연을 갖은 메모리 접근 명령어(instruction)을 실행할 때 

그 수백 cycle을 기다리는것이 아니라 메모리 명령이 끝날때 까지 다른 warp을 가져와 다른 warp을 실행한다.

기본적으로 context switching을 위해선 register를 저장하고 바꾸는 기능이 필요하다.

그러나 GPU에선 많은 register를 가지고 있어 context switching 할때 overhead가 들지 않는다.

따라서 GPU는 context switching을 자유롭게 가능하다.


그렇기 위해선 warp을 최대한 넣어줘야하는데 그 warp의 수는

GPU마다 다르겠지만, 

1) register의 수 ; 하나의 kernel이 사용하는 register수를 GPU(SM, streamming multiprocessor)가 갖고 있는 register의 수에 나눈 값 만큼 warp의 할당이 가능하다.

2) shared 메모리의 양 ; 이또한 하나의 kernel이 사용하는 shared memory의 양과 GPU의 메모리의 양에 따라 계산된다.

3) GPU가 갖을 수 있는 최대 warp의 수 ; 이건 GPU를 좋은걸 사용하는 수 밖에 없다.

이 세가지 요인에 의해 제한된다. 위 세가지를 모두 만족하는 최대의 숫자가 warp의 수가 된다.





b. GPU 메모리 hierarchy

GPU엔 크게 다음과같은 메모리 hierarchy구조가 있다.

register  - private memory - shared memory - global memory

왼쪽으로 갈 수록 thread하나만 사용가능한 메모리이며, 오른쪽으로 갈수록 모든 thread가 공유가능한 메모리이다.

또한, 왼쪽은 빠르고, 오른쪽으로 갈 수록 느리다.

thread간 공유 하는 메모리를 최소화 하여 최대한 왼쪽의 메모리를 사용하는게 좋다.

하지만, a에서 말했듯 최대한 많은 warp을 할당하기 위해 조절을 해주어야한다.




c. cache

GPU는 한번에 여러개의 데이터를 처리 하기 때문에 보통 한번에 많은 데이터를 가져온다.

예를 들어 warp의 크기가 32라면 보통 4Byte*32의 배수 만큼 한번에 가져 올 수 있다.

그러니깐 128Byte의 메모리를 가져오는게 100cycle이 걸린다면, 129Byte의 메모리를 가져오는데에는 200Cycle이 걸린다.


캐시를 잘 사용하기 위해선 메모리를 가져오는 양과 align을 조심 해야 한다.



3. CPU-GPU와의 통신.


CPU-GPU같의 데이터 전송은 PCI를 이용해 이루워지는데


이 또한 굉장히 느리다.


이 데이터 전송량을 최소화하고,


memory transfer overlapping이란 방법을 사용한 최적화가 필요하다.





저작자 표시 비영리 변경 금지
신고

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

GPGPU Computing의 기본 개념.  (0) 2013.04.19
OpenCL 에러  (0) 2013.01.21
global memory replay overhead  (0) 2012.12.21
[CUDA] occupancy  (0) 2012.11.07
NVIDIA clEnqueueReadBuffer non-blocking bug(?)  (0) 2012.11.02
Intel opencl platform analyzer  (0) 2012.10.15
Posted by Leo 리오 트랙백 0 : 댓글 0

OpenCL 에러

2013.01.21 14:20 from Program Language/OpenCL

void boo()

{

if(get_local_id(2)!=0) return;

foo();

barrier(CLK_LOCAL_MEM_FENCE);

}


작동 안한다.

알아서 작동할 줄 알았는데....


void boo()

{

if(get_local_id(2)!=0) 

foo();

barrier(CLK_LOCAL_MEM_FENCE);

}


이렇게 하자.






atomic 함수는 shared-reg에 써도 느리다.


왠만하면 피하자!


branch divergence를 유발한다!







저작자 표시 비영리 변경 금지
신고

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

GPGPU Computing의 기본 개념.  (0) 2013.04.19
OpenCL 에러  (0) 2013.01.21
global memory replay overhead  (0) 2012.12.21
[CUDA] occupancy  (0) 2012.11.07
NVIDIA clEnqueueReadBuffer non-blocking bug(?)  (0) 2012.11.02
Intel opencl platform analyzer  (0) 2012.10.15
TAG IT, OpenCL
Posted by Leo 리오 트랙백 0 : 댓글 0

쓸모없는 barrier를 제거 하였더니 

global memory replay overhead가 줄었다.

성능이 크게 올라갔다.



//barrier(CLK_GLOBAL_MEM_FENCE);

//mem_fence(CLK_GLOBAL_MEM_FENCE);






저작자 표시 비영리 변경 금지
신고

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

GPGPU Computing의 기본 개념.  (0) 2013.04.19
OpenCL 에러  (0) 2013.01.21
global memory replay overhead  (0) 2012.12.21
[CUDA] occupancy  (0) 2012.11.07
NVIDIA clEnqueueReadBuffer non-blocking bug(?)  (0) 2012.11.02
Intel opencl platform analyzer  (0) 2012.10.15
Posted by Leo 리오 트랙백 0 : 댓글 0

들어가는말,


GPU는 캐시가 없거나 아주 작다. 따라서 Global 메모리를 접근 빈도수가 높은데, latency또한 높기 때문에 GPU의 계산 속도는 엄청나게 느릴 것이다. 이것을 보안하기 위한 방법이 메모리에 접근하는 동안 실행되는 warp을 교체(context switch)해버리는 방법을 사용하게 된다.

또한, 앞뒤 instruction간의 dependency가 있을 경우도 stall을 해야하는데 이것 또한 빠른 context switch로 상쇄(hide)시킬 수 있다. (or ILP를 높임)


Fermi의 경우,

GMEM latency: 400-800 cycles

Arithmetic latency: 18-22cycles


GPU는 SM에서 context switch가 자유롭다. (overhead = 0)

[이유를 간단히 말하자면, Warp의 묶음인 block이 SM에 할당 되게 되면, 각각의 Thread들은 자신만의 자원(register, shared memory)들을 가지고 있다. 따라서 말이 context switch이지 변하는건 없다.]


결론은 SM에 Thread(block)를 많이 넣으면 넣으면 넣을 수록 성능이 좋다는것이다. (block개수)/(최대 block개수) = occupancy라고 한다.

그럼 어떻게 occupancy를 높일 수 있을까??



Geforce 550 TI


CL_DEVICE_COMPUTE_CAPABILITY_NV: 2.1
NUMBER OF MULTIPROCESSORS: 4

NUMBER OF CUDA CORES: 192
CL_DEVICE_REGISTERS_PER_BLOCK_NV: 32768
CL_DEVICE_WARP_SIZE_NV: 32
CL_DEVICE_GPU_OVERLAP_NV: CL_TRUE
CL_DEVICE_KERNEL_EXEC_TIMEOUT_NV: CL_FALSE
CL_DEVICE_INTEGRATED_MEMORY_NV: CL_FALSE
CL_DEVICE_PREFERRED_VECTOR_WIDTH_<t> CHAR 1, SHORT 1, INT 1, LONG 1, FLOAT 1, DOUBLE 1



Physical Limits for GPU Compute Capability: 2.0


Threads per Warp 32

Warps per Multiprocessor 48

Threads per Multiprocessor 1536

Thread Blocks per Multiprocessor 8

Total # of 32-bit registers per Multiprocessor 32768

Register allocation unit size 64

Register allocation granularity warp

Registers per Thread 63

Shared Memory per Multiprocessor (bytes) 49152

Shared Memory Allocation unit size 128

Warp allocation granularity 2

Maximum Thread Block Size 1024




GPU Architecture





여러개의 SP(Stream Processor)가 합쳐서 SM(Stream MultiProcessor)를 이룬다.

SM에는 Registers, Shared Memory를 공유한다.


위의 스펙을 보면

지포스 550ti는 SM이 4개이고,

전체 SP는 192개이고,

하나의 SM당 SP는 192/4=48개가 있다.

Register는 32768개.

Shared Memory는 49152bytes있다.


이 SM에 warp의 모음인 block이 매핑 되게 된다. 

global_work_item이 32*27 이고,

local_work_item이 32*3 이면,

한 block당 thread는 32*3개, 3개의 warp이 매핑 되고,

27/3=9개의 block들이 생기게 된다.

그리고 Thread Blocks per Multiprocessor:8 이기 때문에

SM0에 8개의 block이 매핑 되고

SM1에 1개의 block이 매핑 될 수 있다.

그럼, Warps per Multiprocessor:48이고, SM에 3*8warp이 매핑 됬으므로, occupancy는 24warp/48=50%가 되게 된다.


그럼 50%를 100%로 만들 수 있는 방법은?


우선 SM당 warp수는 

[warps/block] * [block/SM] 으로 계산 가능하다.


[warps/block]은 local_work_item/32이다.


[block/SM]은

SM당 block수를 제한하는 요소는 세가지가 있다.

Thread Blocks per Multiprocessor 8

Total # of 32-bit registers per Multiprocessor 32768

Shared Memory per Multiprocessor (bytes) 49152

이 세가지 중에 가장 작은 값이 개수가 된다.


1. block

block의 수는 constant다.

maxblock/SM = 8, maxwarps/SM = 48 이기 때문에

warps/block이 최소 6(32*6threads/block)이면 된다.


2. register

커널을 컴파일 해보면 커널(thread)당 사용하는 register의 수를 알 수 있다.

만약 32개의 레지스터를 사용한다면 SM당 최대 thread의 수는 32768/20=1024개 이고 이것은 1024/32=32warps/SM이다.


3. shared memory

마찬가지로 컴파일 해보면 커널당 (thread)당 사용하는 shared memory의 양을 알 수있다.


1,2,3을 모두 만족하는 즉, 최소값이 SM당 warps수가 된다.





참고.

http://developer.download.nvidia.com/CUDA/training/NVIDIA_GPU_Computing_Webinars_Further_CUDA_Optimization.pdf

http://www.nvidia.com/content/PDF/sc_2010/CUDA_Tutorial/SC10_Fundamental_Optimizations.pdf

http://developer.download.nvidia.com/CUDA/training/cuda_webinars_WarpsAndOccupancy.pdf

http://nvidia.fullviewmedia.com/gtc2010/0922-a5-2238.html

http://www.cs.berkeley.edu/~volkov/volkov10-GTC.pdf























저작자 표시 비영리 변경 금지
신고

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

OpenCL 에러  (0) 2013.01.21
global memory replay overhead  (0) 2012.12.21
[CUDA] occupancy  (0) 2012.11.07
NVIDIA clEnqueueReadBuffer non-blocking bug(?)  (0) 2012.11.02
Intel opencl platform analyzer  (0) 2012.10.15
DeviceQuery  (0) 2012.10.08
Posted by Leo 리오 트랙백 0 : 댓글 0

clEnqueueReadBuffer를 CL_FALSE로 불러도 계속 read될때 까지 기다린다.

Intel CPU OpenCL은 non-blocking이 잘 된다.

AMD는 안해봤는데 아마 잘 될듯하다.


이걸로 1.5일을 헤매었다.


CUDA 4.2.1

OpenCL 1.1을 사용 하였다.


NVIDIA는 끝까지 OpenCL을 숨기려는듯하다.

OPENCL 페이지도 없어지고,

1.2도 공개안하고, VisualProfiler도 CUDA 5.0에선 작동을 안한다.


째튼 

NVIDIA 예제코드도 확인결과 다 blocking 된다. overrap예제만 빼고.. 

뭐가 다른가

lib? inc? 다 바꿔봤지만 그대로..

context? device? queue? 다 똑같다.



결론은

asynchronous pinned transfers 쓰자.


보통

cl_mem deviceResult;

char* hostResult = malloc();

clEnqueueReadBuffer(hostResult, deviceResult, CL_FALSE);

이렇게 들어가는데 clEnqueueReadBuffer가 계속 기다린(blocking)다.


cl_mem deviceResult;

char* hostResult = clEnqueueMapBuffer(deviceResult,CL_TRUE);

clEnqueueReadBuffer(hostResult, deviceResult, CL_FALSE);

이런식 pinned로 하니깐 non-blocking이 가능했다...


pinned로 하면 특정메모리에 밖에 읽을 수가 없어 구현이 좀 복잡해지긴 하지만,

속도(bandwidth)도 빨라지고 하니 이렇게 구현하는게 좋을 것이다.



저작자 표시 비영리 변경 금지
신고

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

global memory replay overhead  (0) 2012.12.21
[CUDA] occupancy  (0) 2012.11.07
NVIDIA clEnqueueReadBuffer non-blocking bug(?)  (0) 2012.11.02
Intel opencl platform analyzer  (0) 2012.10.15
DeviceQuery  (0) 2012.10.08
OpenCL 포팅  (0) 2012.10.05
Posted by Leo 리오 트랙백 0 : 댓글 0

http://software.intel.com/sites/landingpage/opencl/user-guide/index.htm

http://software.intel.com/sites/landingpage/opencl/user-guide/Using_the_Intel(R)_Graphics_Performance_Analyzers_(Intel(R)_GPA)_Platform_Analyzer.htm



AMD OpenCL컴파일러가 병맛같기 때문에 Intel용 툴을 써봤다.





1. GPA TOOL 다운


http://software.intel.com/en-us/vcsource/tools/intel-gpa







2. 프로파일 파일 생성


2-1. 환경 변수 설정


Environment Variable

  Default Value

CL_GPA_CONFIG_ENABLE_API_TRACING

  True







CL_CONFIG_USE_GPA  = True






2-2. Intel GAP Monitor 실행





2-3. 커맨드 설정



Command Line에 실행 Full Path와 argument씀.

Working Folder 설정.




2-4. 프로파일 설정.


Manage Profiles...



a. Enable trigger (단축키로 실행시킬 수 있지만, 귀찮)

b. Apllication Time, >= 0 ( 시작부터)

c. Frame and Trace Capture (Frame은 생성이 안된다;)

d. Retrigger Delay : 1secs (1초마다 생성, 어플이 끝날 때 한번 하고 싶었지만 모르겠다...)




2-5. Run


아래 내문서 폴더에 생성됬다고 뜬다.


1초마다 생성되어?

여러개가 생성되었다.




<span style="background-color: rgb(124, 67, 177); color: rgb(255, 255, 255); ">에러가 날 때</span >





3. Intel opencl Platform Analyzer


맨 마지막 파일을 실행시키면








4. 분석하기















PS. 근데 위에 환경변수들을 켜놓으면

성능이 엄청 나빠진다.

- 해결











저작자 표시 비영리 변경 금지
신고

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

[CUDA] occupancy  (0) 2012.11.07
NVIDIA clEnqueueReadBuffer non-blocking bug(?)  (0) 2012.11.02
Intel opencl platform analyzer  (0) 2012.10.15
DeviceQuery  (0) 2012.10.08
OpenCL 포팅  (0) 2012.10.05
Release  (0) 2012.09.24
Posted by Leo 리오 트랙백 0 : 댓글 0

DeviceQuery

2012.10.08 19:01 from Program Language/OpenCL

NVIDA OpenCL SDK에 들어있는 


DeviceQuery 프로그램 



oclDeviceQuery32.exe


oclDeviceQuery64.exe



결과들


GTX 550 Ti


Radeon HD 7800


Intel i5-2500












저작자 표시 비영리 변경 금지
신고

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

NVIDIA clEnqueueReadBuffer non-blocking bug(?)  (0) 2012.11.02
Intel opencl platform analyzer  (0) 2012.10.15
DeviceQuery  (0) 2012.10.08
OpenCL 포팅  (0) 2012.10.05
Release  (0) 2012.09.24
clCreateBuffer clEnqueueWriteBuffer clSetKernelArg clEnqueueReadBuffer  (0) 2012.09.24
Posted by Leo 리오 트랙백 0 : 댓글 0

OpenCL 포팅

2012.10.05 19:21 from Program Language/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번 돈다.
결과는 되게 느리다.

더보기



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
OpenCL 포팅  (0) 2012.10.05
Release  (0) 2012.09.24
clCreateBuffer clEnqueueWriteBuffer clSetKernelArg clEnqueueReadBuffer  (0) 2012.09.24
OpenCL 외부파일 쉽게 컴파일하기  (0) 2012.09.21
TAG OpenCL
Posted by Leo 리오 트랙백 0 : 댓글 0

Release

2012.09.24 20:20 from Program Language/OpenCL

clReleaseContext(context);

clReleaseCommandQueue(commandq);

clReleaseProgram(program);

clReleaseKernel(kernel);

저작자 표시 비영리 변경 금지
신고

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

DeviceQuery  (0) 2012.10.08
OpenCL 포팅  (0) 2012.10.05
Release  (0) 2012.09.24
clCreateBuffer clEnqueueWriteBuffer clSetKernelArg clEnqueueReadBuffer  (0) 2012.09.24
OpenCL 외부파일 쉽게 컴파일하기  (0) 2012.09.21
[OCL] #include cannot open source file "......cl"  (0) 2012.09.21
Posted by Leo 리오 트랙백 0 : 댓글 0



http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/clCreateBuffer.html

cl_mem clCreateBuffer (cl_context contextcl_mem_flags flagssize_t sizevoid *host_ptrcl_int *errcode_ret)

flags
CL_MEM_READ_WRITE
CL_MEM_WRITE_ONLY
CL_MEM_READ_ONLY
CL_MEM_USE_HOST_PTR
CL_MEM_ALLOC_HOST_PTR
CL_MEM_COPY_HOST_PTR


cl_mem clEnqueueWriteBuffer (cl_command_queue command_queuecl_mem buffercl_bool blocking_writesize_t offsetsize_t cb, const void *ptrcl_uint num_events_in_wait_list, const cl_event *event_wait_listcl_event *event)


cl_int clSetKernelArg (cl_kernel kernelcl_uint arg_indexsize_t arg_size, const void *arg_value)


cl_int clEnqueueReadBuffer (cl_command_queue command_queuecl_mem buffercl_bool blocking_readsize_t offsetsize_t cbvoid *ptrcl_uint num_events_in_wait_list, const cl_event *event_wait_listcl_event *event)



NOTICE! cl_mem 자체가 포인터이다.
(typedef struct _cl_mem *            cl_mem;)


호스트 

cl_mem clmem_input = clCreateBuffer(context, CL_MEM_WRITE_ONLY, SIZE, 0, err);
clEnqueueWriteBuffer(commandq, clmem_input, CL_TRUE, 0, SIZE, INPUT_ptr, 0, 0);
clSetKernelArg(kernel, arg++, sizeof(cl_mem), clmem_input);

cl_mem clmem_result = clCreateBuffer(context, CL_MEM_READ_ONLY, SIZE, 0, err);
clSetKernelArg(kernel, arg++, sizeof(cl_mem),  clmem_result );

커널
__kernel Goo(global unsigned char* clmem1, local unsigned char* clmem2)
{

}

호스트
clEnqueueReadBuffer(commandq, clmem_result, CL_TRUE, 0, SIZE, OUTPUT_ptr, 0,0,0);

clReleaseMemObject(clmem_input);
clReleaseMemObject(clmem_result);




저작자 표시 비영리 변경 금지
신고
Posted by Leo 리오 트랙백 0 : 댓글 0

OpenCL에서 cl파일을 컴파일하기 위해선


파일을 string으로 읽어와야한다.


왜 이렇게 만들었지;


#include를 이용하면 간단하게 할 수 있다.


 

        char source2[512];

        char* ptr_src2 = source2;

        sprintf(source2, "#include \"%s\"\n",filename);

       

        program = clCreateProgramWithSource(context, 1, &ptr_src2, 0, &err);

 

        char build_option2[512] = "-I . ";

        err = clBuildProgram(program, 0, NULL, build_option2, NULL, NULL);

 


중요한점은 build_option에 "-I . "를 추가해줘야 include가 제대로 된다.


더보기


저작자 표시 비영리 변경 금지
신고
Posted by Leo 리오 트랙백 0 : 댓글 0

INTEL OPENCL 기준


#include 해서 같은 폴더에 넣어 놓더라도  

cannot open source file "......cl" 이라는 에러만 나고

include가 안된다.

Default로 어디를 참조하는지는 모르겠지만,

Build 할때 "-I ." 옵션을 주면 된다.

clBuildProgram(program, 0, NULL, build_option, NULL, NULL);


저작자 표시 비영리 변경 금지
신고
Posted by Leo 리오 트랙백 0 : 댓글 0

http://software.intel.com/sites/landingpage/opencl/user-guide/index.htm




1. Intel OpenCL SDK 받기

http://software.intel.com/en-us/articles/vcsource-tools/

http://software.intel.com/en-us/articles/vcsource-tools-opencl-sdk/





2. visual studio project setting

  1. Open the project property pages by selecting Project > Properties.
  2. In the C/C++ > General property page, under Additional Include Directories, enter the full path to the directory where the OpenCL* header files are located:

    $(INTELOCLSDKROOT)\include.




3.  In the Linker > General property page, under Additional Library Directories, enter the full path to the directory where the OpenCL* run-time import library file is located. For example, for 32-bit application:

$(INTELOCLSDKROOT)\lib\x86




In the Linker > Input property page, under Additional Dependencies, enter the name of the OpenCL* ICD import library fileOpenCL.lib.







저작자 표시 비영리 변경 금지
신고

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

OpenCL 외부파일 쉽게 컴파일하기  (0) 2012.09.21
[OCL] #include cannot open source file "......cl"  (0) 2012.09.21
Intel OpenCL visual studio 설정  (0) 2012.08.08
amd opencl __constant  (0) 2012.06.20
standard functions in OpenCL  (0) 2012.06.04
clBuildProgram  (1) 2012.06.02
Posted by Leo 리오 트랙백 0 : 댓글 0

AMD에서 __kernel argument 를 __constant 로 못받는듯하다;


HOST::

 //create buffer read only

cl_mem cltmp = clCreateBuffer(context, CL_MEM_READ_ONLY, 2048, 0, &err);

clEnqueueWriteBuffer(... cltmp...);

cl_mem clarg = cltmp; //copy buffer


clSetKernelArg(... clarg...);

clEnqueueTask(...);


DEVICE::

__kernel void k(... __const mem...){}



이렇게 하면 Intel 에선 잘 작동하지만.


amd에선 argument가 cl_mem주소가 아닌 0으로 넘어온다..


왜그런진 모르겠지만.



저작자 표시 비영리 변경 금지
신고

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

[OCL] #include cannot open source file "......cl"  (0) 2012.09.21
Intel OpenCL visual studio 설정  (0) 2012.08.08
amd opencl __constant  (0) 2012.06.20
standard functions in OpenCL  (0) 2012.06.04
clBuildProgram  (1) 2012.06.02
OpenCL Address Spaces  (0) 2012.06.02
Posted by Leo 리오 트랙백 0 : 댓글 0

memset in OpenCL


커널에서도 memset이된다!


memset(dst_ptr, value, size);




printf in OpenCL


Intel 에선 그냥 쓰면된다.


AMD에선 

#pragma OPENCL EXTENSION cl_amd_printf : enable




Calling clFinish on a command queue flushes all pending output by printf in previously enqueued and completed commands to the implementation-defined output stream.
근데 잘안되네;









저작자 표시 비영리 변경 금지
신고

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

Intel OpenCL visual studio 설정  (0) 2012.08.08
amd opencl __constant  (0) 2012.06.20
standard functions in OpenCL  (0) 2012.06.04
clBuildProgram  (1) 2012.06.02
OpenCL Address Spaces  (0) 2012.06.02
clFinish, clFlush, block  (0) 2012.05.26
Posted by Leo 리오 트랙백 0 : 댓글 0

clBuildProgram

2012.06.02 13:41 from Program Language/OpenCL

clBuildProgram

Builds (compiles and links) a program executable from the program source or binary.

cl_int clBuildProgram (cl_program program,
 cl_uint num_devices,
 const cl_device_id *device_list,
 const char *options,
 void (*pfn_notify)(cl_program, void *user_data),
 void *user_data)

Parameters

program

The program object

device_list

A pointer to a list of devices that are in program. If device_list is NULL value, the program executable is built for all devices associated withprogram for which a source or binary has been loaded. If device_list is a non-NULL value, the program executable is built for devices specified in this list for which a source or binary has been loaded.

num_devices

The number of devices listed in device_list.

options

A pointer to a string that describes the build options to be used for building the program executable. The list of supported options is described in "Build Options" below.

pfn_notify

A function pointer to a notification routine. The notification routine is a callback function that an application can register and which will be called when the program executable has been built (successfully or unsuccessfully). If pfn_notify is not NULL, clBuildProgram does not need to wait for the build to complete and can return immediately. If pfn_notify is NULL, clBuildProgram does not return until the build has completed. This callback function may be called asynchronously by the OpenCL implementation. It is the application's responsibility to ensure that the callback function is thread-safe.

user_data

Passed as an argument when pfn_notify is called. user_data can be NULL.

Notes

OpenCL allows program executables to be built using the source or the binary.

The build options are categorized as pre-processor options, options for math intrinsics, options that control optimization and miscellaneous options. This specification defines a standard set of options that must be supported by an OpenCL compiler when building program executables online or offline. These may be extended by a set of vendor- or platform-specific options.

Preprocessor Options

These options control the OpenCL preprocessor which is run on each program source before actual compilation. -D options are processed in the order they are given in the options argument to clBuildProgram.
-D name

Predefine name as a macro, with definition 1.

-D name=definition

The contents of definition are tokenized and processed as if they appeared during translation phase three in a `#define' directive. In particular, the definition will be truncated by embedded newline characters.

-I dir

Add the directory dir to the list of directories to be searched for header files.


Math Intrinsics Options

These options control compiler behavior regarding floating-point arithmetic. These options trade off between speed and correctness.
-cl-single-precision-constant

Treat double precision floating-point constant as single precision constant.

-cl-denorms-are-zero

This option controls how single precision and double precision denormalized numbers are handled. If specified as a build option, the single precision denormalized numbers may be flushed to zero and if the optional extension for double precision is supported, double precision denormalized numbers may also be flushed to zero. This is intended to be a performance hint and the OpenCL compiler can choose not to flush denorms to zero if the device supports single precision (or double precision) denormalized numbers.

This option is ignored for single precision numbers if the device does not support single precision denormalized numbers i.e. CL_FP_DENORM bit is not set in CL_DEVICE_SINGLE_FP_CONFIG.

This option is ignored for double precision numbers if the device does not support double precision or if it does support double precison but CL_FP_DENORM bit is not set in CL_DEVICE_DOUBLE_FP_CONFIG.

This flag only applies for scalar and vector single precision floating-point variables and computations on these floating-point variables inside a program. It does not apply to reading from or writing to image objects.


Optimization Options

These options control various sorts of optimizations. Turning on optimization flags makes the compiler attempt to improve the performance and/or code size at the expense of compilation time and possibly the ability to debug the program.
-cl-opt-disable

This option disables all optimizations. The default is optimizations are enabled.

-cl-strict-aliasing

This option allows the compiler to assume the strictest aliasing rules.

The following options control compiler behavior regarding floating-point arithmetic. These options trade off between performance and correctness and must be specifically enabled. These options are not turned on by default since it can result in incorrect output for programs which depend on an exact implementation of IEEE 754 rules/specifications for math functions.

-cl-mad-enable

Allow a * b + c to be replaced by a mad. The mad computes a * b + c with reduced accuracy. For example, some OpenCL devices implement madas truncate the result of a * b before adding it to c.

-cl-no-signed-zeros

Allow optimizations for floating-point arithmetic that ignore the signedness of zero. IEEE 754 arithmetic specifies the behavior of distinct +0.0and -0.0 values, which then prohibits simplification of expressions such as x+0.0 or 0.0*x (even with -clfinite-math only). This option implies that the sign of a zero result isn't significant.

-cl-unsafe-math-optimizations

Allow optimizations for floating-point arithmetic that (a) assume that arguments and results are valid, (b) may violate IEEE 754 standard and (c) may violate the OpenCL numerical compliance requirements as defined in section 7.4 for single-precision floating-point, section 9.3.9 for double-precision floating-point, and edge case behavior in section 7.5. This option includes the -cl-no-signed-zeros and -cl-mad-enable options.

-cl-finite-math-only

Allow optimizations for floating-point arithmetic that assume that arguments and results are not NaNs or ±∞. This option may violate the OpenCL numerical compliance requirements defined in in section 7.4 for single-precision floating-point, section 9.3.9 for double-precision floating-point, and edge case behavior in section 7.5.

-cl-fast-relaxed-math

Sets the optimization options -cl-finite-math-only and -cl-unsafe-math-optimizations. This allows optimizations for floating-point arithmetic that may violate the IEEE 754 standard and the OpenCL numerical compliance requirements defined in the specification in section 7.4 for single-precision floating-point, section 9.3.9 for double-precision floating-point, and edge case behavior in section 7.5. This option causes the preprocessor macro __FAST_RELAXED_MATH__ to be defined in the OpenCL program.


Options to Request or Suppress Warnings

Warnings are diagnostic messages that report constructions which are not inherently erroneous but which are risky or suggest there may have been an error. The following languageindependent options do not enable specific warnings but control the kinds of diagnostics produced by the OpenCL compiler.
-w

Inhibit all warning messages.

-Werror

Make all warnings into errors.

Errors

clBuildProgram returns CL_SUCCESS if the function is executed successfully. Otherwise, it returns one of the following errors:

  • CL_INVALID_PROGRAM if program is not a valid program object.
  • CL_INVALID_VALUE if device_list is NULL and num_devices is greater than zero, or if device_list is not NULL and num_devices is zero.
  • CL_INVALID_VALUE if pfn_notify is NULL but user_data is not NULL.
  • CL_INVALID_DEVICE if OpenCL devices listed in device_list are not in the list of devices associated with program.
  • CL_INVALID_BINARY if program is created with clCreateWithProgramWithBinary and devices listed in device_list do not have a valid program binary loaded.
  • CL_INVALID_BUILD_OPTIONS if the build options specified by options are invalid.
  • CL_INVALID_OPERATION if the build of a program executable for any of the devices listed in device_list by a previous call to clBuildProgram for program has not completed.
  • CL_COMPILER_NOT_AVAILABLE if program is created with clCreateProgramWithSource and a compiler is not available i.e.CL_DEVICE_COMPILER_AVAILABLE specified in the table of OpenCL Device Queries for clGetDeviceInfo is set to CL_FALSE.
  • CL_BUILD_PROGRAM_FAILURE if there is a failure to build the program executable. This error will be returned if clBuildProgram does not return until the build has completed.
  • CL_INVALID_OPERATION if there are kernel objects attached to program.
  • CL_OUT_OF_HOST_MEMORY if there is a failure to allocate resources required by the OpenCL implementation on the host.


저작자 표시 비영리 변경 금지
신고

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

amd opencl __constant  (0) 2012.06.20
standard functions in OpenCL  (0) 2012.06.04
clBuildProgram  (1) 2012.06.02
OpenCL Address Spaces  (0) 2012.06.02
clFinish, clFlush, block  (0) 2012.05.26
clEnqueueCopyBufferRect vs clEnqueueWriteBufferRect  (0) 2012.05.22
Posted by Leo 리오 트랙백 0 : 댓글 1


Address Sapce 종류 


__global 

Global memory


Stores data for the entire device and can be read from and writen to 


디바이스내의 모든 커널들이 공유할 수 있다.

매우 느리다.

사용을 최소화 하자.

GPU에선 GPU DRAM에 위치한다.


__constant 

Constant memory 


Similar to global memory, but is read-only


__global메모리와 같다.

다만, 쓰기가 불가능하다.

플랫폼마다 다를 수 있는데, constant메모리 캐시가 있어 global메모리보다는 빠른듯 하다.

읽기만 가능하기 때문에 동기화가 필요없다(불가능 하다.)



__local 

Local memory


Stores data for the work-items in a work-goup


work-goup간에 공유가 가능하다. 다른 work-group끼리는 공유가 불가능 하다.

global 메모리보다 훨씬 빠르다. 

kernel간의 동기화가 필요하다면, 같은 work-group으로 묶고, shared 메모리를 이용해서 동기화 하자.

__kernel void foo_kernel()

{

__local char larr[N];

}

각 work-group마다 N byte씩 할당된다. 총 N*work-group수 만큼 할당된다.

GPU에선 SM의 shared memory에 위치한다.



__private 

Private memory 


Stores data for an individual work-item


기본적인 변수이다.

각각 커널마다 가지고 있다.

다른 커널과는 동기화가 불가능 하다.

GPU에선 register가 담당하고 모자란다면, shared memory로 spill 된다.




Pointer 포인터의 사용(의미) 


__global float *f;

f가 global메모리에 저장되있다는 뜻이 아니라,

f가 global 메모리의 주소를 가르킨다는 뜻이다.


__global float *f;

__global uint *x = 5;

f = (global float*) x; 

는 가능하다.


__global float *f;

__local uint *l = 5;

f = (global float*) l; 

는 불가능하다.

f와 l이 의 address space가 다르기 때문이다.



그럼 포인터를 __private로 사용하려면??

__global uint *x = 5;

__private int offset = 3;

*(x + offset) = 100;

이런식으로 해야 할 듯하다.


Addtion. 이 부분은 다시 생각해봐야할듯, x라는 값자체는 __private로 가지고 있을지도..




전역변수(global variable)무조건 __constant로만 선언해야 한다.

다시 말하면 '전역 변수'는 사용 불가.

(흠, 포인터 전역변수는 __constant Address Space만 가르 킬 수있다.라고 생각 했지만 안된다.)

다른 Address Space도 컴파일러가 충분히 지원해 줄 수 있을꺼 같은데 현재로썬 안된다.


__constant int c; 만 가능.


__constant int* ptr;

__global int g;

__local int l;

__private int p;

void f(){ c++; }

다다다다다 불가능.


전역변수는 미리 정의된 상수를 제외하면 사용할 수 없다고 보면된다.



__global 메모리


__kernel void foo_kerenl()

{

__global char g_arr[10];

}

error: automatic variable qualified with an address space

라는 에러가 뜬다.

global메모리는 kernel에서 선언이 불가능하다. 

아마 커널마다 따로 실행이 되서 인듯하다.


clCreateBuffer를 통해 생성된 메모리 포인터를 받음으로써 

global메모리를 사용가능하다.

__local 메모리나 __private 메모리를 사용하자.









저작자 표시 비영리 변경 금지
신고

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

amd opencl __constant  (0) 2012.06.20
standard functions in OpenCL  (0) 2012.06.04
clBuildProgram  (1) 2012.06.02
OpenCL Address Spaces  (0) 2012.06.02
clFinish, clFlush, block  (0) 2012.05.26
clEnqueueCopyBufferRect vs clEnqueueWriteBufferRect  (0) 2012.05.22
Posted by Leo 리오 트랙백 0 : 댓글 0


cl_int clFinish (cl_command_queue command_queue)


Blocks until all previously queued OpenCL commands in a command-queue are issued to the associated device and have completed.


Notes

Blocks until all previously queued OpenCL commands in command_queue are issued to the associated device and have completed.

clFinish does not return until all queued commands in command_queue have been processed and completed. clFinish is also a synchronization point.


커맨드큐의 커맨드들이 다 실행되기 전까지 리턴하지 않는다. (블러킹)

동기화포인트로 사용.



cl_int clFlush (cl_command_queue command_queue)


Issues all previously queued OpenCL commands in a command-queue to the device associated with the command-queue.


Notes

Issues all previously queued OpenCL commands in command_queue to the device associated with command_queue.

clFlush only guarantees that all queued commands to command_queue get issued to the appropriate device. There is no guarantee that they will be complete after clFlush returns.



커맨드큐의 커맨드들을 이슈 시킨다.




아마,



Finish는 write, execute, read같은 명령어를 다 실행시킬때 까지 기다리는거고,


Flush는 커맨드큐의 커맨드들을 디바이스에 넣을때 까지 기다림.



예를 들면,


enqueueWrite()에서


Finish는 모든 데이터를 다 쓸 때 까지 기다리는 거고,


Flush는 쓰라는 명령어를 디바이스에 전달할 때 까지 기다리는듯.



여기에 따르면,

https://www.khronos.org/message_boards/viewtopic.php?f=28&t=4335



clEnqueueReadBuffer()같은 함수들은 알아서 블러킹을 하기 때문에 Finish나 Flush를 쓸 필요가 없다.(blocking CL_TRUE를 쓸때)


아래글을 보면 정확히는 

블러킹함수는 clflush를 대체하는듯.


Any blocking commands queued in a command-queue such as clEnqueueRead{Image|Buffer}

with blocking_read set to CL_TRUE, clEnqueueWrite{Image|Buffer} with blocking_write set 

to CL_TRUE, clEnqueueMap{Buffer|Image} with blocking_map set to CL_TRUE or 

clWaitForEvents perform an implicit flush of the command-queue.

To use event objects that refer to commands enqueued in a command-queue as event objects to 

wait on by commands enqueued in a different command-queue, the application must call a 

clFlush or any blocking commands that perform an implicit flush of the command-queue where 

the commands that refer to these event objects are enqueued.



이벤트 오브젝트에 관한 내용은


http://www.khronos.org/registry/cl/specs/opencl-1.x-latest.pdf#page=116


clFlush는 이벤트 오브젝트를 사용할 때만 효용성이 있는듯하다.




블러킹 커맨드가 finish가 아닌 flush를 포함하고 있다는 내용은


커맨드큐가 OOO(out of order)로 정의 되어 있을 때를 말하는 듯하다.


커맨드큐가 IO(inorder)로 정의 되었다면, 블러킹 커맨드는 finish까지 포함할 것이다.






고민 해봤는데 특히 주목해야할 점은


clFinish가 동기화 포인트가 된다는 점이다.


IO일때는 마지막에 블러킹 함수를 큐하면 동기화가 이루어지겠지만,


OOO일때는 clFinish를 넣어야만 그전 함수가 모두 끝났다고 할 수있겠다.







뭘원하는지, 뭘대답했는지 모르게 횡설수설했네;

결국 쓰다보면 당연한것을...















저작자 표시 비영리 변경 금지
신고

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

amd opencl __constant  (0) 2012.06.20
standard functions in OpenCL  (0) 2012.06.04
clBuildProgram  (1) 2012.06.02
OpenCL Address Spaces  (0) 2012.06.02
clFinish, clFlush, block  (0) 2012.05.26
clEnqueueCopyBufferRect vs clEnqueueWriteBufferRect  (0) 2012.05.22
Posted by Leo 리오 트랙백 0 : 댓글 0

cl_int clEnqueueCopyBufferRect (cl_command_queue command_queue,
 cl_mem src_buffer,
 cl_mem dst_buffer,
 const size_t src_origin[3],
 const size_t dst_origin[3],
 const size_t region[3],
 size_t src_row_pitch,
 size_t src_slice_pitch,
 size_t dst_row_pitch,
 size_t dst_slice_pitch,
 cl_uint num_events_in_wait_list,
 const cl_event *event_wait_list,
 cl_event *event)

cl_int clEnqueueWriteBufferRect (cl_command_queue command_queue,
 cl_mem buffer,
 cl_bool blocking_write,
 const size_t buffer_origin[3],
 const size_t host_origin[3],
 const size_t region[3],
 size_t buffer_row_pitch,
 size_t buffer_slice_pitch,
 size_t host_row_pitch,
 size_t host_slice_pitch,
 void *ptr,
 cl_uint num_events_in_wait_list,
 const cl_event *event_wait_list,
 cl_event *event)


모지?

했는데 완전 다른 함수.


copy는 src_buffer(cl_mem)에서 dst_buffer(cl_mem)으로 (device) 복사 하는거고, 


writehost의 ptr(void*)을 device의 buffer(cl_mem)으로 쓰기(transfer).




이게 아마 2point 배열에서만 사용가능한거 같다.


int[][]는 사용불가이고,

int**에서 사용가능 할듯.


int[][]는 clEnqueueR,WBuffer나 clEnqueueR,WImage를 사용해야할듯..




Heterogeneous Computing with OpenCL (Paperback)
외국도서
저자 : Benedict Gaster, David R. Kaeli, Lee Howes, Perhaad Mistry
출판 : Morgan Kaufmann Publishers 2011.08.31
상세보기


Opencl Programming Guide (Paperback / 1st Ed.)
외국도서
저자 : Aaftab Munshi
출판 : Pearson 2011.07.23
상세보기







저작자 표시 비영리 변경 금지
신고

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

amd opencl __constant  (0) 2012.06.20
standard functions in OpenCL  (0) 2012.06.04
clBuildProgram  (1) 2012.06.02
OpenCL Address Spaces  (0) 2012.06.02
clFinish, clFlush, block  (0) 2012.05.26
clEnqueueCopyBufferRect vs clEnqueueWriteBufferRect  (0) 2012.05.22
Posted by Leo 리오 트랙백 0 : 댓글 0