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

function String foo()

{

return this.num.toString();

}



ldarg.0

ldfld num

call toString()

ret


이렇게 하면 될줄 알았는데


계속 죽는다. 이게 몰까??


ldfld num

>> ldflda num


이렇게 했어야한다. 


왜그럴까??

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

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

[C#] toString assembly  (0) 2013.02.11
[코드이야기] 자바 JDK의 버그  (0) 2013.02.07
OpenMP 디버깅 Tips  (0) 2012.07.31
OpenCL Typedef enum  (0) 2012.04.26
yacc, bison highlight plugin for eclipse  (0) 2011.12.19
if 안에 boolean expression Optimization  (0) 2011.09.16
Posted by Leo 리오 트랙백 0 : 댓글 0

Jon Bentley가 CMU에서 박사과정 학생들을 불러서 binary search 알고리즘을 써보라고 하였다.


당연히 binary search 알고리즘은 아주 기본적인 알고리즘이기 때문에 간단하게 구현할 수 있었을 것이다.


하지만, 모든 학생들의 코드에는 버그가 있었다.


그건, 무엇일까?


아래는 java.util.Arrays에 있는 실제 코드이다.

1:     public static int binarySearch(int[] a, int key) {
2: int low = 0;
3: int high = a.length - 1;
4:
5: while (low <= high) {
6: int mid = (low + high) / 2;
7: int midVal = a[mid];
8:
9: if (midVal < key)
10: low = mid + 1
11: else if (midVal > key)
12: high = mid - 1;
13: else
14: return mid; // key found
15: }
16: return -(low + 1); // key not found.
17: }

이 코드는 1946년에 추가 되었으며, 버그가 발견되지 않다가 1962년, 20년이 지난 후에야 버그가 발견되었다.


앞에서 말했듯이 binary search는 기본적인 알고리즘이여서 아주 많이 쓰였을꺼다.


20년동안 계속 쓰였으니 못해도 수백억번은 불렸을 것이다.


그런데 20년이 지나서야 버그가 발견되었다.


어느 부분이 버그인지 알아 차렸나?


힌트1


힌트2


힌트3


힌트4










몇번째 힌트를 보고 알아차렸는가?


버그가 있는 부분은 아래의 6번째 라인이다.

6:             int mid = (low + high) / 2;

이거만 봐도 머가 버그지? 하면서 알아차리지 못한 사람들이 많을 것이다.


low, high, mid는 모두 integer형이다. 그렇다면 low+high가 2^32가 넘는다면?


그렇다. Overflow가 일어나는 것이다. 그럼 mid값은 완전 다른 값(음수)이 나올것이다.


int mid = low + ((high - low) / 2);

이런식으로 더하기 대신 빼기를 사용하면 overflow를 방지 할 수 있다. 그런데 기존 코드보단 좀 느리다.


int mid = (low + high) >>> 1;

나,

int mid = ((unsigned int)low + (unsigned int)high)) >> 1;

이런식으로 성능의 저하 없이 고쳐 쓸 수 있다.


PS. >>> operator는 logical shift이다. 부호 extend를 하지 않는다.



자. 이제 이 코드는 버그가 없을까? bug-free 과연?


그렇다고 느끼겠지만, 모른다.

물론 2^32는 정말 큰 수이다. 


메모리문제등등 때문이더라도, 이 만큼의 배열을 가지고 있을 확률도 정말 낮을 것이다.



모든 버그를 고치기 위해선 모든 input에대한 테스트가 필요하다.


하지만 그건 당연히 거의 불가능한 일이다.


그렇지만, 이 일례를 통해 


어떤 코드든(특히 내가 짠 코드) 버그가 있을 수 밖에 없고, 


이것을 최소화하기 위해선 잘짜여진 테스트도구가 중요할것이다.



참조 : http://googleresearch.blogspot.kr/2006/06/extra-extra-read-all-about-it-nearly.html



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

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

