:)感谢大大!对我帮助太大了!而且答得这么详细!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吗?