kernel函数占用寄存器数量的疑问

system · 2013 年6 月 4 日 09:40

写了一个小的cu测试程序，计算16x16像素块对应的sad值（对应像素相减的绝对值和），用nsight运行发现核函数占用了不少的寄存器，代码如下

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <stdlib.h>

#define SIZE 16

typedef unsigned char pixel;

__device__ int x264_pixel_sad_16x16(pixel *pix1, pixel *pix2, int i_stride_pix ) 
{                                                   
   int i_sum = 0;                                  
	int y;

   int r1 = ((int)pix1) & 7;
   int r2 = ((int)pix2) & 7;
   for (y = 0; y < 16; y++)
   {
   __int64 p1 = *((__int64*)(pix1));
   __int64 p2 = *((__int64*)(pix1+8));
   __int64 p3 = 0;
   __int64 q1 = *((__int64*)(pix2));
   __int64 q2 = *((__int64*)(pix2+8));
   __int64 q3 = 0;
   if (r1 != 0) p3 = *((__int64*)(pix1-r1+16));
   if (r2 != 0) q3 = *((__int64*)(pix2-r2+16));

   p1 = (p1>>r1*8 & ((__int64)1<<(64-r1*8))-1) | (p2&(((__int64)1<<r1*8)-1))<<(64-r1*8);
   p2 = (p2>>r1*8 & ((__int64)1<<(64-r1*8))-1) | (p3&(((__int64)1<<r1*8)-1))<<(64-r1*8);
   q1 = (q1>>r2*8 & ((__int64)1<<(64-r2*8))-1) | (q2&(((__int64)1<<r2*8)-1))<<(64-r2*8);
   q2 = (q2>>r2*8 & ((__int64)1<<(64-r2*8))-1) | (q3&(((__int64)1<<r2*8)-1))<<(64-r2*8);
   i_sum += abs((int)(p1    &0xff) - (int)(q1    &0xff));
   i_sum += abs((int)(p1>>8 &0xff) - (int)(q1>>8 &0xff));
   i_sum += abs((int)(p1>>16&0xff) - (int)(q1>>16&0xff));
   i_sum += abs((int)(p1>>24&0xff) - (int)(q1>>24&0xff));
   i_sum += abs((int)(p1>>32&0xff) - (int)(q1>>32&0xff));
   i_sum += abs((int)(p1>>40&0xff) - (int)(q1>>40&0xff));
   i_sum += abs((int)(p1>>48&0xff) - (int)(q1>>48&0xff));
   i_sum += abs((int)(p1>>56&0xff) - (int)(q1>>56&0xff));
   i_sum += abs((int)(p2    &0xff) - (int)(q2    &0xff));
   i_sum += abs((int)(p2>>8 &0xff) - (int)(q2>>8 &0xff));
   i_sum += abs((int)(p2>>16&0xff) - (int)(q2>>16&0xff));
   i_sum += abs((int)(p2>>24&0xff) - (int)(q2>>24&0xff));
   i_sum += abs((int)(p2>>32&0xff) - (int)(q2>>32&0xff));
   i_sum += abs((int)(p2>>40&0xff) - (int)(q2>>40&0xff));
   i_sum += abs((int)(p2>>48&0xff) - (int)(q2>>48&0xff));
   i_sum += abs((int)(p2>>56&0xff) - (int)(q2>>56&0xff)); 
   
   pix1 += i_stride_pix;                      
   pix2 += i_stride_pix; 
   }  
   return i_sum;                                   
}

