cuda 内核执行失败 为什么呢？

[lf864]

联系归约算法时，出现错误，很是不解，下面是程序代码（部分）；望高手指点一二#include<stdio.h>
#include<math.h>
#include<direct.h>
#include <io.h>
#include <stdlib.h>
#include <windows.h>
#include <cutil_inline.h>
__global__ void kernel1(float *idata,float *odata)
{
int id=blockIdx.x*blockDim.x+threadIdx.x;
int index=threadIdx.x;
extern __shared__ float sp[];
sp[index]=idata[id];
__syncthreads();
for(int i=1;i<blockDim.x;i*=2)
{
if(index%(2*i)==0)
{
sp[index]+=sp[index+i];
}
__syncthreads();
}
if(threadIdx.x==0)
odata[blockIdx.x]=sp[0];

}
main()
{
float m[10000]={1.0};
float s=0.0;
float *d_p1;
int size=sizeof(m);
cudaMalloc((void **)&d_p1,size);
cudaMemcpy(d_p1,m,size,cudaMemcpyHostToDevice);
float *d_p2;
cudaMalloc((void **)&d_p2,size);
dim3 dimGrid(32,1);
dim3 dimBlock(512,1,1);
kernel1<<<dimGrid,dimBlock>>>(d_p1,d_p2);
cutilCheckMsg("Kernel execution failed\n");
float sum[32];
cudaMemcpy(sum,d_p2,32*sizeof(float),cudaMemcpyDeviceToHost);
for(int i=0;i<32;i++)
{
s+=sum[i];
}
printf("%f",sum);
}
执行结果如下:
cutilCheckMsg cudaThreadSynchronize error:Kernel execution failed : unspecified launch failure.

[xiaolanyu]

你那个dim3 dimGrid(32,1)，你尝试改下dimGrid(32,1,1).看报错时提示，好像是启动失败，一般是都是配置参数出问题了。

[violetrees]

应该是你for循环中对thread的操作导致分支过乱而引起的问题。你的目的应该是实现shared memory内数据的相加。因为一个block不会太大，你可以分开写
if（threadIdx.x<(blockDim.x<<1)) sp[threadIdx.x]+=sp[threadIdx.x+(blockDim.x<<1)];
__syncthreads();
if(（threadIdx.x<(blockDim.x<<2)) sp[threadIdx.x]+=sp[threadIdx.x+(blockDim.x<<2)];
。。。。。。
如果是在一个warp里，可以省略同步函数。另外，在gpu实现取余耗时较多，建议用位操作

[zas123]

extern __shared__ float sp[];
sp[index]=idata[id];
这两句怎么样

[alpha.5]

extern __shared__ float sp[];没有外部定义；
去掉 extern.

[agathah]

老大……你忘记配置shared memory的大小了……
extern的shared memory需要在launch的时候指定大小

kernel1<<<dimGrid,dimBlock,shared_size>>>(d_p1,d_p2);
尖括号里的第三个参数，看kernel函数应该是sizeof(float)*dimBlock.x