GPU 커널에 적용 가능한 CPU 병렬 처리 교훈: 의존성 체인, 부동소수점 비결합성, 가짜 공유, 메모리 바운드 루프 [P]

제가 진행하는 무료 GPU 커널 엔지니어링 시리즈의 세 번째 세션(제 채널에서 업로드했으며, 이전 게시물: [LINK TO FIRST POST])에서는 GPU 코드는 사용하지 않고 CPU 병렬 처리에 대해 다룹니다. 대부분의 GPU 성능 아이디어는 증폭된 CPU 아이디어이기 때문에, 제가 가장 유용하다고 생각한 부분들을 여기에 정리했습니다. 단일 누산기(accumulator)는 의존성 체인입니다. sum += a[i] * b[i]와 같은 경우, 모든 덧셈은 이전 덧셈을 기다리기 때문에 루프가 코어의 일부만 사용합니다. 비행 중 필요한 독립적인 연산(Independent ops needed in flight)은 지연 시간(latency) x 처리량(throughput)에 관한 것이며, 네 개의 누산기는 동일한 산술 연산을 가진 네 개의 독립적인 체인을 제공합니다. 이를 캐시 내 배열과 -fno-tree-vectorize를 사용하여 측정하거나, 두 가지 효과를 동시에 타이밍하는 실수를 하지 마십시오. 컴파일러가 자동으로 고쳐주지는 않습니다. 부동소수점(FP) 덧셈은 결합법칙(associative)을 따르지 않으므로, 달리 허용하지 않는 한 순서를 유지해야 합니다(-ffast-math 또는 루프당 #pragma omp simd reduction(+:sum)). 동일한 비결정성(non-determinism)이 GPU 리덕션에서도 나타납니다. 분기 비용은 예측 가능성 문제입니다. 정렬된 데이터 대 무작위 데이터에 대한 루프는 제 컴퓨터에서 -O0일 때 [X]배의 격차를 보여주었고([CPU], GCC [version]), 컴파일러가 해당 분기를 제거했기 때문에 대부분 -O2에서는 사라졌습니다. 하드웨어를 탓하기 전에 생성된 코드를 확인해 보십시오. 스레드에 관해서는, Amdahl은 8개 코어에서 90% 병렬화가 8배가 아니라 약 4.7배의 성능을 준다고 말합니다. 작은 배열 내에서의 스레드별 부분 계산(Per-thread partials)은 가짜 공유(false sharing)를 유발하는데, 이는 여러 코어 사이에서 오가는 단일 64바이트 캐시 라인에 위치하기 때문입니다. 패딩(Padding)을 적용하니 저에게는 해결책이 되었고([N] 스레드에서 [X]배 더 빨랐습니다), 다만 이 격차는 스레드가 별도의 물리적 코어에 배치되는지에 따라 달라집니다. 내적 곱셈(dot product)은 메모리 바운드(memory-bound)입니다. 8바이트당 2 FLOPs를 수행하므로, [BW] GB/s에서는 아무리 많은 SIMD나 스레딩을 추가해도 [Y] GFLOP/s 근처에서 제한됩니다. 메모리 지연 시간 은폐(Hiding memory latency)에는 비행 중 지연 시간 x 대역폭 바이트가 필요하며, 이것이 GPU가 수천 개의 스레드를 실행하는 이유입니다. 영상:

주의사항: 제시된 수치는 한 컴퓨터에서 나온 것이며, 내장 함수(intrinsics)는 AVX2를 가정하고 작성되었고, 지연 시간/처리량 수치는 일반적인 인텔 값을 기준으로 합니다. 특히 가짜 공유 설명에 오류가 있거나 지나치게 단순화한 부분이 있습니까?

그리고 어떤 CPU 관련 교훈이 GPU 작업에 가장 잘 적용되었나요? /u/Winter_Mistake_3185 님이 제출 [link] [comments]