如果数据数量小于线程数量会怎样？

[hesicong]

我才刚CUDA入门，正在试验中。下面有一段CUDA执行的让当前位置赋值为当前位置坐标的一段程序。如果我的数据只有3*3=9个，而只开了一个block，一个block开了4*4=16个线程。那么在GPU执行中，是执行到9个线程呢，还是16个线程都执行呢？
如果16个线程都执行，那么最后7个线程不是会产生越界错误？（我在CPU上执行是这样的，GPU上会得到准确答案）
如果这是一个异常的情况，数据不一定能够刚好等于线程数量，如何避免这种情况呢？谢谢！

__global__ void assignNumber(float* d_out, float dim)
{
    int x=blockIdx.x*blockDim.x+threadIdx.x;
    int y=blockIdx.y*blockDim.y+threadIdx.y;
    int pos=y*dim+x;
    //printf("%d\t%d\t%d\t%d\t%d\t%d\t%d\n",gridDim.x, gridDim.y, blockIdx.x, blockIdx.y, threadIdx.x, threadIdx.y, pos);
    d_out[pos]=pos;
}

void main()
{
    int dim=3;
    float* d_out;
        cudaMalloc((void**)&d_out, dim*dim*sizeof(float)));

    dim3 thread(4,4);
    dim3 grid( (dim-1)/thread.x+1, (dim-1)/thread.y+1);

    assignNumber<<<grid, thread>>>(d_out,dim);

    float* h_out=new float[dim*dim];
    cudaMemcpy(h_out, d_out, dim*dim*sizeof(float), cudaMemcpyDeviceToHost);

//Check if pos equals the value
    bool rst=true;
    for(int i=0;i<dim;i++)
    {
        for(int j=0;j<dim;j++)
        {
            int pos=i*dim+j;
            if(h_out[pos]!=pos)
            {
                rst=false;
                break;
            }
        }
        if(!rst) break;
    }
    if(rst)
    {
        cout<<"Success"<<endl;
    }
    else
    {
        cout<<"Failed"<<endl;
    }        
}

[无心人_过过小日子]

1)16个同时运行.
2)会越界错误,但不一定是致命错误(gpu上没有内存保护机制)
3)加if语句if(threadIdx.x<9){...}

[hesicong]

[Quote=引用 1 楼 l7331014 的回复:]
1)16个同时运行.
2)会越界错误,但不一定是致命错误(gpu上没有内存保护机制)
3)加if语句if(threadIdx.x <9){...}
[/Quote]

谢谢。
我想if执行的话会影响效率，能不能将数据装满呢？做padding是不是一个很好的方法呢？

[无心人_过过小日子]

[Quote=引用 2 楼 hesicong 的回复:]
我想if执行的话会影响效率，能不能将数据装满呢？做padding是不是一个很好的方法呢？
[/Quote]

是的,如果"多余"数据不多的话是个好方法,一般到32(warp大小)的倍数就可以了.
超过的32的整数倍的边界完全可以用减少block尺寸的办法来避免多余的计算.

[hesicong]

另外，当dim=4096的时候，程序依然可以正确运行，但是当dim小于或等于4096的时候，cudaMalloc执行正常，kernel也执行正常。GPU执行结果不正确。
在CPU执行检查的时候i=4095,j=2,pos=16777217的时候，值会是0。请问这是怎么回事呢？为什么会是4096以上就不行了呢？莫非有什么限制？

[无心人_过过小日子]

4096x4096x4=64MB.LZ显卡的显存有多大?

[hesicong]

[Quote=引用 5 楼 l7331014 的回复:]
4096x4096x4=64MB.LZ显卡的显存有多大?
[/Quote]
我也计算的64M，我用的是G80，8800GTS, 320M，应该够用了。而且cudaMalloc返回success

[无心人_过过小日子]

__global__ void assignNumber(float* d_out, float dim)
改为
__global__ void assignNumber(float* d_out, int dim)
试试?

[hesicong]

[Quote=引用 7 楼 l7331014 的回复:]
__global__ void assignNumber(float* d_out, float dim)
改为
__global__ void assignNumber(float* d_out, int dim)
试试?

[/Quote]
一样的结果~很郁闷。

[hesicong]

16777217这个数很奇怪，16777216=256*256*256~莫非有啥子关系？和静态寻址有关？我晓得之前OpenGL里面贴图只能到4096*4096，相乘也得这个数。

[hesicong]

大哥，谢谢了。犯了一个错误，CPU验证的时候用的是int和float比较。当数据太大以后float和int比较就会出问题。将gpu代码换成int就对了。