取模、乘法和除法运算在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上的实现代码如下：

[cpp] view
plaincopy

// 算法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上的实现代码如下：

[cpp] view
plaincopy

// 算法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，采用浮点运算最快，其次是取模运算，整数除法最慢。

/article/7103821.html