取模、乘法和除法运算在CPU和GPU上的效率

问题:

    将整数n分解为i和j，满足下面关系：

n = j*idim + i

    其中idim为常量。

    以下为三种算法实现：
    1） i = n%idim，j = (n-i)/idim

    2） j = n*ridim，i ＝ n - j*idim，其中ridim = 1.0f/idim，为浮点数。

    3） i = n%idim，j = (n-i)*ridim，其中ridim = 1.0f/idim，为浮点数。

 

    CPU上的实现代码如下：

// 算法1
for(int i=0,ii=0;i<size;i++)
{
    ii=N[i]%IDIM;
    I[i] = ii;
    J[i] = (N[i]-ii)/IDIM;
}

// 算法2：R1 = 1.0f/IDIM
for(int i=0,j=0;i<size;i++)
{
    j = floor(N[i]*R1);
    I[i] = N[i] - j*IDIM;
    J[i] = j;
}

// 算法3：R1 = 1.0f/IDIM
for(int i=0,ii=0;i<size;i++)
{
    ii=N[i]%IDIM;
    I[i] = ii;
    J[i] = (N[i]-ii)*R1;
}

    GPU上的实现代码如下：

// 算法1
__global__ void kernel1(int *N,int *I,int *J,int IDIM,int JDIM)
{
    int tid = blockIdx.x*blockDim.x+threadIdx.x;
    if(tid<IDIM*JDIM)
    {
        int n = N[tid];
        int i = n%IDIM;
        I[tid] = i;
        J[tid] = (n-i)/IDIM;
    }
}

// 算法2：R1 = 1.0f/IDIM
__global__ void kernel2(int *N,int *I,int *J,int IDIM,int JDIM)
{
    int tid = blockIdx.x*blockDim.x+threadIdx.x;
    int n,j;
    if(tid<IDIM*JDIM)
    {
        n = N[tid];
        j = floor(n*R1);
        I[tid] = n-j*IDIM;
        J[tid] = j;
    }
}

// 算法3：R1 = 1.0f/IDIM
__global__ void kernel3(int *N,int *I,int *J,int IDIM,int JDIM,float R1)
{
    int tid = blockIdx.x*blockDim.x+threadIdx.x;
    if(tid<IDIM*JDIM)
    {
        int n = N[tid];
        int i = n%IDIM;
        I[tid] = i;
        J[tid] = (n-i)*R1;
    }
}

计算效率如下：

N = 1000000, IDIM = 1000, JDIM = 1000

Core2 Q6600:

    算法1：  17 ms

    算法2：  34 ms

    算法3：  16 ms

GTX280：

    算法1：   0.36 ms

    算法2：   0.14 ms

    算法3：   0.23 ms

CUDA Visual Profiler的检测结果显示： 算法1的指令数高达98xxx，而算法2指令数仅为29xxx，算法3的指令数为65xxx。整数除法再一次应验了手册上的那句话：

Integer division and modulo operation are particularly costly and should be avoided...

但是好像取模运算并没有想象中的那么慢。

 

结论：

对于CPU，最好采用取模运算，整数除法和单精度乘法的效率差不多。

对于GPU，采用浮点运算最快，其次是取模运算，整数除法最慢。