取模、乘法和除法运算在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. // 算法1  
2. for(int ii, i = 0; i < size; i++)  
3. {  
4.     ii = N[i] % IDIM;  
5.     I[i] = ii;  
6.     J[i] = (N[i] - ii) / IDIM;  
7. }  
8.   
9. // 算法2：R1 = 1.0f/IDIM  
10. for(int i=0,j=0;i<size;i++)  
11. {  
12.     j = floor(N[i]*R1);  
13.     I[i] = N[i] - j*IDIM;  
14.     J[i] = j;  
15. }  
16.   
17. // 算法3：R1 = 1.0f / IDIM  
18. for(int i = 0, ii = 0; i < size; i++)  
19. {  
20.     ii = N[i] % IDIM;  
21.     I[i] = ii;  
22.     J[i] = (N[i] - ii) * R1;  
23. }  

GPU上的实现代码如下：

[cpp] view plaincopy

1. // 算法1  
2. __global__ void kernel1(int *N, int *I, int *J, int IDIM, int JDIM)  
3. {  
4.     int tid = blockIdx.x * blockDim.x + threadIdx.x;  
5.     if(tid < IDIM * JDIM)  
6.     {  
7.         int n = N[tid];  
8.         int i = n % IDIM;  
9.         I[tid] = i;  
10.         J[tid] = (n - i) / IDIM;  
11.     }  
12. }  
13.   
14. // 算法2：R1 = 1.0f/IDIM  
15. __global__ void kernel2(int *N, int *I, int *J, int IDIM, int JDIM)  
16. {  
17.     int tid = blockIdx.x * blockDim.x + threadIdx.x;  
18.     int n, j;  
19.     if(tid < IDIM * JDIM)  
20.     {  
21.         n = N[tid];  
22.         j = floor(n*R1);  
23.         I[tid] = n - j * IDIM;  
24.         J[tid] = j;  
25.     }  
26. }  
27.   
28. // 算法3：R1 = 1.0f / IDIM  
29. __global__ void kernel3(int *N, int *I, int *J, int IDIM, int JDIM, float R1)  
30. {  
31.     int tid = blockIdx.x * blockDim.x + threadIdx.x;  
32.     if(tid < IDIM * JDIM)  
33.     {  
34.         int n = N[tid];  
35.         int i = n % IDIM;  
36.         I[tid] = i;  
37.         J[tid] = (n - i) * R1;  
38.     }  
39. }  

计算效率如下：

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