OpenCL Address Spaces

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 메모리를 사용하자.