未指定的发射失败-数据自动化系统中的并行扫描

Question

我使用GeForce GT 520 (计算能力v2.1)来运行一个程序，该程序对一个int元素数组执行扫描操作。下面是代码：

/*
This is an implementation of the parallel scan algorithm.
Only a single block of threads is used. Maximum array size = 2048
*/

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

#define errorCheck(ans) { gpuAssert((ans), __FILE__, __LINE__); }

inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
if (code != cudaSuccess) 
{
    fprintf(stderr,"GPUassert: %s, file: %s line: %d\n", cudaGetErrorString(code), file, line);
    if (abort) exit(code);
}
}

__global__ void blelloch_scan(int* d_in, int* d_out, int n) 
{
    extern __shared__ int temp[];// allocated on invocation

    int thid = threadIdx.x;
    int offset = 1;

    temp[2*thid] = d_in[2*thid]; // load input into shared memory
    temp[2*thid+1] = d_in[2*thid+1];

    // build sum in place up the tree
    for (int d = n>>1; d > 0; d >>= 1)
    {
        __syncthreads();
        if (thid < d)
        {
            int ai = offset*(2*thid+1)-1;
            int bi = offset*(2*thid+2)-1;
            temp[bi] += temp[ai];
        }
        offset *= 2;
    }

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

    // traverse down tree & build scan
    for (int d = 1; d < n; d *= 2)
    {
        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();  

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

int main(int argc, char **argv)
{
int ARRAY_SIZE;
if(argc != 2)
{
    printf("Input Syntax: ./a.out <number-of-elements>\nProgram terminated.\n");
    exit (1);        
}
else
ARRAY_SIZE = (int) atoi(*(argv+1));

int *h_in, *h_out, *d_in, *d_out, i;
h_in = (int *) malloc(sizeof(int) * ARRAY_SIZE);
h_out = (int *) malloc(sizeof(int) * ARRAY_SIZE);

cudaSetDevice(0);
cudaDeviceProp devProps;
if (cudaGetDeviceProperties(&devProps, 0) == 0)
{
    printf("Using device %d:\n", 0);
    printf("%s; global mem: %dB; compute v%d.%d; clock: %d kHz\n",
           devProps.name, (int)devProps.totalGlobalMem, 
           (int)devProps.major, (int)devProps.minor, 
           (int)devProps.clockRate);
}

for(i = 0; i < ARRAY_SIZE; i++)
{
    h_in[i] = i;    
}

errorCheck(cudaMalloc((void **) &d_in, sizeof(int) * ARRAY_SIZE));
errorCheck(cudaMalloc((void **) &d_out, sizeof(int) * ARRAY_SIZE));    
errorCheck(cudaMemcpy(d_in, h_in, ARRAY_SIZE * sizeof(int), cudaMemcpyHostToDevice));

blelloch_scan <<<1, ARRAY_SIZE / 2, sizeof(int) * ARRAY_SIZE>>> (d_in, d_out, ARRAY_SIZE);
cudaDeviceSynchronize();
errorCheck(cudaGetLastError());

errorCheck(cudaMemcpy(h_out, d_out, ARRAY_SIZE * sizeof(int), cudaMemcpyDeviceToHost));

printf("Results:\n");    
for(i = 0; i < ARRAY_SIZE; i++)
{
    printf("h_in[%d] = %d, h_out[%d] = %d\n", i, h_in[i], i, h_out[i]);    
}
return 0;
}

在使用nvcc -arch=sm_21 parallel-scan.cu -o parallel-scan编译时，我得到一个错误：GPUassert: unspecified launch failure, file: parallel-scan-single-block.cu line: 106

当我们使用errorCheck检查错误时，第106行是内核启动后的行。

这是我计划实施的：

• 从内核中可以看出，如果一个块有1000个线程，那么它可以对2000元素进行操作。因此，blockSize = ARRAY_SIZE / 2。
• 和，共享内存= sizeof(int) * ARRAY_SIZE
• 所有的东西都被装进了共享的我。然后，完成向上扫描，将最后一个元素设置为0。最后，进行向下扫描，对元素进行独占扫描。

我使用这个文件作为参考来编写这段代码。我不明白我的代码中有什么错误。任何帮助都将不胜感激。

Answer 1

您正在像这样启动内核

blelloch_scan <<<1, ARRAY_SIZE / 2, sizeof(int) * ARRAY_SIZE>>>

这意味着witihin然后是内核0 < thid < int(ARRAY_SIZE/2)。

但是，内核需要最少的可用共享内存的(2 * int(ARRAY_SIZE/2)) + 1字才能正常工作，否则：

temp[2*thid+1] = d_in[2*thid+1];

将产生一个不受限制的共享内存访问。

如果我的整数数学技巧不是太生疏，这应该意味着如果ARRAY_SIZE是奇数，代码将是安全的，因为对于任何奇数整数，ARRAY_SIZE == (2 * int(ARRAY_SIZE/2)) + 1都是安全的。但是，如果ARRAY_SIZE是偶数，那么ARRAY_SIZE < (2 * int(ARRAY_SIZE/2)) + 1和您就有问题了。

可能是共享内存页大小粒度节省了一些理论上应该失败的ARRAY_SIZE的偶数值，因为硬件总是将动态共享内存分配到大于请求大小的下一个页面大小。但是，应该有许多偶数的ARRAY_SIZE值，但这是失败的。

我不能评论内核的其余部分是否正确，但是使用共享内存大小的sizeof(int) * size_t(1 + ARRAY_SIZE)应该可以解决这个特定的问题。