타일 단계 반복하기

전체 행과 전체 열의 곱은 타일 하나에 담기에는 너무 큽니다. 따라서 긴 합을 너비가 TILE인 조각으로 나누고, 단계마다 조각 하나를 처리합니다.

행렬의 너비가 N이고 타일의 너비가 TILE이라면, 내부 차원 전체를 처리하려면 N / TILE개의 단계가 필요합니다.

로드-동기화-계산 과정을 단계에 대한 반복문으로 감쌉니다. 각 반복에서 타일 창이 A의 행과 B의 열을 따라 더 멀리 이동합니다.

지역 sum 변수는 반복문 밖에 있으므로 계속 커집니다. 각 단계에서 내적의 해당 조각을 현재 합계에 더합니다.

각 단계에서 A에서 읽는 열과 B에서 읽는 행을 phase * TILE만큼 이동합니다. 이렇게 창이 앞으로 이동합니다.

앞에서와 마찬가지로 로드가 끝난 후 __syncthreads()를 호출하여, 어떤 스레드도 타일을 계산하기 전에 타일 전체가 준비되도록 합니다.

내부 반복문은 방금 로드한 타일만 사용하여 TILE개의 곱을 sum에 더합니다. 이는 최종 내적의 한 조각을 계산합니다.

느린 스레드가 아직 타일을 읽는 동안 다른 스레드가 타일을 덮어쓰지 않도록 각 단계를 또 다른 __syncthreads()로 끝냅니다. 🚧

한 장벽은 로드 후 읽기를 보호하고, 다른 장벽은 읽기 후 로드를 보호합니다. 두 장벽이 함께 모든 스레드를 단계마다 보조를 맞추게 합니다.

모든 단계가 끝나면 누적된 sum이 완전한 내적이 됩니다. 범위 확인으로 보호한 뒤 C에 한 번 저장합니다.

N이 TILE의 정확한 배수가 아니면 범위를 벗어난 원소를 0으로 로드하여 추가 곱셈이 합계에 영향을 주지 않도록 합니다.

단계를 따라 반복하면서 타일을 이동하고, 두 번 동기화하며, 부분 합계를 하나의 총합에 누적했습니다. 이제 얼마나 빨라졌는지 측정해 보겠습니다. 📈