__device__ int x264_pixel_sad_16x16_L( pixel *pix1, pixel *pix2, int i_stride ) 
{                                                   
   int i_sum = 0;                                  
	int y;

   int r1 = ((int)pix1) & 3;
   int r2 = ((int)pix2) & 3;
   for (y = 0; y < 16; y++)
   {
   int p1 = *((int*)(pix1-r1));
   int p2 = *((int*)(pix1-r1+4));
   int p3 = *((int*)(pix1-r1+8));
   int p4 = *((int*)(pix1-r1+12));
   int p5 = *((int*)(pix1-r1+16));

   int q1 = *((int*)(pix2-r2));
   int q2 = *((int*)(pix2-r2+4));
   int q3 = *((int*)(pix2-r2+8));
   int q4 = *((int*)(pix2-r2+12));
   int q5 = *((int*)(pix2-r2+16));


   p1 = (p1>>r1*8 & (1<<(32-r1*8))-1) | (p2&((1<<r1*8)-1))<<(32-r1*8);
   p2 = (p2>>r1*8 & (1<<(32-r1*8))-1) | (p3&((1<<r1*8)-1))<<(32-r1*8);
   p3 = (p3>>r1*8 & (1<<(32-r1*8))-1) | (p4&((1<<r1*8)-1))<<(32-r1*8);
   p4 = (p4>>r1*8 & (1<<(32-r1*8))-1) | (p5&((1<<r1*8)-1))<<(32-r1*8);

   q1 = (q1>>r2*8 & (1<<(32-r2*8))-1) | (q2&((1<<r2*8)-1))<<(32-r2*8);
   q2 = (q2>>r2*8 & (1<<(32-r2*8))-1) | (q3&((1<<r2*8)-1))<<(32-r2*8);
   q3 = (q3>>r2*8 & (1<<(32-r2*8))-1) | (q4&((1<<r2*8)-1))<<(32-r2*8);
   q4 = (q4>>r2*8 & (1<<(32-r2*8))-1) | (q5&((1<<r2*8)-1))<<(32-r2*8);


   i_sum += abs((p1    &0xff) - (q1    &0xff));
   i_sum += abs((p1>>8 &0xff) - (q1>>8 &0xff));
   i_sum += abs((p1>>16&0xff) - (q1>>16&0xff));
   i_sum += abs((p1>>24&0xff) - (q1>>24&0xff));
   i_sum += abs((p2    &0xff) - (q2    &0xff));
   i_sum += abs((p2>>8 &0xff) - (q2>>8 &0xff));
   i_sum += abs((p2>>16&0xff) - (q2>>16&0xff));
   i_sum += abs((p2>>24&0xff) - (q2>>24&0xff));
   i_sum += abs((p3    &0xff) - (q3    &0xff));
   i_sum += abs((p3>>8 &0xff) - (q3>>8 &0xff));
   i_sum += abs((p3>>16&0xff) - (q3>>16&0xff));
   i_sum += abs((p3>>24&0xff) - (q3>>24&0xff));
   i_sum += abs((p4    &0xff) - (q4    &0xff));
   i_sum += abs((p4>>8 &0xff) - (q4>>8 &0xff));
   i_sum += abs((p4>>16&0xff) - (q4>>16&0xff));
   i_sum += abs((p4>>24&0xff) - (q4>>24&0xff)); 
   
   pix1 += i_stride;                      
   pix2 += i_stride; 
   }                                                    

   return i_sum;                                   
}


__device__ int x264_pixel_sad_16x16_LL( pixel *pix1, pixel *pix2, int i_stride_pix ) 
{                                                   
   int i_sum = 0;                                  
	int x, y;                                      
   for( y = 0; y < 16; y++ )                      
   {                                              
   for( x = 0; x < 16; x++ )                   
   {                                          
   i_sum += abs( pix1[x] - pix2[x] );      
   }                                         
   pix1 += i_stride_pix;                     
   pix2 += i_stride_pix;                      
   }                                              
   return i_sum;                                  
}


__global__ void cal_sad(int *sad, pixel * fenc, pixel * fref, int w, int h)
{
   int id = blockIdx.x*w+threadIdx.x;
   int offset = blockIdx.x*w*16*16 + threadIdx.x*16;
   sad[id] = x264_pixel_sad_16x16(fenc+offset, fref+offset, w*16);    

   
}

