OpenCL - Tại sao sử dụng bộ đệm READ_ONLY hoặc WRITE_ONLY

Question

Trong OpenCL, có bất kỳ lợi ích hiệu suất nào để gắn cờ bộ đệm là READ_ONLY hoặc WRITE_ONLY?

kernel Đây là những gì tôi thường thấy (a là READ_ONLY và b được WRITE_ONLY):

__kernel void two_buffer_double(__global float* a, __global float* b) 
{ 
    int i = get_global_id(0); 
    b[i] = a[i] * 2; 
} 

kernel Điều này có vẻ tốt hơn, vì nó sử dụng ít bộ nhớ toàn cầu (a là READ_WRITE):

__kernel void one_buffer_double(__global float* a) 
{ 
    int i = get_global_id(0); 
    a[i] = a[i] * 2; 
} 

Làm READ_ONLY và WRITE_ONLY cờ chỉ tồn tại để giúp gỡ lỗi và bắt lỗi?

Answer 1

Để trả lời thẳng về phía trước cho câu hỏi của bạn tôi muốn nói: Không, những lá cờ không chỉ tồn tại nhằm giúp gỡ rối và lỗi đánh bắt. Tuy nhiên thật khó để đưa ra bất kỳ tham chiếu nào về cách những lá cờ này được sử dụng bởi bất kỳ việc thực hiện nào và chúng tác động như thế nào đến buổi biểu diễn.

Sự hiểu biết của tôi (rất tiếc là không được sao lưu bằng bất kỳ tài liệu nào), khi bạn sử dụng những lá cờ này, bạn sẽ có nhiều ràng buộc hơn về cách sử dụng bộ đệm và do đó bạn có thể trợ giúp thời gian chạy/trình điều khiển/trình biên dịch có thể cải thiện màn trình diễn. Ví dụ Tôi tưởng tượng rằng không nên lo lắng về tính nhất quán của bộ nhớ với bộ đệm chỉ đọc trong khi hạt nhân đang sử dụng nó vì các workitems không được phép ghi vào nó. Do đó một số kiểm tra có thể được bỏ qua ... mặc dù trong Opencl bạn đang giả sử để chăm sóc bản thân này bằng cách sử dụng các rào cản và như vậy.

Cũng lưu ý rằng kể từ khi Opencl 1.2 một số cờ khác đã được giới thiệu liên quan đến thời gian này để cách máy chủ cần truy cập bộ đệm. Có:

CL_MEM_HOST_NO_ACCESS, 
CL_MEM_HOST_{READ, WRITE}_ONLY, 
CL_MEM_{USE, ALLOC, COPY}_HOST_PTR 

Tôi đoán đó một lần nữa nó phải giúp người dân thực hiện OpenCL để nâng cao hiệu suất, nhưng tôi đoán chúng tôi cần sự đóng góp của một số chuyên gia AMD hoặc NVIDIA.

Xin lưu ý rằng tất cả những gì tôi đã nói cho đến nay là chỉ những suy nghĩ của tôi và không dựa trên bất kỳ tài liệu nghiêm trọng nào (Tôi không tìm thấy bất kỳ tài liệu nào).

Mặt khác, tôi có thể cho bạn biết chắc chắn rằng tiêu chuẩn không buộc bộ đệm chỉ đọc nằm trong không gian cố định như @Quonux đã nêu. Nó có thể là một số triển khai thực hiện điều này cho bộ đệm nhỏ. Đừng quên rằng bộ nhớ không gian liên tục nhỏ nên bạn có thể chỉ đọc bộ đệm quá lớn để vừa với.Cách duy nhất để đảm bảo rằng bộ đệm nằm trong bộ nhớ không gian liên tục là sử dụng từ khóa liên tục trong mã hạt nhân của bạn như được giải thích here. Tất nhiên ở phía máy chủ, nếu bạn muốn sử dụng bộ đệm liên tục, bạn phải sử dụng cờ chỉ đọc.

Answer 2

Nó phụ thuộc,

một vị trí bộ nhớ READ_ONLY __global được lưu trữ trong "Global/liên tục Bộ nhớ cache dữ liệu" mà là nhiều nhanh hơn so với bộ nhớ cache thông thường hoặc RAM trên một GPU (xem here), trên một CPU nó không quan trọng.

Tôi không biết bất kỳ lợi thế nào của WRITE_ONLY, có thể nó cũng giúp ích bởi vì GPU biết rằng nó có thể truyền dữ liệu ra khỏi nhu cầu lưu trữ bộ nhớ đệm.

Chỉ cần đi và đo nó nếu không chắc chắn của bạn ...

