【IT168 技术】前面介绍的计算平方和的程序，似乎没有什么实用价值。所以我们的第二个 CUDA 程序，要做一个确实有(某些)实用价值的程序，也就是进行矩阵乘法。而且，这次我们会使用浮点数。
虽然矩阵乘法有点老套，不过因为它相当简单，而且也可以用来介绍一些有关 CUDA 的有趣性质。
矩阵乘法
为了单纯起见，我们这里以方形的矩阵为例子。基本上，假设有两个矩阵 A 和 B，则计算 AB = C 的方法如下：
for(j = 0; j < n; j++) {
C[i][j] = 0;
for(k = 0; k < n; k++) {
C[i][j] += A[i][k] * B[k][j];
}
}
}
一开始，我们先准备好产生数据、设定 CUDA 等等的工作：
{
float *a, *b, *c, *d;
int n = 1000;
if(!InitCUDA()) return 0;
a = (float*) malloc(sizeof(float) * n * n);
b = (float*) malloc(sizeof(float) * n * n);
c = (float*) malloc(sizeof(float) * n * n);
d = (float*) malloc(sizeof(float) * n * n);
srand(0);
matgen(a, n, n);
matgen(b, n, n);
clock_t time = matmultCUDA(a, n, b, n, c, n, n);
matmult(a, n, b, n, d, n, n);
compare_mat(c, n, d, n, n);
double sec = (double) time / CLOCKS_PER_SEC;
printf("Time used: %.2f (%.2lf GFLOPS)\n", sec,
2.0 * n * n * n / (sec * 1E9));
return 0;
}
InitCUDA 函式和第一个 CUDA 程序一样，可以直接参考前面的文章。以下是上面用到的一些其它的函式：
产生矩阵：
{
int i, j;
for(i = 0; i < n; i++) {
for(j = 0; j < n; j++) {
a[i * lda + j] = (float) rand() / RAND_MAX +
(float) rand() / (RAND_MAX * RAND_MAX);
}
}
}
这个函式只是利用随机数生成器把矩阵填满 0 ~ 1 之间的数字。特别注意到因为 C 语言中无法声明变动大小的二维矩阵，所以我们使用 i * lda + j 的方式。
进行矩阵乘法：
float* c, int ldc, int n)
{
int i, j, k;
for(i = 0; i < n; i++) {
for(j = 0; j < n; j++) {
double t = 0;
for(k = 0; k < n; k++) {
t += a[i * lda + k] * b[k * ldb + j];
}
c[i * ldc + j] = t;
}
}
}
这是以 CPU 进行矩阵乘法、用来进行验证答案正确与否的程序。特别注意到它用 double 来储存暂时的计算结果，以提高精确度。
验证结果：
const float* b, int ldb, int n)
{
float max_err = 0;
float average_err = 0;
int i, j;
for(i = 0; i < n; i++) {
for(j = 0; j < n; j++) {
if(b[i * ldb + j] != 0) {
float err = fabs((a[i * lda + j] -
b[i * ldb + j]) / b[i * ldb + j]);
if(max_err < err) max_err = err;
average_err += err;
}
}
}
printf("Max error: %g Average error: %g\n",
max_err, average_err / (n * n));
}
这个函式计算两个矩阵的最大相对误差和平均相对误差，并把结果印出来。
最后是 CUDA 的矩阵乘法的部份：
clock_t matmultCUDA(const float* a, int lda,
const float* b, int ldb, float* c, int ldc, int n)
{
float *ac, *bc, *cc;
clock_t start, end;
start = clock();
cudaMalloc((void**) &ac, sizeof(float) * n * n);
cudaMalloc((void**) &bc, sizeof(float) * n * n);
cudaMalloc((void**) &cc, sizeof(float) * n * n);
cudaMemcpy2D(ac, sizeof(float) * n, a, sizeof(float) * lda,
sizeof(float) * n, n, cudaMemcpyHostToDevice);
cudaMemcpy2D(bc, sizeof(float) * n, b, sizeof(float) * ldb,
sizeof(float) * n, n, cudaMemcpyHostToDevice);
int blocks = (n + NUM_THREADS - 1) / NUM_THREADS;
matMultCUDA<<<blocks * n, NUM_THREADS>>>
(ac, n, bc, n, cc, n, n);
cudaMemcpy2D(c, sizeof(float) * ldc, cc, sizeof(float) * n,
sizeof(float) * n, n, cudaMemcpyDeviceToHost);
cudaFree(ac);
cudaFree(bc);
cudaFree(cc);
end = clock();
return end - start;
}
这个函式相当单纯，就是在显卡内存中配置存放矩阵的内存，然后把主内存中的矩阵数据复制到显卡内存上。不过，因为我们的矩阵乘法函式可以指定 pitch(即 lda、ldb、和 ldc)，所以如果用一般的 cudaMemcpy 函式来复制内存的话，会需要每个 row 都分开复制，那会需要呼叫很多次 cudaMemcpy 函式，会使效率变得很差。因此，在这里我们用了一个新的 cudaMemcpy2D 函式，它是用来复制二维数组，可以指定数组的 pitch。这样就可以透过一次函数调用就可以了。
进行计算的 kernel 如下：
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;
}
}
这个函式一开始先从 bid 和 tid 计算出这个 thread 应该计算的 row 和 column，在判断 row 和 column 在范围内之后，就直接进行计算，并把结果写到 c 矩阵中，是非常单纯的函式。
在 GeForce 8800GT 上实际执行的结果如下：
Max error: 2.01484e-006 Average error: 3.36637e-007
Time used: 1.1560 (1.73 GFLOPS)
可以看到两个问题：
很明显的，执行效率相当低落。
最大相对误差偏高。理想上应该要低于 1e-6。
计算结果的误差偏高的原因是，在 CPU 上进行计算时，我们使用 double(即 64 bits 浮点数)来累进计算过程，而在 GPU 上则只能用 float(32 bits 浮点数)。在累加大量数字的时候，由于累加结果很快会变大，因此后面的数字很容易被舍去过多的位数。
由于 CUDA 的浮点数运算，在进行加、减、乘法时是符合 IEEE 754 规定的精确度的，因此，我们可以利用 Kahan's Summation Formula 来提高精确度。把程序改成：
float t = 0;
float y = 0;
for(i = 0; i < n; i++) {
float r;
y -= a[row * lda + i] * b[i * ldb + column];
r = t - y;
y = (r - t) + y;
t = r;
}
}
修改后的程序的执行结果是：
Max error: 1.19209e-007 Average error: 4.22751e-008
Time used: 1.1560 (1.73 GFLOPS)
可以看到相对误差有很大的改善，效率则没什么变化。
由于 Kahan's Summation Formula 需要的运算量提高，但是效率却没有什么改变，可以看出这个 kernel 主要的瓶颈应该是在内存的存取动作上。这是因为有大量的内存读取是重复的。例如，矩阵 a 的一个 row 在每次进行计算时都被重复读入，但这是相当浪费的。这样的计算方式，总共需要读取 2*n3 次内存。如果让一个 row 只需要读入一次的话，就可以减到为 n3+n2 次。