__global__ void cal_sad_L(int *sad, pixel * fenc, pixel * fref, int w, int h)
{
   int id = blockIdx.x*w+threadIdx.x;
   int offset = blockIdx.x*w*16*16 + threadIdx.x*16;
   sad[id] = x264_pixel_sad_16x16_L(fenc+offset, fref+offset, w*16);    
}
__global__ void cal_sad_LL(int *sad, pixel * fenc, pixel * fref, int w, int h)
{
   int id = blockIdx.x*w+threadIdx.x;
   int offset = blockIdx.x*w*16*16 + threadIdx.x*16;
   sad[id] = x264_pixel_sad_16x16_LL(fenc+offset, fref+offset, w*16);    
}

int main()
{    
   cudaError_t err;
   int w = 256, h = 64;
   pixel* fenc_c, *fref_c;
   pixel* fenc_g, *fref_g;
   int *sad_c, *sad_g, *sad_s;
   int size = w*h*SIZE*SIZE;

   cudaHostAlloc(&fenc_c, size, cudaHostAllocDefault);
   cudaHostAlloc(&fref_c, size, cudaHostAllocDefault);
   cudaHostAlloc(&sad_c,  w*h*sizeof(int), cudaHostAllocDefault);
   cudaHostAlloc(&sad_s,  w*h*sizeof(int), cudaHostAllocDefault);

   cudaMalloc(&fenc_g, size);
   cudaMalloc(&fref_g, size);
   cudaMalloc(&sad_g,  w*h*sizeof(int));

   int i, j, k;
   for (i = 0; i < size; i++)
   {
   fenc_c[i] = rand();
   fref_c[i] = rand();
   }
   cudaMemcpy(fenc_g, fenc_c, size, cudaMemcpyHostToDevice);
   cudaMemcpy(fref_g, fref_c, size, cudaMemcpyHostToDevice);    
   
   cal_sad_LL<<< h, w >>>(sad_g, fenc_g, fref_g, w, h);
   cal_sad_L<<< h, w >>>(sad_g, fenc_g, fref_g, w, h);
   cal_sad<<< h, w >>>(sad_g, fenc_g, fref_g, w, h);    

   cudaMemcpy(sad_c, sad_g, w*h*sizeof(int), cudaMemcpyDeviceToHost);

   for (i = 0; i < w*h; i++)
   {
   int offset = (i>>8)*w*16*16 + (i&0xff)*16;
   sad_s[i] = 0;
   for (j = 0; j < 16; j++)
   for (k = 0; k < 16; k++)
   sad_s[i] += abs((int)fenc_c[offset+j*256*16+k] - (int)fref_c[offset+j*256*16+k]);
   }

   for (i = 0; i < w*h; i++)
   if (sad_c[i] != sad_s[i])
   break;
   if (i < w*h)
   printf("%d diff\n", i);
   else
   printf("same\n");
   err = cudaDeviceReset();
   return 0;
}

这里直接写出一些nsight给出的值
[attach]3170[/attach]
cal_sad {64, 1, 1} {256, 1, 1} 239.530(us) 75.00% 35
cal_sad_L {64, 1, 1} {256, 1, 1}, 365.536(us) 75.00% 39
cal_sad_LL {64, 1, 1} {256, 1, 1}, 1137.984(us) 100.00% 25
最后一列是每个线程的寄存器数量，倒数第三列是执行时间
程序里的 cal_sad是每次都8个字节的方式、cal_sad_L是每次读4个字节的方式，r1，r2是为了让都值时的对齐，这里因为数据比较整齐都为0，很多时候不为0。

环境
cuda sdk5.0 + vs2010 + nsight3.0 + gtx650ti

我的疑问是代码比较简单，不应该能占用那么多的寄存器。

每个线程占用寄存器数量比较多，是什么原因造成的？
有什么办法能控制每个线程的寄存器数吗？（不是设置一个上限，那样可能会影响效率）寄存器这块怎样的选择性能比较好。
现在学cuda编程，ptx有必要学吗？

