CUDA 개요, 섹션 3: 메모리 정렬

안녕하세요, 제 큐다 시리즈로 돌아온 것을 환영합니다.이번 주에는 CUDA의 장치 메모리 정렬에 대해 논의해 봅시다. 이것은 자주 오해받는 개념이자 코드 속도가 느린 원인입니다.

CUDA 메모리 정렬
GPU 메모리는 32비트, 64비트 및 128비트(4바이트, 8바이트, 16바이트)의 그룹으로 액세스할 수 있습니다.만약 이동 중인 데이터의 크기가 이 값의 배수가 아니라면, 예를 들어 1, 2, 3개의 문자열을 이동하면, 이 수조는 적당한 수량의 빈 값으로 채워져 그 중의 배수에 반올림할 것이다.CUDA가 실행될 때 메모리 접근을 최대한 줄일 것입니다. 더 많은 메모리 접근은 한 번에 실행할 수 있는 이동과 복사 명령의 수량 (흡수량) 을 줄일 수 있기 때문입니다.
따라서 그룹 포인터가 정렬되지 않을 때 메모리 접근이 느려질 수 있습니다.64비트와 128비트를 정렬하면 오류가 발생할 수도 있습니다.cudaMalloc()이 되돌아오는 바늘이 맞춰져 있는 이상 수조가 어떻게 이럴 수 있느냐고 스스로에게 물어볼 수도 있다.우리 예를 하나 봅시다.대형 메모리 구조의 수조나 다른 정렬되지 않은 원소의 분배 공간을 가정하고 바늘을 늘려 다른 수조 원소의 핸들을 되돌려줍니다.포인터가 정렬되지 않으므로 액세스하면 처리량이 감소합니다.이것은 분배된 원소의 크기를 16, 8, 4바이트의 가장 가까운 배수로 반올림함으로써 완화시킬 수 있다.
또한 모든 위젯의 전체 치수가 한 치수의 배수에 맞지 않으면 구조 패턴도 정렬되지 않은 액세스 속도에 영향을 받습니다.__align__(n) 한정부호는 구조와 클래스 이전에 사용할 수 있으며 모든 구성원을 n자리로 강제로 정렬할 수 있다.CUDA의 경우 n은 4의 배수여야 한다. 예를 들어 4, 8, 12, 16...문제의 규모가 크지 않으면 16바이트로 맞추는 것이 좋다.예를 들면 다음과 같습니다.

struct __align__(16) {
    float x;
    float y;
    float z;
};

글로벌 선언 변수의 주소(CUDA 파일):

   int my_array { 1, 2, 3, 4};

또는 cudaMalloc()을 사용하여 할당된 메모리는 항상 32바이트 또는 256비트 경계에 맞지만 512비트 또는 1024비트와 같은 더 큰 경계에 맞출 수 있습니다.
함수에 정의된 일부 국부 변수는 GPU 레지스터를 너무 많이 사용하기 때문에 메모리에 저장됩니다.예를 들어 매우 큰 수조와 로컬 정의의 구조.이 변수들은 정렬된 메모리 접근을 사용해야 합니다. 그렇지 않으면 전역 메모리와 같은 지연 처벌을 받을 것입니다.
CUDA 프로그래머 설명서에 따르면:

Local memory is however organized such that consecutive 32-bit words are accessed by consecutive thread IDs. Accesses are therefore fully coalesced as long as all threads in a warp access the same relative address (e.g., same index in an array variable, same member in a structure variable).

내가 아는 바에 의하면, 이것은 당신이 이렇게 조직된 그룹을 가지고 있다면,

         int a[6]
    Thread
       0     1       2
    +-----+-----+
a[0]|  1  |  1  |
    +-----+-----+
a[1]|  2  |  2  |    ....
    +-----+-----+
a[2]|  3  |  3  |
    +-----+-----+

         ....

