Cách thích hợp để thông báo hạt nhân OpenCL của nhiều đối tượng bộ nhớ?

Question

Trong chương trình OpenCL của tôi, tôi sẽ kết thúc với 60+ bộ nhớ đệm toàn cầu mà mỗi hạt nhân sẽ cần để có thể truy cập. Cách được khuyến nghị để cho mỗi hạt nhân biết vị trí của từng bộ đệm này là gì?

Bản thân bộ đệm ổn định trong suốt thời gian của ứng dụng - tức là chúng tôi sẽ cấp phát bộ đệm khi bắt đầu ứng dụng, gọi nhiều hạt nhân, sau đó chỉ deallocate bộ đệm ở cuối ứng dụng. Tuy nhiên, nội dung của chúng có thể thay đổi khi các hạt đọc/ghi từ chúng.

Trong CUDA, cách tôi đã làm điều này là tạo ra 60 biến phạm vi toàn cầu của chương trình trong mã CUDA của tôi. Sau đó, trên máy chủ lưu trữ, hãy ghi địa chỉ của bộ đệm thiết bị mà tôi đã phân bổ vào các biến toàn cục này. Sau đó, nhân sẽ đơn giản sử dụng các biến toàn cục này để tìm bộ đệm cần thiết để làm việc.

Cách tốt nhất để thực hiện điều này trong OpenCL là gì? Có vẻ như các biến toàn cục của CL khác một chút so với CUDA, nhưng tôi không thể tìm ra câu trả lời rõ ràng nếu phương pháp CUDA của tôi hoạt động, và nếu có, làm thế nào để chuyển các con trỏ đệm thành các biến toàn cầu. Nếu điều đó không hoạt động, cách tốt nhất là gì?

Answer 1

60 biến toàn cầu chắc chắn là rất nhiều! Bạn có chắc chắn không có cách nào để tái cấu trúc thuật toán của bạn một chút để sử dụng các khối dữ liệu nhỏ hơn không? Hãy nhớ rằng, mỗi hạt nhân phải là một đơn vị làm việc tối thiểu, không phải là một cái gì đó khổng lồ!

Tuy nhiên, có một giải pháp khả thi. Giả sử 60 mảng của bạn có kích thước đã biết, bạn có thể lưu trữ tất cả chúng vào một bộ đệm lớn, và sau đó sử dụng bù đắp để truy cập các phần khác nhau của mảng lớn đó. Dưới đây là một ví dụ rất đơn giản với ba mảng:

A is 100 elements 
B is 200 elements 
C is 100 elements 

big_array = A[0:100] B[0:200] C[0:100] 
offsets = [0, 100, 300] 

Sau đó, bạn chỉ cần chuyển big_array và bù đắp cho hạt nhân, và bạn có thể truy cập từng mảng. Ví dụ:

A[50] = big_array[offsets[0] + 50] 
B[20] = big_array[offsets[1] + 20] 
C[0] = big_array[offsets[2] + 0] 

Tôi không biết điều này sẽ ảnh hưởng đến bộ nhớ đệm trên thiết bị cụ thể của bạn như thế nào, nhưng dự đoán ban đầu của tôi là "không tốt". Kiểu truy cập mảng này cũng hơi khó chịu. Tôi không chắc liệu nó có hợp lệ không, nhưng bạn có thể bắt đầu từng hạt nhân của bạn với một số mã trích xuất từng phần bù và thêm nó vào một bản sao của con trỏ ban đầu. Ở phía máy chủ, để giữ cho các mảng của bạn dễ tiếp cận hơn, bạn có thể sử dụng clCreateSubBuffer: http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateSubBuffer.html, nó cũng cho phép bạn chuyển các tham chiếu đến các mảng cụ thể mà không có mảng offset.

Tôi không nghĩ rằng giải pháp này sẽ tốt hơn so với truyền 60 đối số hạt nhân, nhưng tùy thuộc vào clSetKernelArgs của OpenCL, nó có thể nhanh hơn. Nó chắc chắn sẽ giảm độ dài của danh sách đối số của bạn.

Answer 2

Bạn cần làm hai việc. Thứ nhất, mỗi hạt nhân có sử dụng mỗi bộ nhớ đệm toàn cầu nên tuyên bố một cuộc tranh cãi cho mỗi một, một cái gì đó như thế này:

kernel void awesome_parallel_stuff(global float* buf1, ..., global float* buf60) 

để mỗi đệm sử dụng cho hạt nhân được liệt kê. Và sau đó, trên máy chủ máy chủ, bạn cần phải tạo từng bộ đệm và sử dụng clSetKernelArg để đính kèm bộ đệm đã cho vào đối số hạt nhân đã cho trước khi gọi clEnqueueNDRangeKernel để bắt đầu nhóm.

Lưu ý rằng nếu hạt nhân sẽ tiếp tục sử dụng cùng một bộ đệm với mỗi lần thực thi hạt nhân, bạn chỉ cần thiết lập đối số hạt nhân một thời gian. Một lỗi phổ biến tôi thấy, có thể chảy máu hiệu suất phía máy chủ, là liên tục gọi clSetKernelArg trong các tình huống mà nó hoàn toàn không cần thiết.