[C#] toString assembly  (0) 2013.02.11
[코드이야기] 자바 JDK의 버그  (0) 2013.02.07
OpenMP 디버깅 Tips  (0) 2012.07.31
OpenCL Typedef enum  (0) 2012.04.26
yacc, bison highlight plugin for eclipse  (0) 2011.12.19
if 안에 boolean expression Optimization  (0) 2011.09.16
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

http://kerneltrap.org/node/4705


아래처럼 define 되는 likely라는 매크로가 있다.

#define likely(x)       __builtin_expect((x),1)
#define unlikely(x)     __builtin_expect((x),0)



이것은 어떤 기능을 할까?

 if (__builtin_expect (x, 0))
                foo ();

위와같은 코드가 있다고 할 때, 컴파일러는 x가 0이 될 확률이 높다는걸 '알게'된다.

그럼 foo 함수를 실행할 가능성이 적어지고, 컴파일러가 최적화를 할 가능성이 커진다.



구체적인 예를 보자.

int
testfun(int x)
{
        if(x) {
                x = 5;
                x = x * x;
        } else {
                x = 6;
        }
        return x;
}

함수 testfun은 인자 x로 어떤값이 들어올지 모른다. 

레귤러하게 분포된다면, 

x = 6이 될 확률은 1/2^32 이 될 것이다.



그러나, __builtin_expect()을 씀으로써, x는 '거의' 0임을 알려준다.

[kedar@ashwamedha ~]$ cat abc.c
int
testfun(int x)
{
        if(__builtin_expect(x, 0)) {
                              ^^^--- We instruct the compiler, "else" block is more probable
                x = 5;
                x = x * x;
        } else {
                x = 6;
        }
        return x;
}
 
[kedar@ashwamedha ~]$ gcc -O2 -c abc.c
[kedar@ashwamedha ~]$ objdump  -d abc.o
 
abc.o:     file format elf32-i386
 
Disassembly of section .text:
 
00000000 :
   0:   55                      push   %ebp
   1:   89 e5                   mov    %esp,%ebp
   3:   8b 45 08                mov    0x8(%ebp),%eax
   6:   85 c0                   test   %eax,%eax
   8:   75 07                   jne    11 < testfun+0x11 >
                                ^^^ --- The compiler branches the "if" block and keeps "else" sequential
   a:   b8 06 00 00 00          mov    $0x6,%eax
   f:   c9                      leave
  10:   c3                      ret
  11:   b8 19 00 00 00          mov    $0x19,%eax
  16:   eb f7                   jmp    f < testfun+0xf >

생성된 어셈블러를 보면 x가 0이 아닐 때 점프(jne)를 한다.



[kedar@ashwamedha ~]$ cat abc.c
int
testfun(int x)
{
        if(__builtin_expect(x, 1)) {
                              ^^^ --- We instruct the compiler, "if" block is more probable
                x = 5;
                x = x * x;
        } else {
                x = 6;
        }
        return x;
}
                                                                                                    
[kedar@ashwamedha ~]$ gcc -O2 -c abc.c
[kedar@ashwamedha ~]$ objdump  -d abc.o
                                                                                                    
abc.o:     file format elf32-i386
                                                                                                    
Disassembly of section .text:
                                                                                                    
00000000 :
   0:   55                      push   %ebp
   1:   89 e5                   mov    %esp,%ebp
   3:   8b 45 08                mov    0x8(%ebp),%eax
   6:   85 c0                   test   %eax,%eax
   8:   74 07                   je     11 < testfun+0x11 >
                                ^^^ --- The compiler branches the "else" block and keeps "if" sequential 
   a:   b8 19 00 00 00          mov    $0x19,%eax
   f:   c9                      leave
  10:   c3                      ret
  11:   b8 06 00 00 00          mov    $0x6,%eax
  16:   eb f7                   jmp    f < testfun+0xf >

반대로, 위와 같은 경우엔 x가 0일 때 점프하게 된다.



Basic Block(코드의 순서)을 어떻게 배치하느냐에 따라서 성능은 달라 질 수있다.

(branch predictor가 없다면,) CPU는 PIPELINE으로 되있기 때문에 우선 branch는 무시하고, 있는 코드를 순서대로 실행하다 branch가 되야되면 파이프라인을 FLUSH시킨다. 그럼 성능이 떨어질 수 밖에 없다.


조그마한 성능도 중요한 임베디드나 커널 코드에 많이 쓰이는듯 하다.















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

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

[C] likely/unlikely macros  (0) 2012.12.14
LNK2005 already defined in LNK1169: one or more multiply defined symbols found  (0) 2012.11.27
다차원 배열  (0) 2012.10.03
#ifdef #else #endif #elifdef  (0) 2012.08.09
printf debugging  (0) 2012.06.19
sizeof  (0) 2012.06.04
Posted by Leo 리오 트랙백 0 : 댓글 0

LNK2005: __invalid_parameter already defined in

LNK1169: one or more multiply defined symbols found


되게 귀찮게한다. 비쥬얼 스튜디오..

gcc같은 경우엔 strength 따지거나 아니면 그냥 warning으로 넘어가는데













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

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

[C] likely/unlikely macros  (0) 2012.12.14
LNK2005 already defined in LNK1169: one or more multiply defined symbols found  (0) 2012.11.27
다차원 배열  (0) 2012.10.03
#ifdef #else #endif #elifdef  (0) 2012.08.09
printf debugging  (0) 2012.06.19
sizeof  (0) 2012.06.04
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

다차원 배열

2012.10.03 18:30 from Program Language/C

array에는 크게 두 가지가 있다.


1.

int arr1[5];


2.

int *arr2;

arr2 = (int*) malloc(sizeof(int)*5);


배열의 데이터가 1번은 스택에 저장되고 2번은 힙에 저장된다.


이건 1차원 배열에선 크게 상관이 없다.


1차원 배열에선 두 가지를 혼용하여 사용하여도 큰 상관은 없지만


2차원 이상 다차원 배열에선 문제가 된다.





2차원 배열을 보면


1.

int arr1[2][5];



2.

int **arr2;

arr2 = (int**) malloc(sizeof(int*) * 2);

arr2[0] = (int*) malloc(sizeof(int)*5);

arr2[1] = (int*) malloc(sizeof(int)*5);

 or

int **arr2;

arr2 = (int**) malloc(sizeof(int*) * 2);

int brr0[5], brr1[5];

arr2[0] = brr0;

arr2[1] = brr1;




우선 type을 보면

1번에서 arr1은 int*형 type을 갖고

2번에서 arr2는 int**형 type을 갖는다.


저장되는 공간을 보면

1번 배열은 2*5 int형 데이터가 스택에 10개가 쫘르륵 배치된다.

2번 배열은 한 줄(5개)이 각각 다른 장소에 지정되고, arr은 int 5개가 저장되있는 각각의 장소를 가르킨다.


컴파일러 입장에서보면

1번은 컴파일 타임에 위치를 알 수 있고,

2번은 실행시간에야 위치를 알 수 있다.

(다음을 봐보자)


액세스 방식을 보면 (1, 3)액세스

1번에서 arr1[1][3]은 arr + 1 * 5 + 3의 값 같은 의미를 갖는다. = *( arr + 1 * 5 + 3 )

2번에선 arr2[1][3]은 arr[1] + 3과 같은 의미를 갖는다. = *( arr[1] + 3 )

우선 arr[1]주소로 가서 두번째 줄 값이 저장된 장소를 알아낸후 (int*)

5개의 값 중 3번째 값을 가져온다. *( arr[1] + 3)


다시 말하면

1번의 arr1은 바로 값이 있는 위치를 알고 있고 (int*)

2번의 arr2는 값이 있는 위치를 알고있는 그 위치를 알고 있는것이다.(int**)



중요한 다른 함수로 배열 값을 넘길 때를 보자.


1.

선언

void boo( int arr1[2][5] ) or void boo( int arr1[][5] )

호출

boo(arr1);


2.

선언

void foo(int **arr2)

호출

foo(arr2)


이런 식으로 선언, 호출을 해야한다.


1번에선 이 배열은 5개씩 묶음이 2개 있다 읽을 때 5개씩 띄어서 읽어라 라고 알려주기 위한 선언이고

2번에선 이 배열은 2차원 포인터 배열이니 읽을 때는 점프 두번해서 읽어라 라고 알려주기 위한 선언이다.


1번에선 점프하면 값을 읽을 수 있지만

2번에서 값을 읽으려면 점프 + 점프 해야 값을 읽을 수 있는 것이다.



결론


1.

int arr[i][j];

이런 2차원 배열은 int*형을 같는다.

도 arr의 자료형은 int* 이지만

1차원은 int*

2차원도 int *

3차원도 int *

4차원도 int *


2.

int **arr;

이런 다이나믹한 배열에선 배열 변수도 같은 차원을 같는다.

1차원은 int*

2차원은 int **

3차원은 int ***

4차원은 int ****


함수하나에서만 데이터를 사용한다면 큰 무리는 없지만,

다른 함수로 데이터를 넘기기 위해선 두 배열의 차이를 알아야한다.



예제를 들자면








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

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

[C] likely/unlikely macros  (0) 2012.12.14
LNK2005 already defined in LNK1169: one or more multiply defined symbols found  (0) 2012.11.27
다차원 배열  (0) 2012.10.03
#ifdef #else #endif #elifdef  (0) 2012.08.09
printf debugging  (0) 2012.06.19
sizeof  (0) 2012.06.04
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


#ifdef A

...

#endif



#ifdef A

...

#else

...

#endif



#ifdef A

...

#else

#ifdef B

...

#else

...

#endif

#endif



#elifdef (else ifdef)는 어떻게 쓰는건지 모르겠음;

#ifdef A

...

#elifdef B

...

#else

...

#endif

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

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

LNK2005 already defined in LNK1169: one or more multiply defined symbols found  (0) 2012.11.27
다차원 배열  (0) 2012.10.03
#ifdef #else #endif #elifdef  (0) 2012.08.09
printf debugging  (0) 2012.06.19
sizeof  (0) 2012.06.04
extern static 정리  (1) 2012.04.25
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

OpenMP 디버깅 Tips

2012.07.31 19:39 from Program Language

1. printf 

printf로 디버깅 메시지 쓸때는

  1. #pragma omp critical
  2.                 printf("s %d, %d, MBA=%d....

이런식으로 critical section을 정해주자.

printf가 IO라서 오래걸려서 print가 안될때가 있는데

printf동안 다른 쓰레드가 작동 못하게 해준다.


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

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

[C#] toString assembly  (0) 2013.02.11
[코드이야기] 자바 JDK의 버그  (0) 2013.02.07
OpenMP 디버깅 Tips  (0) 2012.07.31
OpenCL Typedef enum  (0) 2012.04.26
yacc, bison highlight plugin for eclipse  (0) 2011.12.19
if 안에 boolean expression Optimization  (0) 2011.09.16
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

printf debugging

2012.06.19 15:06 from Program Language/C

printf로 디버깅 할때 끝에 \n newline을 넣어도


print가 제대로 안 될때가 있다.


이땐 printf후에


fflush(stdout); 을 해주면 출력하고 버퍼를 비우게된다.


(당연히 속도는 늦어진다.)

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

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

다차원 배열  (0) 2012.10.03
#ifdef #else #endif #elifdef  (0) 2012.08.09
printf debugging  (0) 2012.06.19
sizeof  (0) 2012.06.04
extern static 정리  (1) 2012.04.25
caller 찾기  (0) 2012.01.27
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

sizeof

2012.06.04 12:56 from Program Language/C


main()

{

        int arr[16];

        int* nptr;

        int* ptr;

        ptr = malloc(sizeof(int)*16);

        printf("SIZE arr=%d , nptr=%d, ptr=%d\n", sizeof(arr), sizeof(nptr), sizeof(ptr));

}

 

결과 : SIZE arr=64 , nptr=4, ptr=4






main()

{

        char arr[16];

        char* nptr;

        char* ptr;

        ptr = malloc(sizeof(char)*16);

        printf("SIZE arr=%d , nptr=%d, ptr=%d\n", sizeof(arr), sizeof(nptr), sizeof(ptr));

}

 

결과 : SIZE arr=16 , nptr=4, ptr=4 



배열 크기는 하나크기*배열크기

포인터 크기는 한 Word(32bits, 4bytes)



main()

{

        char arr[16];

        char* nptr;

        char* ptr;

        ptr = malloc(sizeof(char)*16);

        printf("SIZE arr=%d , nptr=%d, ptr=%d\n", sizeof(int[16]), sizeof(0), sizeof(int*));

}

 

 결과 : SIZE arr=64 , nptr=4, ptr=4


0은 null pointer를 의미하는듯.






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

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

#ifdef #else #endif #elifdef  (0) 2012.08.09
printf debugging  (0) 2012.06.19
sizeof  (0) 2012.06.04
extern static 정리  (1) 2012.04.25
caller 찾기  (0) 2012.01.27
printf uint64_t uint32_t  (0) 2011.12.01
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

OpenCL Typedef enum

2012.04.26 11:40 from Program Language

이렇게 쓰면 에러가 난다.

intel opencl 컴파일러가 이상한지;

typedef enum {

PRIORITY_HIGHEST     = 3,

 NALU_PRIORITY_HIGHEST     = 3,

 NALU_PRIORITY_HIGH        = 2,

 NALU_PRIORITY_LOW         = 1,

 NALU_PRIORITY_DISPOSABLE  = 0

} NalRefIdc;

아래처럼 치환.

typedef enum NalRefIdc{

PRIORITY_HIGHEST     = 3,

 NALU_PRIORITY_HIGHEST     = 3,

 NALU_PRIORITY_HIGH        = 2,

 NALU_PRIORITY_LOW         = 1,

 NALU_PRIORITY_DISPOSABLE  = 0

} NalRefIdc;


일일히 하기 귀찮다.

regex를 이용한다.

typedef *enum *{([^}|^;]*)} *([^ |^;]*);

typedef enum \2{\1} \2; 로 치환.




vi나 regex replace를 지원하는 에디터를 이용.

대량파일의 경우는 script를 짤 수 있을것 같지만 모르겠다;


visual studio에선 regex가 일반과 다르다;

http://msdn.microsoft.com/query/dev10.query?appId=Dev10IDEF1&l=KO-KR&k=k(VS.REGULAREXPRESSIONBUILDER)&rd=true


신고
Posted by Leo 리오 트랙백 0 : 댓글 0

http://blog.naver.com/entraiger?Redirect=Log&logNo=120122764718



Extern, Static 변수와 함수

Extern, Static 키워드는 변수와 함수에서 쓰일 때 그 의미가 조금 다르다. 형식과 쓰임새가 차이가 있기 때문이다.

 

 

<Extern 전역변수, 지역변수>

extern 변수는 다른 파일에서 변수를 공유해서 쓰기 위해 있는 키워드인데, 전역변수는 키워드를 생략해도 기본으로 extern선언이 되는 성질이 있다.

아래 코드를 보면,

-----------------------

/*main.c*/

#include <stdio.h>

#include "fun1.h"

int a=1; // 변수 선언

void main(){

f();

printf("%d\n",a); // 여기서의 a는 위의 a

}

-----------------------

/*fun1.c*/

#include <stdio.h>

#include "fun1.h"

int a; // 또 다른 선언문

void f(){

a = 7;

printf("%d\n",a); //여기서의 a는 바로 위의 a

}

-----------------------

/*fun1.h*/

void f();

-----------------------

결과는: 7

7

-----------------------

fun1.c에서 바꾼 a값이 main.c의 a에도 적용되어있다.

따라서 main.c와 fun1.c의 전역변수 a는 동일한 변수이거나 항상 싱크가 되어있다는 것을 알 수 있다.

여기에서 main.c와 fun1.c의 선언문에 extern을 붙여보자.

(main.c : int a=1; -> extern int a=1;)

(fun1.c : int a; -> extern int a;)

한쪽파일에만 extern키워드를 쓰든 양쪽에 다 쓰든, 모든 경우에서 결과는 같다.

-----------------------

//간략화한 코드//

/*main.c*/

extern int a=1;

/*fun1.c*/

int a;

-----------------------

/*main.c*/

int a=1;

/*fun1.c*/

extern int a;

-----------------------

/*main.c*/

extern int a=1;

/*fun1.c*/

extern int a;

-----------------------

결과는 항상

7

7

-----------------------

이 된다. 즉, 전역변수로 선언할 경우 extern를 넣든 안 넣든 결과는 똑같아 보인다. extern키워드를 생략해도, 전역변수는 기본으로 extern이기 때문일 것이다. 그런데 초기화구문의 extern과 초기화가 없는 extern선언문은 의미에 차이가 있다. 초기화에 쓰이는 extern은 변수를 새로 정의하는 선언문이라는 의미의 extern이고 실제로 메모리를 할당해서 변수를 정의한다. 그런데 초기화가 없는 선언문은 변수를 새로 정의하는 것이 아닌, “다른 파일(혹은 같은 파일 내의 다른 부분)에서 찾아서 있으면 그것을 가져다 쓰겠다”는 ‘참조선언’의 의미를 갖는 extern이다. 그렇다면 어떤 것이 정의선언이고 어떤 것이 참조선언일까? 아래의 경우를 보자. 아래의 경우는 모두 전역변수로 선언된 경우이다.

-----------------------

전역변수 선언문 종류

1. extern int a=1 -> extern변수 정의 및 초기화

2. int a =1; -> extern 변수 정의 및 초기화

3. extern int a; -> extern 변수(새로 정의하지 않고)참조선언

4. int a; -> extern 변수 정의(0으로초기화) 또는 참조선언

-----------------------

4의 경우는 그 쓰임새가 애매하다. 1이나 2와 쓰일 때는 참조선언이 되고, 3과 쓰일 때는 정의구문이 된다.(컴파일러에서 자동으로 처리함) 애매한만큼 지양해야 할 코드이다.

따라서

-----------------------

/*main.c*/

(extern생략가능)int a=1;

/*fun1.c*/

(extern생략가능)int a=2;

-----------------------

같은 식으로 여러번 초기화를 해주면 둘 다 정의 선언문이므로 ‘여러번 정의했다’는 에러가 뜨게 된다. 같은 a를 두 번 정의했으니 에러가 뜨는 것이다. 아까처럼

-----------------------

/*main.c*/

(extern생략가능)int a=1; //extern변수 정의 및 초기화

/*fun1.c*/

(extern생략가능)int a; //a라는 extern변수를 다른 곳에서 찾아서 쓰겠다는 의미

-----------------------

로 코딩을 했을 경우, 에러가 없다. 이 관계는 다음의 예시에서 더욱 명백해진다.

-----------------------

/*main2.c*/

int a=1;

int a;

void main(){

int b=1; (1)

int b; (2)

}

-----------------------

새로운 main2.c파일에서 a는 전역변수, b는 지역변수로 선언되었다. 처음 int a=1;로 선언과 동시에 초기화되었고, 바로 아래줄에서 int a;는 선언이 아닌, “다른 곳에서 a를 찾아본 뒤에 그것을 쓰겠다”는 표현이 되므로 에러가 나지 않는다.(extern int a;의 생략형이라고 이해하면 쉬울 것이다.) 그러나 지역 변수 b의 경우 “여러번 재정의”에러가 뜬다. 잠깐 지역변수부분만 살펴보면, extern키워드를 통해 초기화하는 것이 아예 불가능하다.

-----------------------

void main(){

extern int b=1;

}

-----------------------

처럼 코딩을 하면, “블록 범위를 사용하여 extern변수를 초기화할 수 없습니다.”라는 에러메세지가 뜬다. 블록 범위, 즉 로컬변수선언시엔 무조건 그 블록 범위 내에서만 사용하게 되어있으므로, 다른 곳에서 이 변수를 사용할 수 있도록 하는 extern초기화 구문을 사용 못하게 막아놓은 것이다. 로컬변수를 다른 곳에서 사용하려면, 함수 결과값으로 끄집어내서 사용하는 수 밖에는 없다. 대신 초기화가 아닌 extern참조선언문은 함수 내에서도 전역에서 쓰듯이 사용이 가능하다.

-----------------------

// 간략화한 코드 //

/*main.c*/

(extern생략가능)int a=1; //extern변수 정의 및 초기화

/*fun1.c*/

void fun1(){

extern int a; // 다른 곳의 extern변수 a를 찾아 쓰겠다는 의미

}

-----------------------

그런데 지역변수에서 extern변수를 참조선언해서 가져다 쓸 때 주의할 점이 있다. 함수 밖에서 참조선언을 할 때는 참조에 실패하면 컴파일 에러가 나지만, 함수 내에서 참조선언을 하면 참조할 것이 없어도 컴파일 상의 에러가 없다.(헉쓰...) 아래의 코드를 보자.

-----------------------

/*main.c*/

int a=1; //extern변수 정의 및 초기화

/*fun1.c*/

extern int b; //extern변수 b참조선언.

-----------------------

결과: “b의 외부기호를 확인할 수 없다”는 에러.

-----------------------

/*main.c*/

int a=1; //extern변수 정의 및 초기화

/*fun1.c*/

void fun1(){

extern int b; //extern변수 b참조선언.

}

-----------------------

결과: no error

-----------------------

아래의 경우처럼 에러가 없으면 나중에 있지도 않은 b변수를 사용할 때 에러가 난다.

-----------------------

/*main.c*/

int a=1; //extern변수 정의 및 초기화

/*fun1.c*/

void fun1(){

extern int b; //extern변수 b참조선언.

printf("%d",b); //변수 b 사용

}

-----------------------

결과: “b의 외부기호를 확인할 수 없다”는 에러.

-----------------------

따라서 결국 b를 안전하게 사용하고 싶으면 전역변수로 참조선언 한 뒤에 함수에서 쓰는 것이 좋은 코딩이라는 이야기다. 아래의 코드가 그렇다.

-----------------------

/*main.c*/

int a=1; //extern변수 정의 및 초기화

int b=2; //extern변수 b를 추가해보았다.

/*fun1.c*/

extern int b; //extern변수 b참조선언은 밖에

void fun1(){

printf("%d",b); //변수 b 사용

}

-----------------------

결과: 2

-----------------------

정리하자면, 전역변수는 초기화인 경우 항상 쓰지 않아도 extern으로 정의가 되며, (즉 필연적으로 모든 전역으로 정의되는 변수는 이름 하나당 변수 하나가 최대다. 단, static변수로 만들지 않을 경우에 한해서) 초기화구문이 아닌 경우엔 extern키워드를 쓰면 확실히 참조선언문이 되지만 extern을 생략하면 경우에 따라 참조선언도 되지만 정의가 되어있지 않았다면 정의구문이 되어버린다. 따라서 extern은 기본적으로 생략 가능하지만 헷갈림을 방지하고 코딩의 가독성을 위해서 아래와 같이 여러 파일에서 사용할 때는 extern키워드를 명시해주는 것이 좋다.

-----------------------

/*main.c*/

extern int a=1; // extern변수를 선언하겠다는 의미. 다른 파일에서도 사용하겠다는 뜻.

생략해도 크게 상관은 없다.(외부에서 쓸 일이 있을 땐 생략 삼가)

/*fun1.c*/

extern int a; //다른 곳에 있는 extern변수를 가져다 쓰겠다는 의미.

생략하면 위의 경우와 구분이 어렵다.(쓰는 것을 필수)

-----------------------

덧붙이자면, 아래의 경우는 어떨까?

-----------------------

/*main.c*/

extern int a;

/*fun1.c*/

extern int a;

-----------------------

둘 모두 참조선언문이라 어디에서도 정의가 되어있지 않기 때문에, 외부 기호를 확인할 수 없다는 에러가 뜨는 것은 당연하다.(전문용어로 찾다가 GG) 그럼 아래의 경우는 또 어떨까?

-----------------------

/*main.c*/

int a;

/*fun1.c*/

int a;

-----------------------

둘 중 하나의 구문의 정의 선언문이 되고 나머지가 참조선언문이 되는 것은 확실하다. 그러나 그게 어느것인지는 나도 모르겠다.

 

 

<Static 전역, 지역변수>

아까 전역변수의 경우에는 extern키워드를 생략해도 자동으로 extern선언이 된다고 했다. 그렇다면 한 파일 내에서 전역변수로 선언해서 쓰고는 싶은데 다른 파일에서 참조하는 것은 막고 싶을 경우(전역변수지만 extern은 아닌)는 어떻게 할까? 그럴 땐 static선언을 하면 된다.

-----------------------

/*main.c*/

static int a=1; //static은 생략하면 안 된다.

/*fun1.c*/

int a; // int a=2, 혹은 extern int a=2도 가능

-----------------------

위의 경우, static변수로 선언한 main.c의 a는 main.c파일 안에서만 사용 가능하고 다른 파일에서는 extern키워드로 참조할 수 없다. fun1.c에서 선언한 a는 main.c의 a와는 다른 메모리에 있는 a이다. 이 경우는 fun1.c의 a가 extern변수이고 main.c의 a는 extern변수가 아닌 mian.c에 속박된 static변수가 되는 것이다. 따라서 아래와 같이 쓰면 에러가 난다.

-----------------------

/*main.c*/

static int a=1; //static은 생략하면 안 된다.

/*fun1.c*/

extern int a;

-----------------------

extern int a;에서 참조선언을 했으므로 다른 곳에서 a를 찾아서 가져다 써야할텐데, 유일한 a는 static으로 묶여있으니 참조를 할 수 없다. 그래서 외부기호 확인 불가능 에러가 뜨는 것이다. 지역변수의 경우엔 extern변수 정의가 불가능하므로 static으로 따로 정의할 필요가 없을 것이다.

<extern, static 함수>

함수의 경우 extern, static개념이 간단하다. 함수는 쓰임새에 따라서 외부함수, 내부함수로 나눠 부르지만 사실 전역, 지역의 개념이 없고 굳이 말하자면 모두 전역이기 때문이다.

기본적으로 함수선언은 생략해도 모두 extern이다.

아래 코드를 보자.

-----------------------

/*main.c*/

#include <stdio.h>

void f(); (1)

void main(){

f();

}

-----------------------

/*fun1.c*/

#include <stdio.h>

#include "fun1.h"

void f(){

printf("fun1.c\n");

}

-----------------------

/*fun1.h*/

void f();

-----------------------

결과: fun1.c

-----------------------

main.c에서 헤더파일을 추가하지도 않았는데, fun1.c의 함수를 가져다 쓴 것을 알 수 있다. fun1.c의 f()함수 선언 자체가 extern으로 되어있고, main.c의 (1)에서도 extern으로 선언해서 다른 파일의 함수를 자동으로 검색해서 call했기 때문이다. 이 점은 위 코드 어느 부분에 extern을 넣어도(심지어 함수정의부분에서도) 똑같은 결과가 나온다는 점에서 알 수 있다. 심지어 (1)을 생략해도 경고가 뜨지만 실행은 된다.(하지만 피해야 할 코딩방법) 함수는 기본적으로 extern으로 선언된다는 것을 알 수 있는 부분이다. 아래와 같이, main.c에서도 같은 이름의 함수 f()를 정의하면 어떻게 될까?

-----------------------

/*main.c*/

#include <stdio.h>

void f();

void main(){

f();

}

void f(){

printf("main.c\n");

}

-----------------------

/*fun1.c*/

#include <stdio.h>

#include "fun1.h"

void f(){

printf("fun1.c\n");

}

-----------------------

/*fun1.h*/

void f();

-----------------------

결과: 함수 재정의 오류

-----------------------

main.c에서 이미 정의한 f()함수가 자동으로 extern이기 때문에, fun1.c에서 또 정의를 해주면 이것도 자동으로 extern이고, extern함수는 고유해야하기 때문에 같은 함수를 두 번 정의한 것으로 인식해서 에러가 나게 된다. 이것을 막고, 두 함수를 모두 사용하고 싶다면

그 때 static 키워드로 선언을 하면 된다.

-----------------------

/*main.c*/

#include <stdio.h>

static void f(); (1)

void main(){

f();

}

void f(){

printf("main.c\n");

}

-----------------------

/*fun1.c*/

#include <stdio.h>

#include "fun1.h"

void f(){

printf("fun1.c\n");

}

-----------------------

/*fun1.h*/

void f();

-----------------------

결과: main.c

-----------------------

잘 보면 mian.c의 f()함수는 static으로 선언되었고, fun1.c의 함수는 extern으로 선언되었으므로, main.c에서 실행한 f()함수는 mian.c에서 정의한 함수가 되며 출력도 그렇게 나온다. 반대로 mian.c의 함수를 extern으로, fun1.c의 함수를 static으로 정의해도 결과는 똑같다. 둘 다 static으로 정의해도 마찬가지다. 이유는 생각해보시길...

신고

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

printf debugging  (0) 2012.06.19
sizeof  (0) 2012.06.04
extern static 정리  (1) 2012.04.25
caller 찾기  (0) 2012.01.27
printf uint64_t uint32_t  (0) 2011.12.01
ignoring return value of 'system', declared with attribute warn_unused_result  (0) 2011.11.11
Posted by Leo 리오 트랙백 0 : 댓글 1

티스토리 툴바