理解 CUDA 设备代码中 memset 的使用答案

【问题标题】：Understanding the use of memset in CUDA device code理解 CUDA 设备代码中 memset 的使用
【发布时间】：2019-01-05 19:47:24
【问题描述】：

我有一个线性 int 数组 arr，它位于 CUDA 全局内存上。我想将arr 的子数组设置为定义的值。子数组的起始索引由starts 数组给出，而每个子数组的长度在counts 数组中给出。

我想要做的是设置子数组i 的值从starts[i] 开始，一直到counts[i] 到值starts[i]。也就是操作是：

arr[starts[i]: starts[i]+counts[i]] = starts[i]

我想在内核中使用memset() 来设置值。但是，它没有被正确写入（数组元素被分配了一些随机值）。我使用的代码是：

#include <stdlib.h>
__global__ void kern(int* starts,int* counts, int* arr,int* numels)
{
    unsigned int idx = threadIdx.x + blockIdx.x*blockDim.x;

    if (idx>=numels[0])
        return;

    const int val = starts[idx];
    memset(&arr[val], val, sizeof(arr[0])*counts[idx]) ;
    __syncthreads();
}

请注意numels[0] 包含starts 数组中的元素数。

我已经使用cuda-memcheck() 检查了代码，但没有收到任何错误。如果相关，我正在使用PyCUDA。由于我正在学习CUDA，我可能在这里误解了memset的用法。

您能否建议一种纠正此问题的方法？或其他高效方式进行此操作。

P.S：我知道thrust::fill() 可能可以很好地做到这一点，但是由于我正在学习 CUDA，所以我想知道如何在不使用外部库的情况下做到这一点。

【问题讨论】：

memset 作用于unsigned chars，而不是ints。
@molbdnilo 所以，我需要先将其转换为unsigned char？
有没有其他有效的方法在CUDA上做同样的操作？

标签： c++ cuda pycuda

【解决方案1】：

CUDA 设备代码中的 memset 和 memcpy 实现发出简单的串行字节值操作（请注意，memset 不能设置字节值以外的任何内容，这可能会导致您看到的问题是您尝试的值设置大于 8 位）。

您可以将 memset 调用替换为以下内容：

const int val = starts[idx];
//memset(&arr[val], val, sizeof(arr[0])*counts[idx]) ;
for(int i = 0; i < counts[idx]; i++)
    arr[val + i] = val;

该代码的性能可能会比内置的 memset 更好。

还要注意，内核末尾的__syncthreads() 调用既是不必要的，也是潜在的死锁来源，应该删除。请参阅here 了解更多信息。

【讨论】：