扫描阵列CUDA

Question

我试图扫描一个简单的数组使用CUDA，但它似乎有错误的代码在下面..我试图找到我做错了什么，但我can't.Can任何人都可以请帮我？

#include <stdio.h> 
#include <stdlib.h> 


__global__ void prescan(int *g_odata, int *g_idata, int n){ 

    extern __shared__ int temp[];// allocated on invocation 
    int thid = threadIdx.x; 
    int offset = 1; 
    temp[2*thid] = g_idata[2*thid]; // load input into shared memory 
    temp[2*thid+1] = g_idata[2*thid+1]; 


    for (int d = n>>1; d > 0; d >>= 1){ // build sum in place up the tree 
     __syncthreads(); 
     if (thid < d){ 

      int ai = offset*(2*thid+1)-1; 
      int bi = offset*(2*thid+2)-1; 
      temp[bi] += temp[ai]; 
     } 
     offset *= 2; 
    } 

    if (thid == 0) { temp[n - 1] = 0; } // clear the last element 

    for (int d = 1; d < n; d *= 2){ // traverse down tree & build scan 
     offset >>= 1; 
     __syncthreads(); 
     if (thid < d){ 

      int ai = offset*(2*thid+1)-1; 
      int bi = offset*(2*thid+2)-1; 
      int t = temp[ai]; 
      temp[ai] = temp[bi]; 
      temp[bi] += t; 
     } 
    } 

    __syncthreads(); 

    g_odata[2*thid] = temp[2*thid]; // write results to device memory 
    g_odata[2*thid+1] = temp[2*thid+1]; 
} 

int main(int argc, char *argv[]){ 

    int i; 
    int *input = 0; 
    int *output = 0; 
    int *g_idata = 0; 
    int *g_odata = 0; 

    int numblocks = 1; 
    int radix = 16; 

    input = (int*)malloc(numblocks*radix*sizeof(int)); 
    output = (int*)malloc(numblocks*radix*sizeof(int)); 

    cudaMalloc((void**)&g_idata, numblocks*radix*sizeof(int)); 
    cudaMalloc((void**)&g_odata, numblocks*radix*sizeof(int)); 

    for(i=0; i<numblocks*radix; i++){ 
      input[i] = 1 + 2*i; 
    } 

    for(i=0; i<numblocks*radix; i++){ 
     printf("%d ", input[i]); 
    } 

    cudaMemcpy(g_idata, input, numblocks*radix*sizeof(int), cudaMemcpyHostToDevice); 

    prescan<<<1,8>>>(g_odata, g_idata, numblocks*radix); 

    cudaThreadSynchronize(); 

    cudaMemcpy(output, g_odata, numblocks*radix*sizeof(int), cudaMemcpyDeviceToHost); 

    for(i=0; i<numblocks*radix; i++){ 
     printf("%d ", output[i]); 
    } 

    free(input); 
    free(output); 
    cudaFree(g_idata); 
    cudaFree(g_odata); 

    return 0; 
} 

输出是这样的：1 3 5 7 9 11 13 15 17 19 21 23 25 27 29 31 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0。I希望有这输出：1 3 5 7 9 11 13 15 17 19 21 23 25 27 29 31 0 1 4 9 16 25 36 49 64 81 100 121 144 169 196 225

Answer 1

只需通过此代码即可在并行环境中执行扫描。 我在这里实现的算法是Hillis Steele独家扫描。我通过共享内存实现算法，它肯定会提高大数据集的执行时间。

#include<stdio.h> 
#include<math.h> 
__global__ void scan(int *d_in,int *d_out,int n) 
{ 
    extern __shared__ int sdata[]; 
    int i; 
    int tid = threadIdx.x; 
    sdata[tid] = d_in[tid]; 
    for (i = 1; i <n; i <<= 1) 
    { 

     if (tid>=i) 
     { 
      sdata[tid] +=sdata[tid-i]; 
     } 
     __syncthreads(); 
    } 
    d_out[tid] = sdata[tid]; 
    __syncthreads(); 
} 

int main() 
{ 

    int h_in[16],h_out[16]; 
    int i,j; 
    for (i = 0; i < 16; i++) 
     h_in[i] = 2*i+1; 
    for (i = 0; i < 16; i++) 
     printf("%d ", h_in[i]); 
    int *d_in; 
    int *d_out; 
    cudaMalloc((void**)&d_in, sizeof(int)* 16); 
    cudaMalloc((void**)&d_out, sizeof(int)* 16); 
    cudaMemcpy(d_in, h_in, sizeof(int) * 16, cudaMemcpyHostToDevice); 

    scan <<<1, 16, sizeof(int)*16 >>>(d_in,d_out, 16); 

    cudaMemcpy(h_out, d_out, sizeof(int) * 16, cudaMemcpyDeviceToHost); 
    for (i = 0; i < 16; i++) 
     printf("%d ", h_out[i]); 
    return 0; 
}