int의 접근은 32비트 (4바이트) 로 정렬되며, 같은 그룹 요소에 동시에 접근하는 모든 라인은 정렬된 접근을 사용합니다.만약에 그룹 원소가 구성원 변수로 교체된다면 구조도 이렇게 말할 수 있다. 또한 32비트 너비의 배수가 아닌 그룹 원소(예를 들어 char 또는 그 중 일부 int과 char의 구조를 포함하는 것)는 그것들이 교체될 때까지 채워져서 그것들도 정렬 접근을 할 수 있게 한다.
내 추측에 의하면 로컬 메모리 귀속 변수를 맞추면 내부에 8개의 128비트 메모리 접근을 허용하고 비뚤어진 32개의 라인의 4바이트 요소의 모든 32개의 복사본을 접근할 수 있다. 블록의 라인은 어떠한 상황에서도 로컬 변수의 같은 복사본을 공유할 수 없다는 것을 기억하십시오. 이것은 다음 주제의 주제입니다.

서로 다른 블록의 라인이 다른 블록의 변수에 접근할 수 있습니까?
답은 정해지지 않았습니다. 다른 블록의 라인이 다른 블록에 있는 변수에 접근하려면 이 변수를 전역 공간에 넣고 & 인용 기호를 사용하여 바늘을 가져와야 합니다.

동기화 스레드 실행
같은 격자에 있는 모든 스레드(즉 모든 블록에 있는 모든 스레드)는 함수의 임의의 위치에 __syncthreads()이라는 CUDA 함수를 삽입하여 동기화할 수 있다.__syncthreads()은 서열화점을 충당하고 모든 라인이 이 함수 이전에 코드의 집행을 완성한 후에야 이 함수 이후의 코드를 계속 집행할 수 있도록 한다.이것은 일부 스레드가 다른 스레드가 읽고 있는 전역 메모리에 쓸 때 경합을 방지할 수 있다.

2D 및 3D 동적 패턴cudaMallocPitch(void **mem, size_t *pitch, size_t width, size_t height)과 cudaMalloc3D(cudaPitchPtr *ptr, cudaExtent extent) 함수는 각각 2차원과 3차원 그룹을 편리하고 안전하게 분배하여 메모리 CUDA 변수의 정렬 요구를 충족시킬 수 있도록 합니다.뒤의 함수 cudaMalloc3D()에서 cudaExtent 유형은 make_cudaExtent(size_t width, size_t height, size_t depth)을 호출하여 되돌아오는 유형이다.이 예에서 패턴과 피치는 각각 ptr->ptr과 ptr->pitch에 있습니다.폭, 높이 및 깊이 지수는 a[x + y*pitch + z*pitch*height]을 사용하여 스토리지에 액세스할 수 있습니다. 여기서 pitch은 cudaMallocPitch()과 cudaMalloc3D()이 반환하는 값입니다.이 x의 값은 너비와 관련이 있기 때문에 인덱스가 1개의 위치를 증가하면 그룹이 하나의 너비 단위로 전진하고 y의 값은 높이와 관련이 있기 때문에 인덱스가 pitch개의 위치를 증가하면 그룹이 하나의 높이 단위로 전진한다. 만약에 3D 그룹이 분배된다면z은 깊이에 대응하고 인덱스를 pitch*height개의 위치로 늘리면 수조를 깊이 단위로 전진시킬 수 있다.
이 인덱스 방안은 기억하기 쉬우며, 바늘로 인덱스를 할 수 있습니다. 바늘은 더 많은 함수에 의해 되돌아오고 전달되는 것 외에, 그룹을 사용하는 것과 같습니다.지침 산술과 지침 a을 사용하는 등효 인덱스 방안은 (a + x + y*pitch + z*pitch*height)이다.물론, 이런 식으로 음조를 사용하면 메모리 접근의 정렬을 확보할 수 있다.

결론
현재 CUDA가 GPU 메모리 접근을 어떻게 정렬하고, 2차원과 3차원 진열을 안전하게 분배하는지, 여러 개의 스레드가 같은 공유 메모리 위치를 동시에 읽거나 쓰는 것을 방지하기 위해 스레드를 동기화하는 방법을 알고 있습니다.
만약 당신이 이 문장에서 어떤 잘못을 보았다면, 내가 그것들을 바로잡을 수 있도록 나에게 알려 주십시오.