CUDA - Parallel Reduction Summe

Question

Ich versuche eine parallele Reduktion Summe in CUDA 7.5 zu implementieren. Ich habe versucht, die NVIDIA PDF zu folgen, die Sie durch den anfänglichen Algorithmus und dann ständig mehr optimierte Versionen führt. Ich mache derzeit ein Array, das als Wert in jeder Array-Position mit 1 gefüllt ist, so dass ich die Ausgabe überprüfen kann, aber ich bekomme einen Wert von -842159451 für ein Array der Größe 64. Ich erwarte, dass der Kernel-Code richtig ist, wie ich den genauen Code von NVIDIA für sie verfolgt habe, aber hier ist mein kernel:

__global__ void reduce0(int *input, int *output) { 
    extern __shared__ int sdata[]; 

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

    sdata[tid] = input[i]; 

    __syncthreads(); 

    for (unsigned int s = 1; s < blockDim.x; s *= 2) { 
     if (tid % (2 * s) == 0) { 
      sdata[tid] += sdata[tid + s]; 
     } 

     __syncthreads(); 
    } 

    if (tid == 0) output[blockIdx.x] = sdata[0]; 
} 

hier mein Code der Kernel aufrufe, das ist, wo ich mein Problem zu erwarten:

int main() 
{ 
    int numThreadsPerBlock = 1024; 

    int *hostInput; 
    int *hostOutput; 
    int *deviceInput; 
    int *deviceOutput; 

    int numInputElements = 64; 
    int numOutputElements; // number of elements in the output list, initialised below 

    numOutputElements = numInputElements/(numThreadsPerBlock/2); 
    if (numInputElements % (numThreadsPerBlock/2)) { 
     numOutputElements++; 
    } 

    hostInput = (int *)malloc(numInputElements * sizeof(int)); 
    hostOutput = (int *)malloc(numOutputElements * sizeof(int)); 


    for (int i = 0; i < numInputElements; ++i) { 
     hostInput[i] = 1; 
    } 

    const dim3 blockSize(numThreadsPerBlock, 1, 1); 
    const dim3 gridSize(numOutputElements, 1, 1); 

    cudaMalloc((void **)&deviceInput, numInputElements * sizeof(int)); 
    cudaMalloc((void **)&deviceOutput, numOutputElements * sizeof(int)); 

    cudaMemcpy(deviceInput, hostInput, numInputElements * sizeof(int), cudaMemcpyHostToDevice); 

    reduce0 << <gridSize, blockSize >> >(deviceInput, deviceOutput); 

    cudaMemcpy(hostOutput, deviceOutput, numOutputElements * sizeof(int), cudaMemcpyDeviceToHost); 

    for (int ii = 1; ii < numOutputElements; ii++) { 
     hostOutput[0] += hostOutput[ii]; //accumulates the sum in the first element 
    } 

    int sumGPU = hostOutput[0]; 

    printf("GPU Result: %d\n", sumGPU); 

    std::string wait; 
    std::cin >> wait; 

    return 0; 
} 

Ich habe auch versucht, größere und kleinere Array-Größen für die Eingabe und ich bekomme das gleiche Ergebnis eines sehr großen negativen Wertes, unabhängig von der Größe des Arrays.

Answer 1

scheint, dass Sie einen dynamisch zugewiesenen Shared-Array verwenden:

reduce0 <<<gridSize, blockSize >>>(deviceInput, deviceOutput); 

Sie haben zwei Möglichkeiten:

Option

extern __shared__ int sdata[]; 

aber Sie es nicht in den Kernel-Aufruf werden die Zuteilung 1

Ordnen Sie den gemeinsamen Speicher statisch im Kernel zu, z.

constexpr int threadsPerBlock = 1024; 
__shared__ int sdata[threadsPerBlock]; 

Mehr als oft nicht finde ich das die sauberste Ansatz, wie es ohne ein Problem funktioniert, wenn Sie mehrere Arrays im gemeinsam genutzten Speicher haben. Der Nachteil ist, dass, während die Größe normalerweise von der Anzahl der Threads im Block abhängt, die Größe zur Kompilierungszeit bekannt sein muss.

Option 2

Geben Sie die Menge des dynamisch in den Kernel-Aufruf gemeinsam genutzten Speicher zugeordnet.

reduce0 <<<gridSize, blockSize, numThreadsPerBlock*sizeof(int) >>>(deviceInput, deviceOutput); 

Dies wird für jeden Wert von numThreadsPerBlock arbeiten (sofern es innerhalb des zulässigen Bereichs natürlich ist). Der Nachteil ist, dass wenn Sie mehrere extern geteilte Arrays haben, Sie herausfinden müssen, wie Sie sie dann selbst in den Speicher legen, damit der andere nicht überschrieben wird.

---

Beachten Sie, dass in Ihrem Code möglicherweise andere Probleme auftreten. Ich habe es nicht getestet. Das ist etwas, was ich sofort nach einem Blick auf Ihren Code gesehen habe.