메모리 타입.

- 메모리 타이별 특성.

  1. 레지스터 

     - 쓰래드 내부에서 몇 개를 사용하는지는 알 수 없다. 다만 툴을 이용해서 확인은 가능.

     - 단항연산을 할수록, 사용하는 레지스터 갯수는 줄어든다.

     

              

  2. 로컬메모리

    - 쓰레드 로컬에서 사용

    -  제한된 레지스터 개수를 초과하면 로컬 메모리에 할당된다.

    - 레지스터냐, 로컬이냐 할당 기준은 명확하지 않다.

  3. 공유메모리 Shared Memory

       

      - 소개 : 블록내 할당된 쓰레드 간에 공유되는 메모리.

      - 특성 :

           - SM과 인접하기 때문에, 메모리양이 적고, 속도가 빠르다.

           - 워프에 의해 영향을 받는다.

           - 같은 크기의 메모리를 32개 가진다.  뱅크 갯수는 Compute ability 테이블을 참조한다.  = 32개. 보통 warp 갯수와 동일하다고 한다.

           - 병행성을 위해, 메모리 뱅크 개념을 알아야 함.

           - 뱅크 내 메모리 단위는 4byte/8byte 로 구성할 수 있고. API로 설정 가능하다.

               cudaDeviceSetSharedMemConfig or cudaFuncSetSharedMemConfig.

           - 동기화. 필요.

               __synchthread();

      - 선언.

         1. __shared__ float tile[size_y][size_x]; 

         2. extern __shared__ int tile[];  

             - 사이즈가 결정되지 않았을 때 사용한다.

             - 커널 함수를 호출할 때 명시하면 된다.  kernel<<<grid, block, shared_memory_size>>>(....);                            

         3. 커널내부, 외부에서 선언할 수 있다. 

    

      - 메모리 뱅크 

        - 메모리 밴드를 최대한 활용 목적으로  병행성을 극대화 하기 위해, 같은 크기의 메모리 블록을[64K/ 32[뱅크갯수] ] 32개 가진 메모리 모듈를 뱅크라고 한다.       

        - 계산식 : 뱅크 수 32개. 

                 1 열은 32 * 4 바이트 = 128 BYTE

                 2. 한 열의 뱅크 크기[4바이트 / 1블록  크기] =  1024 * 64  / 128  = 512개.     

                   0번 스레드 -->뱅크 0   [ 0 BYTE] [128BYTE]  ~~~~

                   1번 쓰레드 --> 뱅크 1  [4 BYTE][132BYTE ]

                     ~~

                  31번 쓰레드 --> 뱅크 31 [ 124 BYTE][252BYTE]

         - 관련 이슈: 뱅크 충돌.

             - 2개 이상의 쓰레드가 동시에 하나의 뱅크에 접근하려고 했을 때 뱅크 충돌이라고 한다.

             - 뱅크 충돌시, 처리는 순차적으로 이루어지기 때문에,   1/ 충돌횟수 만큼 성능이 떨어지게 된다.

      - 동기화.

          1. 공유메모리는 비동기적으로 처리가 되기 때문에, 동기화 작업이 필요하다.

          2. __synchthreads() 를 통해 동기화가 이루어진다.                     

              

  4. 글로벌 메모리.

     -  디바이스의 DRAM을 말한다.

     -  cudaMalloc() 호출시 할당되는 메모리.

  5. 상수 메모리.

     -읽기 전용. / 캐쉬 지원 

     - 호스트영역: 쓰기전용

     - 디바이스영역: 읽기 전용.

     쓰기 예)

       __constant__ int cData[6];

       

       int aData[6] = {1, 2, 3, 4, 5, 6};

      cudaMemcpyToSymbol ("cData",&aData,sizeof(aData));

       

  6. 텍스트 메모리.

    2.2 이후로 지원안함.