CUDA chỉ tắt bộ đệm L1 cho một biến

Question

Có cách nào trên các thiết bị CUDA 2.0 để vô hiệu hóa bộ đệm L1 chỉ cho một biến cụ thể không? Tôi biết rằng người ta có thể vô hiệu hóa bộ nhớ cache L1 tại thời gian biên dịch thêm cờ -Xptxas -dlcm=cg đến nvcc cho tất cả các hoạt động bộ nhớ. Tuy nhiên, tôi muốn vô hiệu hóa bộ nhớ cache chỉ cho bộ nhớ đọc trên một biến toàn cầu cụ thể để tất cả phần còn lại của bộ nhớ đọc để đi qua bộ đệm L1.

Dựa trên tìm kiếm tôi đã thực hiện trên web, giải pháp có thể là thông qua mã lắp ráp PTX.

Answer 1

Như đã đề cập ở trên, bạn có thể sử dụng inline PTX, đây là một ví dụ:

__device__ __inline__ double ld_gbl_cg(const double *addr) { 
    double return_value; 
    asm("ld.global.cg.f64 %0, [%1];" : "=d"(return_value) : "l"(addr)); 
    return return_value; 
} 

Bạn có thể dễ dàng thay đổi này bằng cách trao đổi .f64 cho .f32 (float) hoặc .s32 (int) vv, ràng buộc của return_value "= d" cho "= f" (phao) hoặc "= r" (int) etc. Lưu ý rằng ràng buộc cuối cùng trước (addr) - "l" - biểu thị địa chỉ 64 bit, nếu bạn đang sử dụng địa chỉ 32 bit, nó phải là "r".

Answer 2

PTX nội tuyến có thể được sử dụng để tải và lưu trữ biến. Lệnh ld.cg và st.cg chỉ lưu dữ liệu trong L2. Các toán tử bộ nhớ đệm được mô tả trong phần 8.7.8.1 Các toán tử bộ nhớ cache của tài liệu PTX ISA 2.3. Hướng dẫn hoặc sở thích là ld và st. Nội tuyến PTX được mô tả trong Using Inline PTX Assembly in CUDA.

Answer 3

Nếu bạn khai báo biến là volatile, thì nó sẽ chỉ được lưu trong bộ nhớ cache L2 trên GPU Fermi. Lưu ý rằng một số tối ưu hóa trình biên dịch, chẳng hạn như loại bỏ các tải lặp lại, không được thực hiện trên các biến dễ bay hơi vì trình biên dịch giả sử chúng có thể được viết bởi một luồng khác.