system · 2013 年6 月 4 日 10:06

LZ您好：

1：您的代码中每个线程使用多少寄存器是您的具体的代码实现在编译器尽最大努力优化后的结果。一般来说，编译器编译下来用多少就是多少了。（不过注意需要是release编译，debug模式编译会多用一些寄存器的）

2：您可以强制设置每个线程不要超过多少寄存器用量，但此时，会使用local memory，可能会变的更慢。一个线程使用二三十个寄存器是常见的情况，如果不是因为寄存器数量影响occupancy，那么用多用少没什么的。
当然，如果您能够修改代码实现或者修改算法，确实可能少用一些寄存器。但鉴于在编译的时候，编译器是努力优化过的，所以仅修改具体实现未必能变得更加优化。

3：这个根据您具体的需要决定。PTX是具备更强大更底层的功能，但是一般用途，CUDA C也够用了。

大致如此，供您参考。

祝您编码顺利~

system · 2013 年6 月 4 日 10:14

楼主您好，您的写法未能充分利用您的卡。

（1）所有的卡都提供了32位的sad能力(sum of absolute difference)，并且已经充分导出。
您应该使用__sad()来计算|a - b| + c的，您的写法(s += abs(x - y))将无法正确的指导编译器使用贵卡的32位SAD能力的。

（2）3.0+卡的提供了x264的软件上的加速处理（以及硬件上的，这里不提)，这个加速处理能将x264的1B的sad再次提速4倍，但没有导出到CUDA C。
(如果您需要您可以继续跟帖，我将提供您一个__device__函数的源代码，供您使用此能力。)

（3）您使用移位+&运算，会对编译器造成充分的困扰。建议您使用union来充分提醒编译器。

（4）您使用8B读取挺好的。但建议4B足矣。甚至写的好普通的1B访问也行。这里没什么意见。

（5）在3.0+上，使用小于32个4B寄存器无意义。（除非您决定为其他的同时运行的kernel出让寄存器，但这个一般无必要）。这里建议您停止纠结。

（6）ptx可学可不学。

system · 2013 年6 月 4 日 10:32

楼主您未能及时提供响应。

因为我要吃饭了，我先给出一个kepler的辅助264加速处理的函数，您可以直接使用：
device int absdiff4(unsigned int a, unsigned int b, int previous_sum)
{
int sum;
asm volatile (“vabsdiff4.s32.u32.u32.add %0, %1, %2, %3;”:“=r”(sum):“r”(a),“r”(b),“r”(previous_sum));
return sum;
}
该语句等价于：
int sum = previous_sum;
sum += abs((a >> 0 & 0xff) - (b >> 0 & 0xff));
sum += abs((a >> 8 & 0xff) - (b >> 8 & 0xff));
sum += abs((a >> 16 & 0xff) - (b >> 16 & 0xff));
sum += abs((a >> 24 & 0xff) - (b >> 24 & 0xff));
return sum;

您可以尝试直接用absdiff4(a,b,previous_sum)来代替这个移位/按位and/减法/绝对值/加法/重复4次的整体过程。

因为楼主您未能表示您是否真的需要，
我表示此函数仅供参考，不保证您能顺利编译此函数，也不保证您能成功运行此函数。

system · 2013 年6 月 4 日 10:37

以及，如果您需要让您的代码能在3.0以下的设备上跑。您可以尝试使用32位版本的绝对差函数。这个在所有的卡上都支持。并且CUDA C里已经导出，可以直接用：

unsigned int sum = …;
unsigned int a = (unsigned int)pixel0;
unsigned int b = (unsigned int)pixel1;
sum = __usad(a,b,sum);
虽然此函数一次只能计算一组减法/绝对值/加法, 但依然比您的原始写法提速了3倍。
（仅是最好情况，不考虑访存和其他指令等因素)
以及此函数能在所有的卡上运行。

