取模、乘法和除法运算在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 ii, i = 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，采用浮点运算最快，其次是取模运算，整数除法最慢。

http://www.cnblogs.com/codezhang/archive/2009/06/19/1506532.html