병렬 프로그래밍13 [실습] Nvidia의 MergeSort 예제 분석 - 1 [실습] CUDA MergeSort 구현 (Naive 버전) (tistory.com) [실습] CUDA MergeSort 구현 (Naive 버전) CUDA에서 정렬을 구현한다고 할때, 가장 먼저 드는 생각은 우리가 알고 있는 알고리즘들(Quick Sort, Merge Sort, Buble Sort 등등.. )을 병렬로 바꿔볼까 하는생각들을 하게 될것 입니다. 그러면 이게 생각 hotstone.tistory.com 이전 시간에 기존에 CPU에서 동작하는 방식을 그대로 CUDA에서 Merge Sort를 구현했었습니다. 그런데 당연하게도 std::qsort에 비해 한참 떨어지는 성능을 보여줬었고, 따라서 CUDA에서는 정렬을 구현한다면 어떻게 해야 하는지 Nvidia의 예제를 참고하게 되었습니다. 오늘은 제가.. 2023. 3. 28. [실습] CUDA MergeSort 구현 (Naive 버전) CUDA에서 정렬을 구현한다고 할때, 가장 먼저 드는 생각은 우리가 알고 있는 알고리즘들(Quick Sort, Merge Sort, Buble Sort 등등.. )을 병렬로 바꿔볼까 하는생각들을 하게 될것 입니다. 그러면 이게 생각보다 쉽지 않다는 것을 느낄 수 있고 여기에 성능까지 고려하게 된다면 CUDA를 사용해 병렬로 정렬하다는 것이 의미 없어보이기도 합니다. (왜냐하면, std::sort에 비해 성능이 너무 안나오기 때문이죠.) 단순히 우리가 이번에 이야기 해볼 Merge Sort만 놓고 고민해 본다고 하더라도, 분활하여 비교하는 부분은 병렬로 처리할 수 있다고 생각하지만, 이후에 Merge 하는 과정은 특정 Thread만 수행하게 되어 CUDA를 사용하는 의미를 많이 잃게 됩니다. (일반적인 M.. 2023. 3. 21. [빌드] CUDA에서 Separate Compilation과 Linking 이슈 최근 PTX 어셈블리를 활용하여 실행 중인 SM 확인하기 (tistory.com) 이라는 글을 보고서, Debug 용으로 실행중인 Kernal의 Warp나 Thread 정보들을 출력할 수 있는 유틸 함수들을 추가하면 좋겠다는 생각을 하게 되었습니다. 그래서 기존에 헤더로만 구성되어 있던 util 코드를 소스코드와 헤더로 분리하고 위에 블로그에서 아래의 코드들을 복사해서 가져왔습니다. // utils.cu #include "utils.cuh" __device__ uint32_t __mysmid() { uint32_t smid; asm volatile("mov.u32 %0, %%smid;" : "=r"(smid)); return smid; } __device__ uint32_t __mywarpid() { u.. 2023. 3. 7. [실습] CUDA의 Cooperative Groups 병렬 알고리즘을 작성하다 보면 Thread들을 그룹화하고 이러한 그룹들간에 동기화를 수행해야 하는 경우가 많이 있습니다. 그런데 CUDA 버전9 이전에는 Thread Block 내의 모든 Thread들을 동기화 하는 __syncthreads( )만 제공 되어졌습니다. 그러다 보니 Thread Block보다 더 작은 그룹을 정의하고 이에 대한 동기화하고자 하는 요구가 많이 존재했었습니다. 왜냐하면 보다 작은 그룹을 구성함으로써 불필요한 동기화를 최소화 할수있고, 이에 더해 설계에 대한 유연성을 높일 수 있기 때문입니다. 그리고 Cooperative Group는 이러한 요구를 충족시킬 수 있게 하위 블록 및 다중 블록 세분성에서 Thread 그룹을 명시적으로 정의하고 Thread에 대한 동기화와 같은 집합적.. 2023. 3. 5. 이전 1 2 3 4 다음
