Giảm song song

Question

Tôi đã đọc bài viết Tối ưu hóa giảm song song trong CUDA bởi Mark Harris, và tôi thấy nó thực sự rất hữu ích, nhưng tôi vẫn thỉnh thoảng không thể hiểu được 1 hoặc 2 khái niệm. Nó được viết trên pg 18:

//First add during load 

// each thread loads one element from global to shared mem 

unsigned int tid = threadIdx.x; 

unsigned int i = blockIdx.x*blockDim.x + threadIdx.x; 

sdata[tid] = g_idata[i]; 
__syncthreads(); 

Tối ưu hóa Code: Với 2 tải và add 1 của giảm:

// perform first level of reduction, 

// reading from global memory, writing to shared memory 
unsigned int tid = threadIdx.x;         ...1 

unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;   ...2 

sdata[tid] = g_idata[i] + g_idata[i+blockDim.x];     ...3 

__syncthreads();             ...4 

Tôi không thể hiểu được dòng 2; nếu tôi có 256 phần tử, và nếu tôi chọn 128 là khối của tôi, thì tại sao tôi nhân nó với 2? Xin giải thích làm thế nào để xác định các khối?

Answer 1

Nó về cơ bản thực hiện hoạt động cho thấy trong hình dưới đây:

Mã này về cơ bản là "nói" rằng một nửa số chủ đề sẽ thực hiện việc đọc từ bộ nhớ toàn cầu và văn bản cho bộ nhớ chia sẻ, như thể hiện trong hình.

Bạn thực thi Kernell và bây giờ bạn muốn giảm một số giá trị, bạn giới hạn quyền truy cập vào mã ở trên chỉ bằng một nửa tổng số chủ đề đang chạy. Tưởng tượng bạn có 4 khối, mỗi một với 512 bài, bạn giới hạn các mã trên chỉ được thực hiện bởi hai khối đầu tiên, và bạn có một g_idate [4 * 512]:

unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x; 

sdata[tid] = g_idata[i] + g_idata[i+blockDim.x]; 

Vì vậy:

thread 0 of block = 0 will copy the position 0 and 512, 
thread 1 of block = 0 position 1 and 513; 
thread 511 of block = 0 position 511 and 1023; 
thread 0 of block 1 position 1024 and 1536 
thread 511 of block = 1 position 1535 and 2047 

BlockDim.x * 2 được sử dụng bởi vì mỗi chuỗi sẽ truy cập vào vị trí i và i+blockDim.x để bạn có nhiều để đảm bảo rằng các chuỗi trên khối id tiếp theo không trùng lặp vị trí giống nhau của g_idata đã được tính toán.

Answer 2

Trong mã được tối ưu hóa, bạn chạy hạt nhân với các khối lớn bằng một nửa so với cách triển khai không được tối ưu hóa.

Hãy gọi kích thước của khối trong mã không được tối ưu hóa work, để một nửa kích thước này được gọi là unit và để các kích thước này có cùng giá trị số cho mã được tối ưu hóa.

Trong mã không được tối ưu hóa, bạn chạy hạt nhân với bao nhiêu chủ đề là work là, là blockDim = 2 * unit. Mã trong mỗi khối chỉ sao chép một phần của g_idata vào một mảng trong bộ nhớ dùng chung, có kích thước 2 * unit.

Trong mã được tối ưu hóa blockDim = unit, vì vậy hiện tại có 1/2 chủ đề và mảng trong bộ nhớ dùng chung nhỏ hơn 2x. Trong dòng 3 summon đầu tiên đến từ các đơn vị thậm chí, trong khi thứ hai từ các đơn vị lẻ. Bằng cách này, tất cả các dữ liệu cần thiết để giảm được tính đến.

Ví dụ: Nếu bạn chạy hạt nhân không được tối ưu hóa với blockDim=256=work (khối duy nhất, unit=128), sau đó mã tối ưu có một khối duy nhất của blockDim=128=unit. Vì khối này nhận được blockIdx=0, các *2 không quan trọng; chủ đề đầu tiên là g_idata[0] + g_idata[0 + 128].

Nếu bạn có 512 phần tử và chạy không được tối ưu hóa với 2 khối kích thước 256 (work=256, unit=128), thì mã được tối ưu hóa có 2 khối, nhưng bây giờ là kích thước 128. Chuỗi đầu tiên trong khối thứ hai (blockIdx=1) g_idata[2*128] + g_idata[2*128+128].