当然，您可以使用上楼层的那个absdiff4, 它可以在此3x的基础上继续提速4x. 如果您运气好（排除其他所有因素），您能看到12x的最大速度提升。但是，您需要3.x卡。

感谢来访，欢迎再来。

system · 2013 年6 月 4 日 10:43

您好，感谢的话就不多说了

（3）使用union能不能给个例子

我现在是在x264上尝试用cuda替换一些部分的实验，目前cuda的api知道的都是C接口的函数，__device__的函数知道的比较少，这种函数有参考手册吗，还是说自己去device_functions.h里自己一个个看。

x264加速代码真的很想知道

我把x264的模式选择都拖到核函数里执行，函数体不小，寄存器早早的到63了，用了很多的local memory，所以想着控制寄存器数量的问题

system · 2013 年6 月 4 日 10:48

您少看了一楼。您还是回头看吧。您可以用vabsdiff4, 它本身就是SIMD的，无需union之类的去提示编译器。

system · 2013 年6 月 4 日 11:26

你们回复速度太快了。。。我会认真多看几遍的

system · 2013 年6 月 4 日 11:33

LZ您好，我稍微说下__device__函数的事情。
__device__函数就是被kernel函数（__global__函数）调用的函数，和kernel一样，每个线程跑一份自己的代码。__device__函数一般是自己写的，比如你将kernel函数里面一部分拿出来写成__device__函数，然后在kernel原来的地方调用该函数即可等价完成原来的任务。

另外，针对某个算法/某个任务，一般需要重写算法，才能在GPU上实现较好的效果，直接将CPU原有函数内容抓到kernel里面，一般无法得到较好的效果。

其他内容请参考横扫斑竹前面的纤细叙述。

祝您好运~

system · 2013 年6 月 4 日 13:43

横扫千军斑竹给的absdiff4，我自己测了测，能得到正确的结果，有不同程度的加速。vabsdiff4在ptx指令文档（ptx_isa_3.1.pdf)里提及，是不是说还是需要了解相应数量的ptx指令？另外absdiff4能影响寄存器的使用数量（有变多也有变少的情况）

之前做的实验中，把像素一个一个读换成一次读8个字节，整个kernel速度确实提升了4倍以上（里面sad的计算占了主要的部分），然后我当时认为读global memory是主要的一个制约，这里简单的例子中似乎也有这个规律 8字节版本比 4字节版本快，一个字节的版本最慢。然后这里写了一个一次读16字节的版本，写成这样都比8字节的快

__device__ int x264_pixel_sad_16x16_LLL( pixel *pix1, pixel *pix2, int i_stride ) 
{                                                   
   int i_sum = 0;                                  
	int y;
   for (y = 0; y < 16; y++)
   {
   int r1 = ((int)pix1) & 15;
   int r2 = ((int)pix2) & 15;
   int4 p = *((int4 *)(pix1-r1));
   if (r1 != 0)
   {
   int4 p1 = *((int4 *)(pix1-r1+16));
   while (r1 > 3)
   {
   p.x = p.y; p.y = p.z; p.z = p.w; p.w = p1.x; p1.x = p1.y; p1.y = p1.z; p1.z = p1.w;
   r1 -= 4;
   }
   if (r1 > 0){
   p.x = (p.x>>r1*8 & (1<<(32-r1*8))-1) | (p.y&((1<<r1*8)-1))<<(32-r1*8);
   p.y = (p.y>>r1*8 & (1<<(32-r1*8))-1) | (p.z&((1<<r1*8)-1))<<(32-r1*8);
   p.z = (p.z>>r1*8 & (1<<(32-r1*8))-1) | (p.w&((1<<r1*8)-1))<<(32-r1*8);
   p.w = (p.w>>r1*8 & (1<<(32-r1*8))-1) | (p1.x&((1<<r1*8)-1))<<(32-r1*8);
   }
   }
   int4 q = *((int4 *)(pix2-r2));
   if (r2 != 0)
   {
   int4 q1 = *((int4 *)(pix2-r2+16));
   while (r2 > 3)
   {
   q.x = q.y; q.y = q.z; q.z = q.w; q.w = q1.x; q1.x = q1.y; q1.y = q1.z; q1.z = q1.w;
   r2 -= 4;
   }
   if (r2 > 0){
   q.x = (q.x>>r2*8 & (1<<(32-r2*8))-1) | (q.y&((1<<r2*8)-1))<<(32-r2*8);
   q.y = (q.y>>r2*8 & (1<<(32-r2*8))-1) | (q.z&((1<<r2*8)-1))<<(32-r2*8);
   q.z = (q.z>>r2*8 & (1<<(32-r2*8))-1) | (q.w&((1<<r2*8)-1))<<(32-r2*8);
   q.w = (q.w>>r2*8 & (1<<(32-r2*8))-1) | (q1.x&((1<<r2*8)-1))<<(32-r2*8);
   }
   }
   i_sum = absdiff4(p.x, q.x, i_sum);
   i_sum = absdiff4(p.y, q.y, i_sum);
   i_sum = absdiff4(p.z, q.z, i_sum);
   i_sum = absdiff4(p.w, q.w, i_sum);        
   pix1 += i_stride;                      
   pix2 += i_stride; 
   } 
   return i_sum;                                   
}

