【发布时间】:2013-08-27 18:52:47
【问题描述】:
我必须计算矩阵的元素(见下图6*8 示例)。矩阵元素具有依赖关系,因此必须首先计算 t1 的值,然后计算依赖于 t1 的 t2,然后计算依赖于 t2 值的 t3 的值,依此类推.
如何计算CUDA中的矩阵元素?我应该只使用一个内核调用并在 ____syncthreads() 的帮助下计算所有值,还是应该在单独的内核中执行次对角线中的矩阵元素的计算?
【问题讨论】:
标签: cuda
我必须计算矩阵的元素(见下图6*8 示例)。矩阵元素具有依赖关系,因此必须首先计算 t1 的值,然后计算依赖于 t1 的 t2,然后计算依赖于 t2 值的 t3 的值,依此类推.
如何计算CUDA中的矩阵元素?我应该只使用一个内核调用并在 ____syncthreads() 的帮助下计算所有值,还是应该在单独的内核中执行次对角线中的矩阵元素的计算?
【问题讨论】:
标签: cuda
正如您所描述的,合乎逻辑的方法是为计算的每个阶段启动一个单独的内核。在非平凡规模的问题中,“计算前沿”的规模会迅速增长,因此当解决方案在域中传播时,可以获得一定程度的计算效率。
“最好”的方法可能不是扫过域,而是迭代地解决整个域,直到解决方案收敛。 Jeong and Whittaker 发表了一篇非常好的论文,介绍了用于求解平稳 Eikonal 方程的迭代标签校正方法(这是一个经典的迎风扫掠计算,类似于您的矩阵图)。在他们的方法中,计算网格被分解为块,并且包含尚未收敛的值的每个块被重新计算,直到它收敛。当一个特征跨越子块边界时,任何依赖于变化值的值都被重新标记为未收敛,并且该过程继续进行,直到整个域收敛。
您可以在 CUDA GPU here 上观看该算法的 Youtube 视频
【讨论】:
您可以使用相同的内核来计算 t1 值,然后根据 t1 计算 t2 值,依此类推。这个内核可以递归调用,根据传递给它的一些参数对不同的值进行操作。
如果 t1 中的元素不相互依赖,则无需在内核内部使用 __syncthreads(),因为在每次内核调用之后都有一个隐式屏障。
但是,如果它们是依赖的,则必须在内核调用中使用 __syncthreads()。
【讨论】: