Tính toán Multi-GPU Cuda

Question

Tôi là người mới trong lập trình đa CPU và tôi có một số câu hỏi về tính toán đa GPU. Ví dụ, hãy lấy ví dụ về sản phẩm chấm. Tôi đang chạy một CPU-thread tạo ra 2 mảng lớn A [N] và B [N]. Do kích thước của các mảng này, tôi cần phải chia tính toán sản phẩm chấm của chúng thành 2 GPU, cả Tesla M2050 (khả năng tính toán 2.0). Vấn đề là tôi cần phải tính toán các sản phẩm này nhiều lần trong một vòng lặp do điều khiển bởi luồng CPU của tôi. Mỗi sản phẩm chấm yêu cầu kết quả của sản phẩm trước đó. Tôi đã đọc về việc tạo ra 2 chủ đề khác nhau để kiểm soát 2 GPU khác nhau một cách riêng biệt (như được mô tả trên cuda bằng ví dụ) nhưng tôi không có đầu mối về cách đồng bộ hóa và trao đổi dữ liệu giữa chúng. Có cách nào khác không? Tôi thực sự đánh giá cao bất kỳ loại trợ giúp/example.Thanks trước!

Answer 1

Trước khi CUDA 4.0, lập trình đa GPU yêu cầu lập trình CPU đa luồng. Điều này có thể khó khăn nhất là khi bạn cần đồng bộ hóa và/hoặc giao tiếp giữa các chủ đề hoặc GPU. Và nếu tất cả sự song song của bạn nằm trong mã GPU của bạn, thì việc có nhiều luồng CPU có thể làm tăng thêm sự phức tạp của phần mềm mà không cải thiện hiệu suất vượt xa những gì GPU làm.

Vì vậy, bắt đầu với CUDA 4.0, bạn có thể dễ dàng lập trình nhiều GPU từ một chương trình máy chủ đơn luồng. Here are some slides I presented last year about this.

trình nhiều GPU có thể đơn giản như thế này:

int numDevs = 0; 
cudaGetNumDevices(&numDevs); 
... 
for (int d = 0; d < numDevs; d++) { 
    cudaSetDevice(d); 
    kernel<<<blocks, threads>>>(args); 
} 

Ví dụ cụ thể của sản phẩm dấu chấm, bạn có thể sử dụng thrust::inner_product như là một điểm khởi đầu. Tôi sẽ làm điều đó để tạo mẫu. Nhưng hãy xem nhận xét của tôi ở phần cuối về tắc nghẽn băng thông.

Vì bạn không cung cấp đủ chi tiết về vòng lặp bên ngoài chạy các sản phẩm chấm nhiều lần, tôi không cố gắng làm bất cứ điều gì với điều đó.

// assume the deviceIDs of the two 2050s are dev0 and dev1. 
// assume that the whole vector for the dot product is on the host in h_data 
// assume that n is the number of elements in h_vecA and h_vecB. 

int numDevs = 0; 
cudaGetNumDevices(&numDevs); 
... 
float result = 0.f; 
for (int d = 0; d < numDevs; d++) { 
    cudaSetDevice(d); 
    device_vector<float> vecA(h_vecA + d*(n/d), h_vecA + (d+1)*(n/d)-1); 
    device_vector<float> vecB(h_vecB + d*(n/d), h_vecB + (d+1)*(n/d)-1); 
    result += thrust::inner_product(vecA.begin(), vecA.end(), vecB.begin(), 0.f); 
} 

(Tôi thừa nhận rằng việc lập chỉ mục trên là không đúng nếu n không là bội chẵn các numDevs, nhưng tôi sẽ để sửa chữa mà như một bài tập cho người đọc. :)

Đây là đơn giản và là một khởi đầu tuyệt vời. Làm cho nó hoạt động trước, sau đó tối ưu hóa.

