CUDA Memory

1. Register

 - 온 칩 프로세서에 있는 메모리. 1개의 크기는 32bit (GTX 770 : Block 당 65536개)
 - 액세스 속도 : 1 GPU Cycle 이내 (가장 빠른 메모리)
 - 커널 함수 안에서 사용하는 로컬 변수
 - 스레드 수에 따라 1개의 스레드당 사용가능한 레지스터 수는 달라짐
 - 너무 많은 로컬 변수, 로컬 변수를 배열로 사용 할 경우 Local 메모리에 할당 됨 (Global Memory)

2. Shared Memory

 - 온 칩 프로세서에 있는 메모리. SM 내의 Thread 들이 공유하여 사용 (GTX 770 : 48 KBytes , 단 16 Bytes는 시스템에서 사용)
 - 액세스 속도 : 1 GPU Cycle (L1 캐시와 동등한 속도)
[-] Collapse

__global__ void Kernel(int* In, int* Out)
{
    __shared__ int SharedMemory[512];

    SharedMemory[threadIdx.x] = In[ThreadIdx.x];
    __syncthreads();

   Out[ThreadIdx.x] = SharedMemory[ThreadIdx.x];
    __syncthreads();
}

int main()
{
    Kernel<<<1,512>>>(dev_In, dev_Out);
}

3. Constant Memory

 - 읽기전용 캐시 (64 KBytes)
 - 액세스 속도 : Write (from DRAM) : 400 ~ 600 Cycles , Read : 레지스터와 동일 속도
[-] Collapse

__constant__ int cData[6]; //  반드시 배열로 선언할 것. cudaMemcpyToSymbol에 Constant 주소 넣는 부분에 & 연산자 안됨
int hData[6] = { 1, 2, 3, 4, 5, 6 };

cudaMemcpyToSymbol (cData, &hData, sizeof(hData));

__global__ void Kernel()
{
    int a = cData[0];
}

4. Global Memory 

 - 비디오 카드에 장착된 DRAM 메모리 (GTX 770 : 2048 MBytes)
 - 액세스 속도 : 400 ~ 600 Cycles (GTX 770 : Memory Clock rate : 3505 Mhz)
[-] Collapse

cudaError_t cudaMalloc(void** devPry, size_t count); //  -할당
cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind); // Global 메모리 <-> Host 메모리
cudaError_t cudaFree(void* devPtr); // 해제

5. Texture Memory

 - 캐시 읽기를 지원하는 읽기 전용 메모리 (GTX 770 : 1D=(65536), 2D=(65536, 65536), 3D=(4096, 4096, 4096))
 - 액세스 속도 : DRAM에서 읽어들인 데이터는 캐시에서 재사용됨
 - 공간 구역성(Spatial locality) 을 자주 드러내는 메모리 접근 패턴을 위해 설계
[-] Collapse

#define X_SIZE 800
#define Y_SIZE 600

texture<float> tex_fData; // 전역으로 선언

__global__ void Kernel(float *dev)
{
     float f = tex1Dfetch(texArray, x * X_SIZE + y); // Texure에서 읽어오기
}

int main()
{
    cudaMalloc((void**)&dev_fData, sizeof(float) * X_SIZE * Y_SIZE); // 메모리 할당
    cudaBindTexture(NULL,tex_fData, dev_fData, sizeof(float) * X_SIZE * Y_SIZE); // Texture로 Bind

    cudaMemcpy(dev_fData, host_fData, sizeof(float) * X_SIZE * Y_SIZE, cudaMemcpyHostToDevice); // Texture에 값 입력

    cudaUnbindTexture(tex_fData); // Texure Unbind
    cudaFree(dev_fData); // Memory 해제
}

[-] Collapse

#define X_SIZE 800
#define Y_SIZE 600

texture<float,2> tex_fData; // 전역으로 선언

__global__ void Kernel(float *dev)
{
     float f =tex2D(texConstSrc, y, x);  // Texure에서 읽어오기
}

int main()
{
    cudaMalloc((void**)&dev_fData, sizeof(float) * X_SIZE * Y_SIZE); // 메모리 할당
    cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
    cudaBindTexture2D(NULL, (NULL,tex_fData, dev_fData, desc, Y_SIZE, X_SIZE, sizeof(float) * Y_SIZE); // Texture로 Bind

    cudaMemcpy(dev_fData, host_fData, sizeof(float) * X_SIZE * Y_SIZE, cudaMemcpyHostToDevice); // Texture에 값 입력

    cudaUnbindTexture(tex_fData); // Texure Unbind
    cudaFree(dev_fData); // Memory 해제
}