Ergebnis der Reduzierung # 1 ist falsch

Question

Ich habe die Reduktion # 1 aus den bekannten Dias von Mark Harris implementiert, aber ich erhalte 0 als Ergebnis. Ich füllte das Input-Array mit den gleichen Werten wie in den Folien. Ich kompilierte mit cuda 7.0 mit dem Befehl nvcc reduction1.cu -o red1. Wo ist der Fehler? Vielen Dank.

#include <stdio.h> 
#include <cuda_runtime.h> 

#define THREADS_PER_BLOCK 16 

__global__ void reduce1(int *g_idata, int *g_odata) { 
    extern __shared__ int sdata[]; 
    // 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(); 

    // do reduction in shared mem 
    for(unsigned int s=1; s < blockDim.x; s *= 2) 
    { 
     if (tid % (2*s) == 0) sdata[tid] += sdata[tid + s]; 
      __syncthreads(); 
    } 

    // write result for this block to global mem 
    if (tid == 0) g_odata[blockIdx.x] = sdata[0]; 
} 

int main() 
{ 
    int inputLength=16; 
    int hostInput[16]={10,1,8,-1,0,-2,3,5,-2,-3,2,7,0,11,0,2}; 
    int hostOutput=0; 
    int *deviceInput; 
    int *deviceOutput; 

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

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

    reduce1<<<1,THREADS_PER_BLOCK>>>(deviceInput, deviceOutput); 

    cudaDeviceSynchronize(); 

    cudaMemcpy(&hostOutput, deviceOutput,sizeof(int), cudaMemcpyDeviceToHost); 

    printf("%d\n",hostOutput); 

    cudaFree(deviceInput); 
    cudaFree(deviceOutput); 

    return 0; 
} 

Answer 1

Wie talonmies sagte, Sie verwenden dynamische Shared Memory, aber Sie sind jeden Speicherplatz für sie nicht zugeordnet werden. Sie müssen die Größe dieses Speichers als drittes Argument Ihres Kernel Ausführungskonfiguration angeben.

reduce1<<<1, THREADS_PER_BLOCK, 64>>>(deviceInput, deviceOutput); 
           ^^ 

Eine andere Möglichkeit, diesen Code zu beheben, ist statische Shared Memory zu verwenden. Erklären Sie Ihre gemeinsam genutzten Speicher wie folgt aus:

__shared__ int sdata[16]; 

Bitte lesen Sie this vor Fragen für CUDA zu fragen.