cuda 矩阵按列乘法答案

【问题标题】：cuda matrix multiplication by columnscuda 矩阵按列乘法
【发布时间】：2012-12-13 00:29:26
【问题描述】：

我正在尝试在 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

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

【问题讨论】：

一些内核代码可能会有所帮助。我不能说我完全理解你的描述。但是如果您有多个线程更新给定的输出点，那么您可能需要考虑使用atomics 来避免线程之间的干扰。
所以你基本上是通过另一个矩阵的对角线条目来缩放一个矩阵的列？为什么需要多个线程来更新单个输出条目并不是很明显。除了它自己的值和第二个矩阵的恒定对角线条目之外，每个输出条目不会独立于任何东西吗？
抱歉描述含糊。我更新了问题以帮助描述问题。我担心的是线程在尝试更新输出向量时会相互干扰。我会看看原子，看看它是关于什么的。感谢您的回复！
最终结果是相同的矩阵乘法。你为什么要重新发明轮子？
这对我来说很方便，因为我的数据存储方式。这些值是从邻接列表中提取的，因此矩阵实际上是用结构和数组表示的。该结构基于列进行索引。如果我想找到一行，我必须使用 for 循环在其中一个列中搜索我想要的行......所以以传统方式执行矩阵计算效率低下，因为我不能直接从我的数据结构。

标签： cuda matrix-multiplication

【解决方案1】：

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

获取之前的a值

它通过执行 3*1 + a 的先前值来更新值

然后它将新值存储到 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
}

【讨论】：