Sử dụng bộ nhớ liên tục trong mã CUDA

Question

Tôi không thể tự mình tìm ra, cách tốt nhất để đảm bảo bộ nhớ được sử dụng trong hạt nhân của tôi là hằng số. Có một câu hỏi tương tự tại số http://stackoverflow...r-pleasant-way. Tôi đang làm việc với GTX580 và chỉ biên dịch cho khả năng 2.0. hạt nhân của tôi trông giống như

__global__ Foo(const int *src, float *result) {...} 

tôi thực hiện đoạn mã sau trên máy chủ:

cudaMalloc(src, size); 
cudaMemcpy(src, hostSrc, size, cudaMemcpyHostToDevice); 
Foo<<<...>>>(src, result); 

cách khác là để thêm

__constant__ src[size]; 

để .cu tập tin, loại bỏ src con trỏ từ hạt nhân và thực hiện

cudaMemcpyToSymbol("src", hostSrc, size, 0, cudaMemcpyHostToDevice); 
Foo<<<...>>>(result); 

Hai cách này tương đương hay cách đầu tiên không đảm bảo việc sử dụng bộ nhớ liên tục thay vì bộ nhớ toàn cục? kích thước thay đổi động nên cách thứ hai không tiện dụng trong trường hợp của tôi.

Answer 1

Cách thứ hai là cách duy nhất để đảm bảo rằng mảng được biên dịch thành bộ nhớ không đổi CUDA và được truy cập chính xác qua bộ đệm bộ nhớ không đổi. Nhưng bạn nên tự hỏi làm thế nào các nội dung của mảng đó sẽ được truy cập trong một khối các chủ đề. Nếu mọi luồng sẽ truy cập vào mảng đồng đều, thì sẽ có lợi thế về hiệu năng khi sử dụng bộ nhớ liên tục, vì có cơ chế phát sóng từ bộ nhớ đệm liên tục (nó cũng tiết kiệm băng thông bộ nhớ toàn cục vì bộ nhớ liên tục được lưu trong DRAM offchip và cache giảm số lượng giao dịch DRAM). Nhưng nếu truy cập là ngẫu nhiên, sau đó có thể được serialization truy cập vào bộ nhớ địa phương mà sẽ ảnh hưởng tiêu cực hiệu suất.

Những điều điển hình có thể phù hợp với bộ nhớ __constant__ sẽ là hệ số mô hình, trọng số và các giá trị không đổi khác cần được đặt ở thời gian chạy. Trên GPU Fermi, danh sách đối số hạt nhân được lưu trữ trong bộ nhớ không đổi, ví dụ. Nhưng nếu các nội dung được truy cập không thống nhất và loại hoặc kích thước của các thành viên không phải là liên tục từ cuộc gọi đến cuộc gọi, thì bộ nhớ toàn cầu bình thường là thích hợp hơn.

Cũng nên nhớ rằng có giới hạn 64kb bộ nhớ không đổi trên mỗi ngữ cảnh GPU, vì vậy không thực tế để lưu trữ lượng dữ liệu rất lớn trong bộ nhớ không đổi. Nếu bạn cần nhiều bộ nhớ chỉ đọc với bộ nhớ cache, có thể bạn nên thử ràng buộc dữ liệu vào một kết cấu và xem hiệu suất là như thế nào. Trên các thẻ pre-Fermi, nó thường mang lại hiệu suất tiện lợi, trên Fermi, kết quả có thể ít có thể dự đoán được so với bộ nhớ toàn cầu vì cải thiện bố cục bộ nhớ cache trong kiến ​​trúc đó.

Answer 2

Phương pháp đầu tiên sẽ đảm bảo bộ nhớ không đổi bên trong chức năng Foo. Cả hai không tương đương, thứ hai đảm bảo nó tiếp tục ở đó sau khi nó khởi tạo. Nếu bạn cần năng động hơn bạn nee để sử dụng một cái gì đó tương tự như cách đầu tiên.