Rào cản liên khối trên CUDA

Question

Tôi muốn thực hiện rào cản Liên khối trên CUDA, nhưng gặp phải một vấn đề nghiêm trọng.

Tôi không thể hiểu tại sao nó không hoạt động.

#include <iostream> 
#include <cstdlib> 
#include <ctime> 

#define SIZE 10000000 
#define BLOCKS 100 

using namespace std; 

struct Barrier { 
    int *count; 

    __device__ void wait() { 
     atomicSub(count, 1); 
     while(*count) 
      ; 
    } 

    Barrier() { 
     int blocks = BLOCKS; 
     cudaMalloc((void**) &count, sizeof(int)); 
     cudaMemcpy(count, &blocks, sizeof(int), cudaMemcpyHostToDevice); 
    } 

    ~Barrier() { 
     cudaFree(count); 
    } 
}; 


__global__ void sum(int* vec, int* cache, int *sum, Barrier barrier) 
{ 
    int tid = blockIdx.x; 

    int temp = 0; 
    while(tid < SIZE) { 
     temp += vec[tid]; 
     tid += gridDim.x; 
    } 

    cache[blockIdx.x] = temp; 

    barrier.wait(); 

    if(blockIdx.x == 0) { 
     for(int i = 0 ; i < BLOCKS; ++i) 
      *sum += cache[i]; 
    } 
} 

int main() 
{ 
    int* vec_host = (int *) malloc(SIZE * sizeof(int));  
    for(int i = 0; i < SIZE; ++i) 
     vec_host[i] = 1; 

    int *vec_dev; 
    int *sum_dev; 
    int *cache; 
    int sum_gpu = 0; 

    cudaMalloc((void**) &vec_dev, SIZE * sizeof(int)); 
    cudaMemcpy(vec_dev, vec_host, SIZE * sizeof(int), cudaMemcpyHostToDevice); 
    cudaMalloc((void**) &sum_dev, sizeof(int)); 
    cudaMemcpy(sum_dev, &sum_gpu, sizeof(int), cudaMemcpyHostToDevice); 
    cudaMalloc((void**) &cache, BLOCKS * sizeof(int)); 
    cudaMemset(cache, 0, BLOCKS * sizeof(int)); 

    Barrier barrier; 
    sum<<<BLOCKS, 1>>>(vec_dev, cache, sum_dev, barrier); 

    cudaMemcpy(&sum_gpu, sum_dev, sizeof(int), cudaMemcpyDeviceToHost); 

    cudaFree(vec_dev); 
    cudaFree(sum_dev); 
    cudaFree(cache); 
    free(vec_host); 
    return 0; 
} 

Trong thực tế, ngay cả khi tôi viết lại wait() như sau

__device__ void wait() { 
     while(*count != 234124) 
      ; 
    } 

Chương trình thoát bình thường. Nhưng tôi mong đợi để có được một vòng lặp vô hạn trong trường hợp này.

Answer 1

Thật không may, những gì bạn muốn đạt được (liên lạc khối/đồng bộ hóa) không hoàn toàn có thể xảy ra trong CUDA. Hướng dẫn lập trình CUDA nói rằng "các khối luồng được yêu cầu thực hiện độc lập: Nó phải có khả năng thực hiện chúng theo bất kỳ thứ tự nào, song song hoặc theo chuỗi." Lý do cho sự hạn chế này là cho phép sự linh hoạt trong bộ lập lịch khối chuỗi và để cho phép mã quy mô theo tỉ lệ với số lượng lõi. Phương thức đồng bộ liên khối duy nhất được hỗ trợ là khởi chạy một hạt nhân khác: khởi chạy hạt nhân (trong cùng một luồng) là các điểm đồng bộ ngầm định.

Mã của bạn vi phạm quy tắc độc lập khối vì nó ngầm giả định rằng các luồng chuỗi hạt nhân của bạn thực thi đồng thời (song song). Nhưng không có gì đảm bảo rằng họ làm. Để xem tại sao điều này quan trọng với mã của bạn, hãy xem xét một GPU giả định chỉ với một lõi. Chúng tôi cũng giả định rằng bạn chỉ muốn khởi chạy hai khối chuỗi. Hạt nhân spinloop của bạn sẽ thực sự bế tắc trong tình huống này. Nếu khối chuỗi 0 được lên lịch trên lõi đầu tiên, nó sẽ lặp lại mãi mãi khi nó tới hàng rào, bởi vì khối chuỗi không bao giờ có cơ hội cập nhật bộ đếm. Bởi vì khối chuỗi 0 không bao giờ được hoán đổi (các khối luồng thực thi để hoàn thành) nó chặn luồng của một trong các lõi trong khi nó quay.

