GPU Gems 3의 병렬 접두사 알고리즘에 사용되는 CONFLICT_FREE_OFFSET 매크로

StackOverflow https://stackoverflow.com//questions/9689185

  •  13-12-2019
  •  | 
  •  

문제

우선, 알고리즘에 대한 링크는 다음과 같습니다.

GPU Gems 3, 39장: CUDA를 사용한 병렬 접두사 합계(스캔).

뱅크 충돌을 피하기 위해 NUM_BANKS(즉, 계산 가능성 2.x 장치의 경우 32) 요소마다 공유 메모리 배열에 패딩이 추가됩니다.이는 다음과 같이 수행됩니다(그림 39-5 참조).

int ai = offset*(2*thid+1)-1
int bi = offset*(2*thid+2)-1
ai += ai/NUM_BANKS
bi += ai/NUM_BANKS
temp[bi] += temp[ai]

ai/NUM_BANKS가 매크로와 어떻게 동일한지 이해가 되지 않습니다.

   #define NUM_BANKS 16  
   #define LOG_NUM_BANKS 4  
   #define CONFLICT_FREE_OFFSET(n) \  
          ((n) >> NUM_BANKS + (n) >> (2 * LOG_NUM_BANKS))

같지 않나요? 

n >> LOG_NUM_BANKS

도움을 주시면 감사하겠습니다.감사해요

도움이 되었습니까?

해결책

해당 코드를 작성하고 기사를 공동 집필했으며, 기사를 스캔 알고리즘에 대한 학습용으로만 사용하고, 거기에 포함된 코드를 사용하지 말 것을 요청합니다.그것은 CUDA가 처음이었을 때 작성되었으며 저는 CUDA를 처음 접했습니다.CUDA에서 최신 스캔 구현을 사용하는 경우 은행 충돌 방지가 필요하지 않습니다.

쉬운 방법으로 스캔을 수행하려면 다음을 사용하십시오. thrust::inclusive_scan 또는 thrust::exclusive_scan.

정말로 스캔을 구현하고 싶다면 이 기사와 같은 최신 기사를 참조하세요. [1].또는 더 빠른 코드를 사용하는 실제 작업의 경우 좀 더 많은 연구가 필요합니다. [2].아니면 읽어보세요 Sean Baxter의 튜토리얼 (후자는 스캔 알고리즘에 대한 중요한 작업에 대한 인용을 포함하지 않지만).

[1] 슈바브라타 센굽타(Shubhabrata Sengupta), 마크 해리스(Mark Harris), 마이클 갈랜드(Michael Garland), 존 D.오웬스."다코어 GPU를 위한 효율적인 병렬 스캔 알고리즘".Jakub Kurzak, David A.Bader 및 Jack Dongarra, 편집자, 멀티코어 및 가속기를 사용한 과학 컴퓨팅, Chapman & Hall/CRC 컴퓨팅 과학, 19장, 413~442페이지.테일러 & 프랜시스, 2011년 1월. http://www.idav.ucdavis.edu/publications/print_pub?pub_id=1041

[2] 메릴, D.그리고 그림쇼, A.스트림 아키텍처를 위한 병렬 스캔.기술 보고서 ​​CS2009-14, 버지니아 대학교 컴퓨터 공학과.12월2009.

라이센스 : CC-BY-SA ~와 함께 속성
제휴하지 않습니다 StackOverflow
scroll top