在 CUDA 中同步多个变量
more than one variable to synchronize in CUDA
我的程序有很多 4 字节的字符串,例如 "aaaa" "bbbb" "cccc"...我需要收集通过 crc 检查的特定字符串。
因为一个字符串通过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_count 线程安全?
我可能会因为这个建议而下地狱,但是在内核中动态分配内存怎么样?请参阅此 question/answer 示例:CUDA allocate memory in __device__ function
然后您将向每个内核传递一个共享内存数组,在内核拥有 运行 之后,数组的每个元素要么指向一块动态分配的内存,要么指向 NULL。因此,在您的线程块具有 运行 之后,您将在单个线程上 运行 一个最终的清理内核,以构建最终的字符串。
您正在寻找的是无锁线性分配器。这样做的通常方法是使用一个原子递增的累加器来索引缓冲区。例如,在您的情况下,以下内容应该有效:
__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" 将是一个有效的输出。
正如 Drop 在他们的评论中提到的那样,您正在尝试做的实际上是流压缩(并且 Thrust 可能已经提供了您所需要的)。
我上面发布的代码可以通过首先通过 warp 聚合输出字符串 而不是直接分配到全局缓冲区中来进一步优化。这将减少全局原子争用并可能导致更好的性能。有关如何执行此操作的说明,我邀请您阅读以下文章:CUDA Pro Tip: Optimized Filtering with Warp-Aggregated Atomics.
我的程序有很多 4 字节的字符串,例如 "aaaa" "bbbb" "cccc"...我需要收集通过 crc 检查的特定字符串。
因为一个字符串通过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_count 线程安全?
我可能会因为这个建议而下地狱,但是在内核中动态分配内存怎么样?请参阅此 question/answer 示例:CUDA allocate memory in __device__ function
然后您将向每个内核传递一个共享内存数组,在内核拥有 运行 之后,数组的每个元素要么指向一块动态分配的内存,要么指向 NULL。因此,在您的线程块具有 运行 之后,您将在单个线程上 运行 一个最终的清理内核,以构建最终的字符串。
您正在寻找的是无锁线性分配器。这样做的通常方法是使用一个原子递增的累加器来索引缓冲区。例如,在您的情况下,以下内容应该有效:
__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" 将是一个有效的输出。
正如 Drop 在他们的评论中提到的那样,您正在尝试做的实际上是流压缩(并且 Thrust 可能已经提供了您所需要的)。
我上面发布的代码可以通过首先通过 warp 聚合输出字符串 而不是直接分配到全局缓冲区中来进一步优化。这将减少全局原子争用并可能导致更好的性能。有关如何执行此操作的说明,我邀请您阅读以下文章:CUDA Pro Tip: Optimized Filtering with Warp-Aggregated Atomics.