Sử dụng hằng số với CUDA

Question

Cách tốt nhất để sử dụng hằng số trong CUDA là gì?

Một cách là để xác định các hằng số trong bộ nhớ liên tục, như:

// CUDA global constants 
__constant__ int M; 

int main(void) 
{ 
    ... 
    cudaMemcpyToSymbol("M", &M, sizeof(M)); 
    ... 
} 

Một cách phai lạt sẽ được sử dụng C Preprocessor:

#define M = ... 

Tôi nghĩ rằng việc xác định các hằng số với tiền xử lý C nhanh hơn nhiều. Vậy thì lợi ích của việc sử dụng bộ nhớ không đổi trên thiết bị CUDA là gì?

Answer 1

1. hằng được biết tại thời gian biên dịch cần được xác định bằng macro tiền xử lý (ví dụ #define) hoặc qua C/C++ const biến ở phạm vi toàn cầu/file.
2. Sử dụng __constant__ memory có thể có lợi cho các chương trình sử dụng các giá trị nhất định không thay đổi trong thời gian hạt nhân và có một số mẫu truy cập nhất định (ví dụ: tất cả các chuỗi truy cập cùng một giá trị cùng một lúc). Điều này không tốt hơn hoặc nhanh hơn các hằng số thỏa mãn các yêu cầu của mục 1 ở trên.
3. Nếu số lựa chọn được thực hiện bởi một chương trình tương đối nhỏ về số lượng, và những lựa chọn này ảnh hưởng đến thực hiện hạt nhân, một cách tiếp cận có thể cho thêm tối ưu hóa thời gian biên dịch sẽ được sử dụng templated code/kernels

Answer 2

Regular C/Hằng số kiểu C++: Trong CUDA C (chính nó là một sửa đổi của C99) hằng số là các thực thể thời gian biên dịch tuyệt đối. Điều này là không đáng ngạc nhiên khi số lượng tối ưu hóa xảy ra trong NVCC là RẤT liên quan đến bản chất của xử lý GPU.

#define: macro luôn kém chất lượng nhưng hữu ích trong một nhúm.

Trình xác định biến số __constant__ là, tuy nhiên một động vật hoàn toàn mới và một cái gì đó của một từ ngữ sai theo ý kiến ​​của tôi. Tôi sẽ đặt xuống những gì Nvidia có here trong không gian dưới đây:

Các __constant__ vòng loại, tùy chọn sử dụng cùng với __device__, khai báo một biến:

• nằm trong không gian bộ nhớ liên tục,
• Có tuổi thọ của một ứng dụng,
• Có thể truy cập từ tất cả các chuỗi trong lưới và từ máy chủ thông qua thư viện thời gian chạy (cudaGetSymbolAddress()/ cudaGetSymbolSize()/cudaMemcpyToSymbol()/cudaMemcpyFromSymbol()).

tài liệu của Nvidia xác định rằng __constant__ có sẵn tại đăng ký mức tốc độ (gần như zero độ trễ) miễn là nó là hằng số tương tự được truy cập bởi tất cả các chủ đề của một warp.

Chúng được khai báo ở phạm vi toàn cầu trong mã CUDA. BAO GIỜ dựa trên kinh nghiệm cá nhân (và hiện đang diễn ra), bạn phải cẩn thận với specifier này khi nói đến việc biên dịch riêng biệt, như tách mã CUDA của bạn (.cu và.tập tin cuh) từ mã C/C++ của bạn bằng cách đặt các hàm bao bọc trong các tiêu đề kiểu C.

Không giống như các biến được chỉ định "cố định" truyền thống tuy nhiên chúng được khởi tạo khi chạy từ mã máy chủ cấp phát bộ nhớ thiết bị và cuối cùng khởi chạy hạt nhân. Tôi lặp lại tôi hiện đang làm việc mã chứng minh rằng này có thể được thiết lập tại thời gian chạy bằng cách sử dụng cudaMemcpyToSymbol() trước khi thực hiện hạt nhân.

Chúng khá tiện lợi khi nói rằng ít nhất là tốc độ cấp bộ nhớ cache L1 được đảm bảo để truy cập.