Giới thiệu chức năng bỏ phiếu dọc

Question

Hướng dẫn lập trình CUDA đã giới thiệu khái niệm chức năng bỏ phiếu dọc, "_ tất cả", " _any" và "__ballot".

Câu hỏi của tôi là: ứng dụng nào sẽ sử dụng 3 chức năng này?

Answer 1

__ballot được sử dụng trong CUDA-histogram và trong thư viện NPP CUDA để tạo nhanh bitmask và kết hợp với __popc nội tại để thực hiện hiệu quả giảm boolean rất hiệu quả.

__all và __any được sử dụng để giảm trước khi giới thiệu __ballot, mặc dù tôi không thể nghĩ ra bất kỳ cách sử dụng nào khác của chúng.

Answer 2

Nguyên mẫu của __ballot là sau

unsigned int __ballot(int predicate); 

Nếu predicate là nonzero, __ballot trả về một giá trị với các thiết lập N chút thứ, nơi N là chỉ số chủ đề.

Kết hợp với atomicOr và __popc, nó có thể được sử dụng để tích lũy số lượng chủ đề trong mỗi sợi dọc có vị từ đúng.

Trên thực tế, nguyên mẫu của atomicOr là

int atomicOr(int* address, int val); 

và atomicOr đọc giá trị được trỏ đến bởi address, thực hiện một bitwise OR tác với val, và ghi giá trị trở lại address và trả về giá trị cũ của nó như là một tham số trả về.

Ở phía bên kia, __popc trả về số bit được đặt với tham số 32-bit.

Theo đó, các hướng dẫn

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])); 

có thể được sử dụng để đếm số lượng các chủ đề mà các vị là đúng.

Để biết thêm chi tiết, xem Shane Cook, CUDA Lập trình, Morgan Kaufmann

Answer 3

Như một ví dụ về thuật toán có sử dụng API __ballot tôi sẽ đề cập đến In-Kernel Suối Đầm bởi D.M Hughes et Al. Nó được sử dụng trong tiền tố tổng hợp một phần của dòng đầm nén để đếm (mỗi sợi dọc) số lượng các phần tử đã chuyển vị từ.

Here the paper. In-k Stream Compaction