各位大大，请教个问题：
假设现在有两个显卡，A卡compute capacity是1.1，B卡的compute capacity是3.0
请问：执行相同的一段kernel代码（比如矩阵相乘），B卡是否会比A卡快很多？
另外请问会不会出现这种情况：
一段kernel代码，在A卡上有很大的优化余地。也就是说我可以利用shared memory等方法优化代码，使其运行速度得到很大的提升。相反，如果在B卡上，则优化余地不是很大，因为本身耗时就比A卡少，再采用同A卡相同的的优化方法，带来的提升效果不是很明显？
请问有这种情况发生吗？
LZ您好：
首先需要说明的是，计算能力版本表示的是特性的支持情况，而不是某卡具体的计算峰值情况。所以无法笼统说某代码是否高计算能力版本的卡一定比低计算能力版本的卡快（或者慢）。
其次，您的代码需要针对目标运行的计算能力版本做一些适应性的调整和优化，这样才能在目标GPU上跑的更好。单纯一段固定的代码，看似公平，其实不公平。
再次，随着计算能力版本的提升，新的特性可能会大为降低程序员的工作强度，减少手工优化的工作量。比如在1.x的硬件中，合并访问的要求很严格，但是在2.0及更高的计算能力版本的硬件中，因为cache的存在以及其他硬件的进化，很多以前不合并效率很低的访存行为，现在直接不改，效率都有很大的提升，这是一个很好的福利。比如1.x时代需要访存对齐才有较好的访存效率，而从fermi开始，对对齐的要求大为下降了，不对齐效率也比较高。此时如果按照以前对齐的处理（比如cudaMallocPitch()），带来的提升效果就不明显了。（换句话说，此时按照不优化的写法，因为硬件能力的提升，可以获得额外的自动的提升）
所以您说的这种情况是有可能的，但是具体什么情况下会发生，则需要具体根据您的问题实现和对应的硬件架构进行分析。
大致如此，祝您编码顺利~
:)感谢大大！对我帮助太大了！而且答得这么详细！OTZ…
小弟我正再看《深入浅出CUDA》，一个30+页的pdf文档。里面介绍了一个矩阵乘法的优化步骤，他的那个显卡compute capacity 1.1，我在3.0的显卡上跑了一下，感觉按着那个优化的步骤似乎没有那么显著的速度提升。
举个例子：
他的初始kernel是这样写的：
global static void matMultCUDA(const float* a, size_t lda, const float* b, size_t ldb, float* c, size_t ldc, int n)
{
const int tid = threadIdx.x;
const int bid = blockIdx.x;
const int idx = bid * blockDim.x + tid;
const int row = idx / n;
const int column = idx % n;
int i;
if(row < n && column < n)
{
float t = 0;
for(i = 0; i < n; i++)
{
t += a[row * lda + i] * b[i * ldb + column];
}
c[row * ldc + column] = t;
}
}
其中，lda,ldb,ldc是a,b,c矩阵的pitch，就是每行的长度。矩阵乘法，A,B矩阵相乘，结果存入C矩阵
他的优化方法是：
global static void matMultCUDA(const float* a, size_t lda, const float* b, size_t ldb, float* c, size_t ldc, int n)
{
extern shared float data; //启用shared memory
const int tid = threadIdx.x;
const int row = blockIdx.x;
int i, j;
for(i = tid; i < n; i += blockDim.x)
{
data[i] = a[row * lda + i];
}
__syncthreads();
for(j = tid; j < n; j += blockDim.x)
{
float t = 0;
float y = 0;
for(i = 0; i < n; i++)
{
float r;
y -= data[i] * b[i * ldb + j];
r = t - y;
y = (r - t) + y;
t = r;
}
c[row * ldc + j] = t;
}
}
他的优化思想是把row的数据存入共享存储器，避免每次都要重新读取。绿色部分是采用了Kahan’s Summation Formula算法，本质没有变，还是计算求矩阵乘积。
按照他上面的优化方法，在他的显卡上提速了一倍，而在我的显卡上基本没变化。不知道是不是大大说的原因？
我用的计时方法是：
//time_begin
cudaEvent_t start,stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start,0);
调用kernel函数…
//time_end
cudaEventRecord(stop,0);
cudaEventSynchronize(stop);
float costtime = 0;
cudaEventElapsedTime(&costtime,start,stop);
printf("Used time: %f\n",costtime);
cudaEventDestroy(start);
cudaEventDestroy(stop);
请问大大：我这中计时方法，得出来的值，也就是print出来的时间，单位是ms吗？
LZ您好：
这里不就某资料里面的具体代码实现分析和解释了，如您有兴趣，可以联系其作者讨论。
说一下一般的情况：
1：1.x下如果使用了cudaMallocPitch申请空间，那么申请的空间要比你需要使用的空间大一些，并同时保证了对齐。在1.x下对齐有利于保证访存速度，在3.0下一般也没有坏处（除非你的规模太小）。以及3.0下不对齐问题也不大。
2：使用shared memory，这个是推荐的，各版本的CUDA C Programming Guide中矩阵相乘的例子都是使用shared memory的（作为讲述shared memory的例子）。从global memory读取一次，并在block内反复使用的数据，一般推荐使用shared memory以减少global memory的访存压力。
3：在您的新卡上速度是否明显变化，还取决于其他因素的，比如计算量大小，比如线程安排是否合适等。换句话说需要针对目标架构做一些优化，才能发挥出GPU的实际效能。
4：如果您不是为了练手，而是为了进行矩阵乘法，那么建议您直接使用相应的线性代数库，那样会有更多针对性的优化，一般会比您自己写更快。
5：event计时单位为ms（毫秒）。
大致如此，祝您好运~