Truy cập bộ nhớ toàn cầu không được hợp nhất do truy cập gián tiếp trong CUDA

Question

Chương trình CUDA của tôi đang bị truy cập bộ nhớ toàn cầu chưa được kết hợp. Mặc dù thread idx-th chỉ đối phó với ô [idx] -throng một mảng, có nhiều truy cập bộ nhớ gián tiếp như được hiển thị bên dưới.

int idx=blockDim.x*blockIdx.x+threadIdx.x; 

.... = FF[m_front[m_fside[idx]]]; 

Đối m_fisde [idx], chúng tôi đã kết hợp từ các truy cập, nhưng những gì chúng tôi thực sự cần là FF [m_front [m_fside [idx]]]. Có một truy cập gián tiếp hai cấp.

Tôi đã cố gắng tìm một số mẫu dữ liệu trong m_front hoặc m_fsied để làm cho điều này trở thành truy cập tuần tự trực tiếp, nhưng phát hiện ra rằng chúng gần như 'ngẫu nhiên'.

Có cách nào để giải quyết vấn đề này không?

Answer 1

Đẩy mạnh bộ nhớ truy cập ngẫu nhiên toàn cầu: làm vô hiệu các dòng bộ nhớ cache L1

Fermi và Kepler kiến ​​trúc hỗ trợ hai loại tải từ bộ nhớ toàn cầu. Bộ nhớ đệm đầy đủ là chế độ mặc định , nó cố gắng nhấn vào L1, sau đó L2, sau đó GMEM và độ chi tiết tải là dòng 128 byte. L2-only cố gắng nhấn trong L2, sau đó GMEM và độ chi tiết tải là 32 byte. Đối với một số mẫu truy cập ngẫu nhiên nhất định, hiệu quả bộ nhớ có thể tăng lên bằng cách làm mất hiệu lực L1 và khai thác độ chi tiết thấp hơn của L2. Điều này có thể được thực hiện bằng cách biên dịch với –Xptxas –dlcm=cg tùy chọn nvcc.

Hướng dẫn chung cho tăng tốc truy cập bộ nhớ toàn cầu: vô hiệu hóa hỗ trợ ECC

Fermi và Kepler GPU hỗ trợ Lỗi Correcting Code (ECC), và ECC được kích hoạt theo mặc định. ECC giảm băng thông bộ nhớ đỉnh và được yêu cầu tăng cường tính toàn vẹn dữ liệu trong các ứng dụng như chụp ảnh y tế và tính toán cụm lớn. Nếu không cần thiết, nó có thể bị tắt để cải thiện hiệu suất bằng tiện ích nvidia-smi trên Linux (xem link) hoặc qua Bảng điều khiển trên hệ thống Microsoft Windows. Lưu ý rằng bật hoặc tắt ECC yêu cầu khởi động lại để có hiệu lực.

Hướng dẫn chung cho tăng tốc truy cập bộ nhớ toàn cầu trên Kepler: bằng cách sử dụng chỉ đọc dữ liệu bộ nhớ cache

Kepler có bộ nhớ cache 48KB cho dữ liệu được biết đến là read-only cho thời gian của hàm. Việc sử dụng đường dẫn chỉ đọc là có lợi vì nó giảm tải đường dẫn bộ nhớ cache chia sẻ/L1 và nó hỗ trợ truy cập bộ nhớ không được ký hiệu một cách nhanh chóng. Việc sử dụng đường dẫn chỉ đọc có thể được trình biên dịch tự động quản lý (sử dụng từ khóa const __restrict) hoặc rõ ràng (sử dụng __ldg() nội tại) của trình lập trình .