Answer 3

Lưu ý rằng thực tế có hai loại. Bạn có CL_MEM_READ_ONLY, CL_MEM_WRITE_ONLY và CL_MEM_READ_WRITE khi phân bổ bộ đệm nhưng sau đó bạn cũng có __read_only, __write_only và __read_write để trang trí con trỏ của bạn bằng mã hạt nhân.

Đây có thể được sử dụng cho cả việc tối ưu hóa và kiểm tra lỗi. Hãy xem hiệu suất trước. Nếu một bộ đệm chỉ ghi được gặp phải, việc ghi không cần phải được lưu trữ (như trong ghi thông qua bộ nhớ đệm), tiết kiệm bộ nhớ cache nhiều hơn cho các lần đọc. Điều này phụ thuộc vào phần cứng GPU rất nhiều và ít nhất là phần cứng NVIDIA có các hướng dẫn cần thiết để thực sự thực hiện điều này (các sửa đổi .cs và .lu). Bạn có thể tham khảo PTX ISA của họ. Tôi chưa thấy bất kỳ bằng chứng về trình biên dịch thực sự thực hiện tối ưu hóa này, ví dụ:

__kernel void Memset4(__global __write_only unsigned int *p_dest, 
    const unsigned int n_dword_num) 
{ 
    unsigned int i = get_global_id(0); 
    if(i < n_dword_num) 
     p_dest[i] = 0; // this 
} 

được biên soạn như:

st.global.u32 [%r10], %r11; // no cache operation specified 

này có ý nghĩa như CUDA không có tương đương đối với những vòng loại để trình biên dịch rất có thể âm thầm bỏ qua chúng. Nhưng nó không đau khi đặt chúng ở đó, chúng ta có thể may mắn hơn trong tương lai. Trong CUDA, một số chức năng này được tiếp xúc bằng cách sử dụng chức năng __ldg và bằng cách sử dụng cờ trình biên dịch để chọn tham gia/thoát khỏi bộ nhớ đệm chuyển toàn bộ bộ nhớ trong L1 (-Xptxas -dlcm=cg). Bạn cũng có thể luôn sử dụng asm nếu bạn thấy rằng bộ nhớ đệm bỏ qua mang lại lợi thế lớn.

Để kiểm tra lỗi, việc ghi vào bộ đệm chỉ đọc có thể tránh được bằng cách sử dụng thông số const trong khai báo hạt nhân. Không cho phép đọc từ bộ đệm chỉ ghi trong "C" thuần túy.

Tối ưu hóa khác có thể xảy ra khi ánh xạ các bộ đệm đó để lưu trữ bộ nhớ. Khi ánh xạ bộ đệm CL_MEM_READ_ONLY, vùng được ánh xạ có thể không được khởi tạo vì máy chủ sẽ chỉ ghi vào bộ nhớ đó, để thiết bị chỉ đọc nó. Tương tự, khi bỏ ánh xạ bộ đệm CL_MEM_WRITE_ONLY, trình điều khiển không cần phải sao chép nội dung (có thể được sửa đổi bởi máy chủ lưu trữ) từ bộ nhớ máy chủ sang bộ nhớ thiết bị. Tôi không đo lường điều này.

Như một mặt lưu ý, tôi đã cố gắng sử dụng:

inline unsigned int n_StreamingLoad(__global __read_only const unsigned int *p_src) 
{ 
#ifdef NVIDIA 
    unsigned int n_result; 
    asm("ld.global.cs.u32 %r0, [%r1];" : "=r" (n_result) : "r" (p_src)); 
    return n_result; 
#else // NVIDIA 
    return *p_src; // generic 
#endif // NVIDIA 
} 

inline void StreamingWrite(__global __write_only unsigned int *p_dest, const unsigned int n_value) 
{ 
#ifdef NVIDIA 
    asm("st.global.cs.u32 [%r0], %r1;" : : "r" (p_dest), "r" (n_value) : "memory"); 
#else // NVIDIA 
    *p_dest = n_value; // generic 
#endif // NVIDIA 
} 

mà cung cấp cho bạn khoảng 15 phụ GB/giây ngay cả trên một hạt nhân memcpy đơn giản với sm_35 thiết bị (thử nghiệm trên GTX 780 và K40). Đã không nhìn thấy tốc độ đáng chú ý trên sm_30 (không chắc chắn nếu nó thậm chí có nghĩa là để được hỗ trợ ở đó - mặc dù các hướng dẫn không bị tước khỏi ptx). Lưu ý rằng bạn cần phải tự xác định NVIDIA chính mình (hoặc xem Detect OpenCL device vendor in kernel code).