Xét đoạn mã sau:Có cách nào để đặt giá trị mặc định cho mảng bộ nhớ dùng chung không?
__global__ void kernel(int *something) {
extern __shared__ int shared_array[];
// Some operations on shared_array here.
}
Có thể thiết lập toàn bộ shared_array một số giá trị - ví dụ 0 - mà không giải quyết rõ ràng từng ô trong một chuỗi?
No. Bộ nhớ chia sẻ chưa được khởi tạo. Bạn cần phải bằng cách nào đó khởi nó cho mình, cách này hay cách khác ...
Từ CUDA C Programming Guide 3.2, Mục B.2.4.2, khoản 2:
__shared__biến không thể có một khởi như một phần tuyên bố của họ.
Điều này cũng loại bỏ các hàm tạo mặc định không cố định cho các biến được chia sẻ.
Bạn có hiệu quả có thể khởi tạo mảng chia sẻ song song như thế này
// if SHARED_SIZE == blockDim.x, eliminate this loop
for (int i = threadIdx.x; i < SHARED_SIZE; i += blockDim.x)
shared_array[i] = INITIAL_VALUE;
__syncthreads();
Vâng, bạn có thể. Bạn có thể xác định rằng các chủ đề đầu tiên trong khối đặt nó, trong khi người kia không Eg .:
extern __shared__ unsigned int local_bin[]; // Size specified in kernel call
if (threadIdx.x == 0) // Wipe on first thread - include " && threadIdx.y == 0" and " && threadIdx.z == 0" if threadblock has 2 or 3 dimensions instead of 1.
{
// For-loop to set all local_bin array indexes to specified value here - note you cannot use cudaMemset as it translates to a kernel call itself
}
// Do stuff unrelated to local_bin here
__syncthreads(); // To make sure the memset above has completed before other threads start writing values to local_bin.
// Do stuff to local_bin here
Lý tưởng nhất là bạn nên làm càng nhiều việc càng tốt trước khi syncthreads gọi, vì điều này cho phép tất cả các khác các chủ đề để thực hiện công việc của họ trước khi bộ nhớ hoàn tất - rõ ràng điều này chỉ quan trọng nếu tác phẩm có khả năng có thời gian hoàn thành luồng hoàn toàn khác nhau, ví dụ nếu có phân nhánh có điều kiện. Lưu ý rằng đối với chuỗi 0 "thiết lập" cho vòng lặp, bạn cần phải vượt qua kích thước của mảng local_bin làm tham số cho hạt nhân để bạn biết kích thước của mảng mà bạn đang lặp lại.
Cảm ơn bạn. Tôi đã sử dụng một cái gì đó tương tự như điều này trong việc thực hiện cuối cùng của tôi. – rayryeng
điều này mất lợi ích của việc song song, luôn luôn cố gắng sử dụng threadIdx và BlockIdx càng nhiều càng tốt – ejectamenta
Tôi nghĩ rằng bạn nói chung, nhưng nó phụ thuộc vào việc giá trị khởi tạo có thể được tính từ những giá trị đó hay không./etc mô hình phân công. – metamorphosis
Đó chỉ là trường hợp nếu bạn có khối 1D, tất nhiên. Nói chỉ vì vậy bất kỳ người mới nào không rơi vào bẫy rõ ràng. Tôi cũng tự hỏi làm thế nào tăng nhiều float4, mà là một thủ thuật vẫn còn đưa ra trên các thiết bị mới hơn và bao nhiêu lợi ích nó cung cấp kết hợp với loại bộ nhớ này phối hợp của init. Sidenote, nếu bạn đang tải bên trong một hạt nhân 2D, hoặc 3D, đó là importrant để biết rằng chúng được phân đoạn thành warps như mảng của [z] [y] [x] là bên trong bộ nhớ. Vì vậy, để cho [x] chủ đề viết gần nhất với nhau và những khác nhau trong [z] xa nhất của họ. –
Vì vậy, tôi đã thử và có, sử dụng reinterpret_cast để sao chép trong các đoạn 16 byte như trong https://devblogs.nvidia.com/parallelforall/cuda-pro-tip-increase-performance-with-vectorized-memory-access/vẫn nhận được kết quả tốt hơn một chút. Ngoài ra, điều quan trọng là phải đặt __restrict__ trên dữ liệu bạn muốn sao chép --- tăng hiệu suất 15% dễ dàng, khoảng chừng 16 byte khối liên kết cung cấp. –
Tâm trí mà bộ nhớ coalescing có thể có nghĩa là một sự khác biệt tốc độ 16x một cách dễ dàng. –