결합 읽기와 보폭 읽기

메모리 병합은 각 스레드가 어느 주소에 접근하는지에 따라 달라집니다. 스레드에서 주소로의 매핑에 따라 하나의 전송으로 워프 전체를 처리할 수 있는지가 결정됩니다.

인접한 스레드가 인접한 주소를 읽으면 워프가 하나의 연속된 구간을 다룹니다. 이렇게 정돈된 배치가 바로 병합된 접근입니다.

스레드 0은 인덱스 0에, 스레드 1은 인덱스 1에 접근하고, 이런 식으로 이어집니다. 32개의 주소가 하나의 정렬된 캐시 라인에 모두 들어가므로 GPU는 단 한 번의 전송만 필요로 합니다.

보폭은 연속된 스레드가 접근하는 주소 사이의 일정한 간격입니다. 그 간격이 요소 하나를 넘어서는 순간 메모리 병합이 깨지기 시작합니다.

보폭이 2이면 스레드 0은 0을, 스레드 1은 2를, 스레드 2는 4를 읽습니다. 이제 워프가 메모리의 두 배에 걸쳐 있으므로 더 많은 전송이 필요합니다.

보폭을 두 배로 늘리면 대략 접근하는 캐시 라인의 수도 두 배가 됩니다. 큰 보폭은 워프가 많은 전송을 하게 만들면서 각 라인의 일부만 사용하게 할 수 있습니다.

행 우선 행렬에서 각 스레드에 열 하나를 통째로 맡기면 보폭이 매우 커집니다. 코드에서는 깔끔해 보이지만 각 워프의 접근을 조용히 흩어 놓습니다.

대개 해결 방법은 스레드에 따라 가장 빠르게 변하는 인덱스를 서로 바꾸는 것뿐입니다. 연속된 스레드가 연속된 메모리를 따라가게 하면 읽기가 다시 병합됩니다.

병합된 패턴이라도 캐시 라인의 중간에서 시작하면 손해를 봅니다. 정렬되지 않은 오프셋은 워프를 경계 너머로 이동시켜 하나의 읽기를 두 번으로 나눕니다.

패턴을 판단할 때 스레드 하나만 떠올리지 마십시오. 32개 레인이 한꺼번에 접근한다고 생각하고, 그 주소들이 함께 몇 개의 라인을 덮는지 살펴보십시오. 🔍

가장 빠르게 변하는 인덱스가 threadIdx.x를 따라가게 하십시오. 이 한 가지 습관만으로도 전역 메모리 읽기의 대부분을 별도 비용 없이 병합된 형태로 유지할 수 있습니다.

가장 잘 병합되는 접근 패턴을 고르십시오.

병합된 읽기는 인접한 요소들을 함께 유지하지만, 보폭이 요소들을 여러 라인에 흩어 놓는다는 것을 배웠습니다. 가장 빠른 인덱스를 threadIdx.x가 결정하도록 하십시오. 🎉