Cuda输出随计算能力而变化

我试图运行类似于稀疏矩阵向量乘法（SpMV）内核的以下内核。

__global__ void cost_spmv(const int *population,const int *row,const int *col,int *out){
    /*Calculate the cost vector for multiplication of the matrices*/
    //int tid=threadIdx.x+blockDim.x*blockIdx.x;
    int lane=threadIdx.x;
    extern __shared__ int val[];
    int r=blockIdx.x;
    int rowStart=row[r];
    int rowEnd=row[r+1];
    val[threadIdx.x]=0;
    for(int i=rowStart+lane;i<rowEnd;i+=32)
        val[threadIdx.x]+=  population[col[i]];
    __syncthreads();
    if(lane<16)
        val[threadIdx.x]+=val[threadIdx.x+16];
    if(lane<8)
        val[threadIdx.x]+=val[threadIdx.x+8];
    if(lane<4)
        val[threadIdx.x]+=val[threadIdx.x+4];
    if(lane<2)
        val[threadIdx.x]+=val[threadIdx.x+2];
    if(lane<1)
        val[threadIdx.x]+=val[threadIdx.x+1];
    if(lane==0)
        out[r]=val[threadIdx.x];
}

内核使用调用

cost_spmv<<<numRows,32,32*sizeof(int)>>>(population,rowArray,colArray, out)

其中numRows是数组总数的大小，out和rowArray(numRows+1 actually) 。 rowArray[i]包含属于第i行的元素的起始索引.colArray的大小为rowArray[numRows] 。 colArray[i]包含使用rowArray描述的行的非零值的列号。

然而， 在Tesla P4上编译计算能力3.5时，与计算能力6.1相比，我得到了不同的答案。 我在Tesla P4上使用计算能力6.1的答案与我在920m上使用计算能力3.5时得到的答案相同。 这可能是什么原因呢？

---

请记住，CUDA编译器有一个单线程的世界视图。 它不知道用于执行代码的运行时配置，这在编译时不可用。

在val[]和之前写入val[]的加载之间的代码中没有表示相关性。 因此，编译器可以自由移动它认为合适的负载。 在某些情况下，它可能会选择尽早发出一些或全部负载来增加代码的负载延迟容限，例如通过如下转换代码：

int __temp0 = val[threadIdx.x+16];
int __temp1 = val[threadIdx.x+ 8];
int __temp2 = val[threadIdx.x+ 4];
int __temp3 = val[threadIdx.x+ 2];
int __temp4 = val[threadIdx.x+ 1];

if(lane<16)
    val[threadIdx.x]+=__temp0;
if(lane<8)
    val[threadIdx.x]+=__temp1;
if(lane<4)
    val[threadIdx.x]+=__temp2;
if(lane<2)
    val[threadIdx.x]+=__temp3;
if(lane<1)
    val[threadIdx.x]+=__temp4;

根据编译器选择放置负载的位置，缩减顺序的结果将有所不同。 CUDA编译器中的代码生成和指令调度特别因GPU架构而异，因此在针对不同GPU架构进行编译时可能会观察到不同的结果。

为了强制加载和存储之间所需的依赖关系，CUDA编程模型认可的方法是在每个缩减步骤后使用__syncthreads()来创建屏障。 通过使用volatile修饰符，代码范围之外的代理可以声明val更改，但实现期望结果的方法可能更快但却很冒险。 这可以防止编译器从val[]移动负载。

链接地址: http://www.djcxy.com/p/80113.html

上一篇: Cuda output varying with compute capability

下一篇: Maximum number of blocks and threads working in parallel for a shared variable