cuda矩阵逐列乘法

Question

我正在尝试在cuda中做矩阵乘法。我的实现与cuda示例不同。

cuda示例(来自cuda样本)通过将第一矩阵行中的每个值乘以第二矩阵列中的每个值来执行矩阵乘法，然后对乘积求和并将其存储在第一矩阵行的索引处的输出向量中。

我的实现将第一个矩阵的列中的每个值乘以第二个矩阵的行的单个值，其中行索引=列索引。然后，它在全局内存中有一个输出向量，该向量的每个索引都已更新。

cuda示例实现可以让单个线程更新输出向量中的每个索引，而我的实现可以让多个线程更新每个索引。

我得到的结果只显示了其中的一些值。例如，如果我让它进行4次迭代更新，它只会进行2次或1次更新。

我认为线程可能相互干扰，因为它们都试图写入全局内存中向量的相同索引。因此，当一个线程写入索引时，另一个线程可能无法插入它的值并更新索引？

只是想知道这个评估是否有意义。

例如。将以下两个矩阵相乘：

[3 0 0 2         [1       [a
 3 0 0 2    x     2   =    b
 3 0 0 0          3        c
 0 1 1 0]         4]       d]

Cuda示例使用4个线程执行矩阵乘法，其中a、b、c、d存储在全局内存中：

Thread 0:   3*1 + 0*2 + 0*3 + 2*4 = a
Thread 1:   3*1 + 0*2 + 0*3 + 2*4 = b
Thread 2:   3*1 + 0*2 + 0*3 + 0*4 = c
Thread 3:   0*1 + 1*2 + 1*3 + 0*4 = d

我的实现如下所示：

a = b = c = d = 0

Thread 0:
3*1 += a
3*1 += b
3*1 += c
0*1 += d

Thread 1:
0*2 += a
0*2 += b
0*2 += c
1*2 += d

Thread 2:
0*3 += a
0*3 += b
0*3 += c
1*3 += d

Thread 3:
2*4 += a
2*4 += b
0*4 += c
0*4 += d

因此，所有四个线程都可能同时尝试更新其中一个索引。

Answer 1

为了解决此问题，我使用atomicAdd执行+=操作。例如，当一个线程执行3*1 += a的操作时，它会做三件事。

1. 它获取a
2. 的上一个值它通过执行3*1 + a的上一个值来更新该值e217
3. 然后将新值存储到a

中

通过使用atomicAdd，它可以保证这些操作可以由线程执行，而不会受到其他线程的干扰。如果未使用atomicAdd，则thread0可以获取a的先前值，并且在thread0更新该值时，thread1可以获取a的先前值并执行自己的更新。这样，+=操作就不会发生，因为线程无法完成它们的操作。

如果使用a += 3*1而不是atomicAdd(&a，3*1)，那么thread1有可能在thread0完成它正在做的事情之前干扰并更改thread0的值。它创建了一个竞态条件。

atomicAdd是+=操作。您将使用以下代码来执行该操作：

__global__ void kernel(){    
int a = 0;   
atomicAdd(&a, 3*1);  //is the same as a += 3*1
}