워프 투표 기능 정보

Question

CUDA 프로그래밍 가이드에는 워프 투표 기능 "_ all", " _any"및 "__ballot"개념이 도입되었습니다.

내 질문은 : 어떤 응용 프로그램에서이 3 가지 기능을 사용할 것입니까?

Answer 1

__ballot는 CUDA-histogram과 비트 마스크의 빠른 생성을위한 CUDA NPP 라이브러리에 사용되며, 고유 __popc과 결합하는 부울 감소의 매우 효율적인 구현을 할 수 있습니다.

__all 및 __any은 __ballot이 도입되기 전에 축소되어 사용되었지만 다른 용도는 생각할 수 없습니다.

Answer 2

__ballot의 원형은 predicate 0이 아닌 경우, __ballotN 스레드 인덱스 인 N 번째 비트 세트를 가진 값을 반환 다음

unsigned int __ballot(int predicate); 

이다.

atomicOr 및 __popc과 결합하여 참 참을 가진 각 워프의 스레드 수를 누적하는 데 사용할 수 있습니다.

사실, atomicOr의 프로토 타입은

int atomicOr(int* address, int val); 

및 atomicOr이 값이 address 가리키는 읽어 val와 비트 OR 작업을 수행하고 address로 다시 값을 기록하고로의 이전 값을 반환 return 매개 변수.

__popc은 32 비트 매개 변수로 설정된 비트 수를 반환합니다.

따라서, 지침

volatile __shared__ u32 warp_shared_ballot[MAX_WARPS_PER_BLOCK]; 

const u32 warp_sum = threadIdx.x >> 5; 

atomicOr(&warp_shared_ballot[warp_num],__ballot(data[tid]>threshold)); 

atomicAdd(&block_shared_accumulate,__popc(warp_shared_ballot[warp_num])); 

은 술어가 true 인 스레드의 수를 계산하는 데 사용할 수 있습니다.

은 자세한 내용은, 내가 D.M 휴즈 등으로 인 - 커널 스트림 압축을 언급 할 __ballot API를 사용하여 알고리즘의 예를 들어 쉐인 쿡, CUDA 프로그래밍, 모건 카우프만

Answer 3

를 참조하십시오. 술어를 전달한 요소의 수를 계산 (워프 당)하기 위해 스트림 압축의 접두사 합 부분에 사용됩니다.

는

Here the paper. In-k Stream Compaction