parallel-processing - 金属非原子平行还原
问题描述
我刚刚进入并行减少的世界。我正在尝试用 Metal 来实现这一点。我已经能够使用原子类型和 atomic_fetch_* 函数成功编写一个简单的版本。
我现在正试图用非原子变量做类似的事情,一个简单的结构。
定义如下:
struct Point2
{
int x;
int y;
};
使用这样的核函数:
kernel void compareX(const device Point2 *array [[ buffer(0) ]],
device Point2 *result [[ buffer(1) ]],
uint id [[ thread_position_in_grid ]],
uint tid [[ thread_index_in_threadgroup ]],
uint bid [[ threadgroup_position_in_grid ]],
uint blockDim [[ threads_per_threadgroup ]]) {
threadgroup Point2 shared_memory[THREADGROUP_SIZE];
uint i = bid * blockDim + tid;
shared_memory[tid] = array[i];
threadgroup_barrier(mem_flags::mem_threadgroup);
// reduction in shared memory
for (uint s = 1; s < blockDim; s *= 2) {
if (tid % (2 * s) == 0 && shared_memory[tid + s].x < shared_memory[tid].x) {
shared_memory[tid] = shared_memory[tid + s];
}
threadgroup_barrier(mem_flags::mem_threadgroup);
}
if (0 == tid ) {
///THIS IS NOT CORRECT
result[0] = shared_memory[0];
}
}
我首先认为内存复制到缓冲区/从缓冲区出现问题,但我已经验证到/从 CPU/GPU 与结构正常工作。然后我意识到它与跨线程组同步有关。
CUDA 有很多示例/文档,但其他任何东西都很少,而且 CUDA 并不总是能很好地转换为 Metal。
在没有原子类型的情况下获得跨线程组同步的方法是什么?
内核正在尝试获取输入数组中的最小点。现在,由于写入顺序,结果在执行过程中会发生变化。
解决方案
这可能不是最正确或最好的解决方案。但这是我在为此苦苦挣扎一段时间后想出的。如果其他人找到更好的解决方案,请发布!这也可能与不同版本的 Metal 过时。
我首先尝试_atomic<T>
在我的结构上使用 Metal 语言中包含的那个。这应该有效。在为此苦苦挣扎之后,我终于检查了文档并意识到模板目前被苹果限制为 bool's、int's 和 uint's。
然后,我尝试使用原子 int 来“锁定”关键比较部分,但实际上并没有成功保护关键部分。我可能在这个实现中做错了什么,并且可以看到它工作。
然后我简化为返回索引而不是点,这允许我再次在结果上使用 atomic_int。有点作弊,并且仍然使用原子来减少。但它有效,所以我可以继续前进。
这是内核现在的样子:
kernel void compareX(const device Point2 *array [[ buffer(0) ]],
device atomic_int *result [[ buffer(1) ]],
uint id [[ thread_position_in_grid ]],
uint tid [[ thread_index_in_threadgroup ]],
uint bid [[ threadgroup_position_in_grid ]],
uint blockDim [[ threads_per_threadgroup ]]) {
threadgroup int shared_memory[THREADGROUP_SIZE];
uint i = bid * blockDim + tid;
shared_memory[tid] = i;
threadgroup_barrier(mem_flags::mem_threadgroup);
for (uint s = 1; s < blockDim; s *= 2) {
if (tid % (2 * s) == 0) {
// aggregate the index to our smallest value in shared_memory
if ( array[shared_memory[tid + s]].x < array[shared_memory[tid]].x) {
shared_memory[tid] = shared_memory[tid + s];
}
}
threadgroup_barrier(mem_flags::mem_threadgroup);
}
if (0 == tid ) {
// get the current index so we can test against that
int current = atomic_load_explicit(result, memory_order_relaxed);
if( array[shared_memory[0]].x < array[current].x) {
while(!atomic_compare_exchange_weak_explicit(result, ¤t, shared_memory[0], memory_order_relaxed, memory_order_relaxed)) {
// another thread won. Check if we still need to set it.
if (array[shared_memory[0]].x > array[current].x) {
// they won, and have a smaller value, ignore our best result
break;
}
}
}
}
}
推荐阅读
- linux - 如何在 PSFTP 中使用 MPUT 来放置具有 777 权限的文件
- c - 如何使用socket可以扩展框架?
- typescript - typescript 静态类型检查 ES 模块导出
- javascript - 如何合并列表中 3 个不同对象的字符串?
- rest - FHIR 服务返回的字符串显示在变量中并且具有长度但在日志记录中不打印(或打印为空格)
- database - PostgreSQL 中的循环外键
- javascript - 来自数据库的自动填充文本框值
- android - Android 11 模拟器使 PCI 总线不适用于 hda
- azure-devops - Azure Visual Studio 2019 无法推送到 Azure DevOps Git
- c# - editorconfig - PascalCase 方法命名的正确 IDE 建议