测试的时候让两个指针偏移了一些，使它们起始地址不是16的倍数（8个字节的版本少写了几个字母），都是用nsight运行release版的。

然后给我的感觉就是内存的影响比指令的影响更直接。

寄存器数量的决定靠编译器？指令的优化和内存的优化如何去均衡？就说这个sad最后应该写成那种形式比较合适？

system · 2013 年6 月 5 日 05:49

以及，楼主您可能还有个问题。昨日没说。

您的访问是具有一定步长的，我看到如下代码：
blockIdx.xw1616 + threadIdx.x16;

您的相邻2个线程间可以因此错开16B的距离。
为了保证在3.x上的较好的读写效果，请不要使用任何<8B的访问类型。
（例如1B,2B,4B)

以及，如果可能，请使用16B访问，以便访问合并。
（这个无法直接做到，因为没有16B的native type, 您可以尝试组合4个int或者2个double在struct里, 但编译器可能会错误的拆分）。

以及，在3.x上，能保证8B的大小其实就可以了。此时虽然不是合并的，但也不怎么浪费。（L2->SMX的一次传输大小为8B, 中间的空白的部分将不会传输）

这个回复是告诉您为何您的<8B访问无法取得较好的效果的可能原因。

system · 2013 年6 月 5 日 09:42

不好意思。上文有个错误。写上文的时候还不甚清醒。导致弄错了。。。。。

您这种访问模型，线程间间距较大，您要使用尽量可能大的访存模型，才能有效的不浪费L2到SMX的port. 这个宽度是64B/port。您如果连续的多个线程在此64B内只能需要NB, 那么只有N/64的cache port效率。您需要尽量加大这个，以便让N尽量增加。

以及，如果可能，建议改写为合并的访问模型。
以及，如果可能，您在您的3.0卡上可以尝试考虑使用texture/surface object。来看看有无效果。

system · 2013 年6 月 5 日 12:33

楼主您好，我和ICE尝试了一下使用vabsdiff4 + 128-bit/thread的合并访问, 基本上得到了峰值。代码如下。您可以直接使用此代码，此代码不含有任何版权，荣誉归于kapa妹子：

（如图）
（附件中)

[attach]3171[/attach]

