多个变量在CUDA中同步

我的程序有很多4字节字符串，比如“aaaa”“bbbb”“cccc”...我需要收集通过crc检查的特定字符串。多个变量在CUDA中同步

因为字符串可以通过crc检查的可能性很小，所以我不想使用非常大的缓冲区来保存所有结果。我喜欢一个接一个的结果，就像输入一样。例如，如果输入是“aaaabbbbcccc”和“BBBB”未通过CRC校验，则输出字符串应该是“aaaacccc”和output_count应该为2

代码看起来像：

__device__ 
bool is_crc_correct(char* str, int len) { 
    return true; // for simplicity, just return 'true'; 
} 

// arguments: 
// input: a sequence of 4-bytes-string, eg: aaaabbbbccccdddd.... 
__global__ 
void func(char* input, int* output, int* output_count) { 
    unsigned int index = blockDim.x*blockIdx.x + threadIdx.x; 

    if(is_crc_correct(input + 4*index)) { 
     // copy the string 
     memcpy(output + (*output_count)*4, 
       input + 4*index, 
       4); 
     // increase the counter 
     (*output_count)++; 
    } 
}

显然内存拷贝不是线程安全的，我知道atomicAdd函数可以用于++操作，但是如何使output和output_count线程安全？

来源

2015-11-10 aj3423

我相信你正在试图重建*流压缩*，尤其是*采集*操作的效率非常低。并行编程通常需要不同的思考。例如，你避免竞争，而不是试图用原子和锁来解决它们（序列化有点违背并行化的目的）。你可以使用[thrust :: copy_if]（https://thrust.github.io/doc/group__stream__compaction.html）。 – Drop

你在找什么是一个无锁的线性分配器。这样做的通常方法是通过原子级增加的累加器来索引缓冲区。例如，你的情况，下面应该工作：

__device__ 
char* allocate(char* buffer, int* elements) { 
    // Here, the size of the allocated segment is always 4. 
    // In a more general use case you would atomicAdd the requested size. 
    return buffer + atomicInc(elements) * 4; 
}

然后可以使用这样的：

__global__ 
void func(char* input, int* output, int* output_count) { 
    unsigned int index = blockDim.x*blockIdx.x + threadIdx.x; 

    if(is_crc_correct(input + 4*index)) { 
     // Reserve the output buffer. 
     char* dst = allocate(output, output_count); 
     memcpy(dst, input + 4 * index, 4); 
    } 
}

虽然这是完全线程安全的，这是不保证保护输入顺序。例如，“ccccaaaa”将是一个有效的输出。

如滴在他们的评论提到，你正在尝试做的是一个有效的数据流压缩（和推力已经可能已经提供了你所需要的）。

我上面发布的代码可以进一步优化，首先将输出字符串以变形聚合，而不是直接分配到全局缓冲区。这将减少全球原子争夺并可能导致更好的表现。有关如何执行此操作的解释，请您阅读以下文章：CUDA Pro Tip: Optimized Filtering with Warp-Aggregated Atomics。

来源

2015-11-10 10:36:38

这很聪明。 – aj3423

我可能最终会暗示这一点，但如何在内核中动态分配内存？例如：CUDA allocate memory in __device__ function

然后，您可以将共享内存数组传递给每个内核，并且在内核运行后，数组的每个元素都将指向一块动态分配的内存，或者空值。因此，在线程块运行后，您将在单个线程上运行最终的清理内核来构建最终的字符串。

来源

2015-11-10 10:17:28 endian

谢谢，我甚至不知道内核可以使用malloc/free。我记得我看到过一些编译错误，比如“不能在设备代码中使用主机函数”，也许它是旧的CUDA 3.x.现在我试着编译CUDA 6.5。 – aj3423

多个变量在CUDA中同步

回答

相关问题