Khi bạn đã làm việc, nếu tất cả những gì bạn đang làm trên thiết bị là sản phẩm chấm, bạn sẽ thấy mình bị ràng buộc băng thông - chủ yếu là bởi PCI-e, và bạn cũng sẽ không nhận được sự tương tranh giữa các thiết bị vì lực đẩy :: inner_product là đồng bộ do đọc lại để trả về kết quả ..Vì vậy, bạn có thể sử dụng cudaMemcpyAsync (các nhà xây dựng device_vector sẽ sử dụng cudaMemcpy). Nhưng cách tiếp cận dễ dàng và có khả năng hiệu quả hơn sẽ là sử dụng "zero copy" - truy cập trực tiếp vào bộ nhớ máy chủ (cũng được thảo luận trong bản trình bày đa-gpu được liên kết ở trên). Vì tất cả những gì bạn đang làm là đọc từng giá trị một lần và thêm nó vào tổng (việc tái sử dụng song song xảy ra trong bản sao bộ nhớ chia sẻ), bạn cũng có thể đọc nó trực tiếp từ máy chủ thay vì sao chép từ máy chủ sang thiết bị, và sau đó đọc nó từ bộ nhớ thiết bị trong hạt nhân. Ngoài ra, bạn sẽ muốn khởi chạy hạt nhân không đồng bộ trên mỗi GPU, để đảm bảo đồng thời tối đa.

Bạn có thể làm một cái gì đó như thế này:

int bytes = sizeof(float) * n; 
cudaHostAlloc(h_vecA, bytes, cudaHostAllocMapped | cudaHostAllocPortable); 
cudaHostAlloc(h_vecB, bytes, cudaHostAllocMapped | cudaHostAllocPortable); 
cudaHostAlloc(results, numDevs * sizeof(float), cudaHostAllocMapped | cudaHostAllocPortable); 
// ... then fill your input arrays h_vecA and h_vecB 


for (int d = 0; d < numDevs; d++) { 
    cudaSetDevice(d); 
    cudaEventCreate(event[d])); 
    cudaHostGetDevicePointer(&dptrsA[d], h_vecA, 0); 
    cudaHostGetDevicePointer(&dptrsB[d], h_vecB, 0); 
    cudaHostGetDevicePointer(&dresults[d], results, 0); 
} 

... 

for (int d = 0; d < numDevs; d++) { 
    cudaSetDevice(d); 
    int first = d * (n/d); 
    int last = (d+1)*(n/d)-1; 
    my_inner_product<<<grid, block>>>(&dresults[d], 
             vecA+first, 
             vecA+last, 
             vecB+first, 0.f); 
    cudaEventRecord(event[d], 0); 
} 

// wait for all devices 
float total = 0.0f; 
for (int d = 0; d < devs; d++) { 
    cudaEventSynchronize(event[d]); 
    total += results[numDevs]; 
} 

Answer 2

Để tạo một số chủ đề, bạn có thể sử dụng OpenMP hoặc pthreads. Để làm những gì bạn đang nói, có vẻ như bạn sẽ cần tạo và khởi chạy hai luồng (phần song song omp, hoặc pthread_create), mỗi phần thực hiện một phần của tính toán và lưu trữ kết quả trung gian của nó trong các biến riêng biệt của process-wIDE (nhớ lại, các biến toàn cầu được tự động chia sẻ giữa các luồng của một tiến trình, do đó, chuỗi ban đầu sẽ có thể thấy các thay đổi được thực hiện bởi hai luồng được sinh ra). Để có được các chủ đề ban đầu chờ cho những người khác hoàn thành, đồng bộ hóa (sử dụng một hàng rào toàn cầu hoặc hoạt động nối luồng) và kết hợp các kết quả trong chuỗi ban đầu sau khi hai luồng được sinh ra hoàn thành (nếu bạn tách mảng thành một nửa và tính toán sản phẩm dấu chấm bằng cách nhân các phần tử tương ứng và thực hiện giảm tổng kết toàn cục trên nửa, chỉ cần thêm hai kết quả trung gian từ hai chuỗi được sinh ra).

Bạn cũng có thể sử dụng MPI hoặc ngã ba, trong trường hợp giao tiếp có thể được thực hiện theo cách tương tự với lập trình mạng ... đường ống/ổ cắm hoặc thông tin liên lạc và đồng bộ hóa (chặn) gửi và nhận.