此kernel在您相同的数据规模(4MBx2读取，64KB写入），完成了和您同样的逻辑功能(对每2个16x16区域进行绝对差求出，并累加）。
此kernel在您的卡(GTX650 Ti)上应该能跑到110us左右。请参考。

system · 2013 年6 月 6 日 03:59

照着kapa函数抄了一遍，确实能跑到110us左右（后面附带代码和执行结果截图，我用的卡GPU clock 928MHz，最基本的那种，一般比110us大）2. kapa函数里vabsdiff4换成abs函数，执行时间变化很小，寄存器从21变成30
我之前写的16字节版本（可以适应起始地址不16字节对齐的，参考宏块一般可能出现这种情况），速度相当，寄存器使用27个，使用了vabsdiff4。把vabsdiff4换成abs，寄存器数不变，速度变成160us多
8字节版本慢一些，用abs比用vabsdiff4用的寄存器少，但是还是vabsdiff4版本速度快。

测试中碰到的问题：只有前面的kernel函数执行失败（偏移了没对齐，16字节读内存会失败），后面的kernel函数也会失败。在device函数前面强行加__noinline__会影响速度。

对寄存器，指令还比较生疏。
感谢提供kapa源代码

[attach]3172[/attach]

cu代码放在附件里，直接贴超过字数限制[attach]3173[/attach]

system · 2013 年6 月 6 日 04:05

LZ您好：
“测试中碰到的问题：只有前面的kernel函数执行失败（偏移了没对齐，16字节读内存会失败），后面的kernel函数也会失败。”

CUDA里面是这样的，kernel一个挂，个个挂。您观察到的是正确的现象。

祝您好运~

system · 2013 年6 月 6 日 04:11

（1）
楼主您好，这个本身就卡在访存上。110us的数据也是根据您的访存计算出来的大致理论极限。
当您改成使用16B读取后，您的代码将是完全合并访存，的确应该展现出和kapa相似的直接结果的。

因为此代码卡在访存上，所以导致您使用普通的分布减法-绝对值-求和，或者__sad或者vabsdiff4, 将运行时间基本不变。

（2）前面的kernel运行失败，后面的也失败。这个没有办法解决。除非您选择cudaDeviceReset()，但这将会摧毁当前的上下文（包括cudaMalloc出来的任何空间), 以及下次需要重新初始化。

（3）对编译器暗示__noinline__是个不好的形式，（实际上包括暗示__forceinline__也是个不好的行为），您应该自己尊重编译器的选择，让它自行决定是否inline. 除否您有重复的把握比编译器决定的更好，并实测得到速度提升。

这也是个老生常谈的问题，很多从C++过来的人总是过度评估了自己的水平，而总是手工inline之类的(或者手工不inline)。实际上往往会取得相反效果。这个不好。

业界的一些趋势也在强制不允许手工作出是否inline的选择。一些编译器/平台不允许手工制订，而将此工作交给了编译器和平台自行决定(例如C#/CLR)。

（4）“对寄存器，指令还比较生疏。”—这个我不懂的。

（5）16B访问内存要求对齐到16B的边界。这个是常规要求。实际上如果您的显存是从cudaMalloc直接出来的，那么此kernel将满足此要求。因为它将将首行对齐到512B的边界。
如果您的行宽不是2的幂或者无法满足您的要求(例如您是1111x1111x4B, 行宽是3333B)，您可以使用cudaMallocPitch(), 它将会将每行对齐到512B。

（6）fermi不支持不对齐读写，2B类型要求对齐到2B边界。4B类型要求对齐到4B边界。8B类型要求对齐到8B边界。16B类型要求对齐到16B边界。这个请注意了。

system · 2013 年6 月 6 日 10:00

absdiff4运用到我的试验程序里有不少的提升，分享一下我的测试结果

500帧视频数据编码（1080P），模式选择部分写成kernel函数（P250，B247，I3），主要计算部分就是sad，时间统计用profile给出的结果

一个一个像素读时 kernel总时间 > 30s
换成8字节读 kernel总时间 = 4.28s
abs里__int64先强转成int kernel总时间 = 2.84s
abs替成absdiff4 kernel总时间 = 1.733s

thanks for everyone

system · 2013 年6 月 6 日 11:03

感谢您的来访。欢迎下次再来。