Tôi biết rằng "mỗi sợi dọc chứa các chuỗi liên tiếp, tăng ID luồng với sợi dọc đầu tiên chứa chuỗi 0" vì vậy 32 luồng đầu tiên phải nằm trong sợi dọc đầu tiên. Ngoài ra tôi biết rằng tất cả các chủ đề trong một dọc được thực hiện đồng thời trên bất kỳ bộ xử lý đa luồng nào có sẵn.CUDA. Làm thế nào để unroll 32 chủ đề đầu tiên để họ sẽ được thực hiện song song?
Như tôi đã hiểu, vì điều đó không cần đồng bộ chuỗi nếu chỉ có một warp đang được thực thi. Nhưng mã bên dưới tạo ra câu trả lời sai nếu tôi xóa bất kỳ số nào trong số __syncthreads() trong áp chót if khối. Tôi đã cố gắng tìm ra nguyên nhân nhưng cuối cùng lại chẳng có gì. Tôi thực sự hy vọng sự giúp đỡ của bạn, vì vậy bạn có thể cho tôi biết những gì là sai với mã này? Tại sao tôi không thể để lại chỉ __syncthreads() cuối cùng và nhận được câu trả lời đúng?
#define BLOCK_SIZE 128
__global__ void reduce (int * inData, int * outData)
{
__shared__ int data [BLOCK_SIZE];
int tid = threadIdx.x;
int i = blockIdx.x * blockDim.x + threadIdx.x;
data [tid] = inData [i] + inData [i + blockDim.x/2 ];
__syncthreads();
for (int s = blockDim.x/4; s > 32; s >>= 1)
{
if (tid < s)
data [tid] += data [tid + s];
__syncthreads();
}
if (tid < 32)
{
data [tid] += data [tid + 32];
__syncthreads();
data [tid] += data [tid + 16];
__syncthreads();
data [tid] += data [tid + 8];
__syncthreads();
data [tid] += data [tid + 4];
__syncthreads();
data [tid] += data [tid + 2];
__syncthreads();
data [tid] += data [tid + 1];
__syncthreads();
}
if (tid == 0)
outData [blockIdx.x] = data [0];
}
void main()
{
...
reduce<<<dim3(128), dim3(128)>>>(dev_data, dev_res);
...
}
P.S. Tôi đang sử dụng GT560Ti