Khi nào sử dụng dễ bay hơi với Bộ nhớ CUDA chia sẻ

Question

Trong trường hợp nào bạn nên sử dụng từ khóa volatile với bộ nhớ chia sẻ của hạt nhân CUDA? Tôi hiểu rằng volatile nói với trình biên dịch không bao giờ để cache bất kỳ giá trị, nhưng câu hỏi của tôi là về hành vi với một mảng chia sẻ:

__shared__ float products[THREADS_PER_ACTION]; 

// some computation 
products[threadIdx.x] = localSum; 

// wait for everyone to finish their computation 
__syncthreads(); 

// then a (basic, ugly) reduction: 
if (threadIdx.x == 0) { 
    float globalSum = 0.0f; 
    for (i = 0; i < THREADS_PER_ACTION; i++) 
     globalSum += products[i]; 
} 

Tôi có cần products là dễ bay hơi trong trường hợp này? Mỗi mục nhập mảng chỉ được truy cập bởi một chuỗi duy nhất, ngoại trừ ở cuối, ở đó mọi thứ được đọc theo chủ đề 0. Có thể trình biên dịch có thể lưu toàn bộ mảng hay không và vì vậy tôi cần nó là volatile hoặc bộ nhớ cache chỉ yếu tố?

Cảm ơn!

Answer 1

Nếu bạn không khai báo mảng chia sẻ là volatile, thì trình biên dịch sẽ tự do tối ưu hóa vị trí trong bộ nhớ chia sẻ bằng cách định vị chúng trong sổ đăng ký (có phạm vi cụ thể cho một chuỗi), cho bất kỳ chuỗi nào . Điều này đúng cho dù bạn truy cập vào phần tử được chia sẻ cụ thể đó chỉ từ một luồng hay không. Do đó, nếu bạn sử dụng bộ nhớ dùng chung làm phương tiện liên lạc giữa các luồng của một khối, tốt nhất là khai báo nó volatile. Rõ ràng là mỗi luồng chỉ truy cập vào các phần tử chia sẻ của chính nó, và không bao giờ được kết nối với một luồng khác, thì điều này không quan trọng, và tối ưu hóa trình biên dịch sẽ không phá vỡ bất cứ thứ gì. Trong trường hợp của bạn, nơi bạn có một phần mã nơi mỗi chuỗi đang truy cập vào các phần tử riêng của bộ nhớ dùng chung và truy cập chỉ luồng chỉ xảy ra ở vị trí được hiểu rõ, bạn có thể sử dụng memory fence function để buộc trình biên dịch để loại bỏ bất kỳ giá trị nào được lưu trữ tạm thời trong thanh ghi, trở lại mảng được chia sẻ. Vì vậy, bạn có thể nghĩ rằng __threadfence_block() có thể hữu ích, nhưng trong trường hợp của bạn, __syncthreads()already has memory-fencing functionality built in. Vì vậy, cuộc gọi __syncthreads() của bạn đủ để buộc đồng bộ hóa luồng cũng như buộc bất kỳ giá trị nào được ghi vào bộ nhớ cache trong bộ nhớ dùng chung được gỡ bỏ trở lại bộ nhớ dùng chung.

Bằng cách này, nếu việc giảm ở cuối mã của bạn có liên quan đến hiệu suất, bạn có thể xem xét sử dụng phương pháp giảm song song để tăng tốc.