Một số người đã thử các đề án như của bạn và đã thấy thành công vì người lên lịch đã tình cờ lập lịch biểu các khối theo cách mà các giả định đã làm việc. Ví dụ, có một thời gian khi tung ra nhiều khối chuỗi như một GPU có SM có nghĩa là các khối thực sự được thực hiện đồng thời. Nhưng họ đã thất vọng khi thay đổi trình điều khiển hoặc thời gian chạy CUDA hoặc GPU đã vô hiệu hóa giả định đó, phá vỡ mã của họ.

Đối với ứng dụng của bạn, hãy thử tìm một giải pháp không phụ thuộc vào đồng bộ liên khối, vì (chặn một thay đổi ký hiệu thành mô hình lập trình CUDA) nó không thể thực hiện được.

Answer 2

Dường như vấn đề tối ưu hóa trình biên dịch. Tôi không tốt với đọc PTX-mã, nhưng có vẻ như trình biên dịch đã bỏ qua các while -loop ở tất cả (ngay cả khi biên soạn với -O0):

.loc 3 41 0 
cvt.u64.u32  %rd7, %ctaid.x; // Save blockIdx.x to rd7 
ld.param.u64 %rd8, [__cudaparm__Z3sumPiS_S_7Barrier_cache]; 
mov.s32  %r8, %ctaid.x; // Now calculate ouput address 
mul.wide.u32 %rd9, %r8, 4; 
add.u64  %rd10, %rd8, %rd9; 
st.global.s32 [%rd10+0], %r5; // Store result to cache[blockIdx.x] 
.loc 17 128 0 
ld.param.u64 %rd11, [__cudaparm__Z3sumPiS_S_7Barrier_barrier+0]; // Get *count to rd11 
mov.s32  %r9, -1; // put -1 to r9 
atom.global.add.s32  %r10, [%rd11], %r9; // Do AtomicSub, storing the result to r10 (will be unused) 
cvt.u32.u64  %r11, %rd7; // Put blockIdx.x saved in rd7 to r11 
mov.u32  %r12, 0; // Put 0 to r12 
setp.ne.u32  %p3, %r11, %r12; // if(blockIdx.x == 0) 
@%p3 bra $Lt_0_5122; 
ld.param.u64 %rd12, [__cudaparm__Z3sumPiS_S_7Barrier_sum]; 
ld.global.s32 %r13, [%rd12+0]; 
mov.s64  %rd13, %rd8; 
mov.s32  %r14, 0; 

Trong trường hợp của mã CPU, hành vi như vậy được ngăn chặn bằng cách khai báo biến với tiền tố volatile. Nhưng ngay cả nếu chúng ta khai báo count như int __device__ count (và thích hợp thay đổi mã), thêm volatile specifier chỉ phá vỡ biên soạn (với các lỗi Loke argument of type "volatile int *" is incompatible with parameter of type "void *")

Tôi đề nghị xem xét threadFenceReduction ví dụ từ CUDA SDK. Ở đó họ đang làm khá giống như bạn làm, nhưng khối để thực hiện tổng kết cuối cùng được chọn trong thời gian chạy, thay vì được xác định trước và vòng loại while bị loại bỏ, vì spin-lock trên biến toàn cầu phải là rất chậm.

Answer 3

Có thể chặn chặn đồng bộ hóa. Xem này paper.
Bài báo không đi sâu vào chi tiết về cách hoạt động của nó, nhưng nó dựa vào hoạt động của __syncthreads(); để tạo khoảng dừng tạm thời cho khối hiện tại, ... trong khi đợi các khối khác để đến điểm đồng bộ.

Một mục không được chú ý trong bài báo là chỉ có thể đồng bộ hóa nếu số lượng khối đủ nhỏ hoặc số lượng SM đủ lớn cho tác vụ trên tay. tức là nếu bạn có 4 SM và đang cố gắng đồng bộ 5 khối, .. hạt nhân sẽ bế tắc.

Với cách tiếp cận của họ, tôi đã có thể truyền một tác vụ nối tiếp dài giữa nhiều khối, dễ dàng tiết kiệm 30% thời gian qua một phương pháp chặn duy nhất. tức là khối đồng bộ hóa đã hoạt động đối với tôi.