CUDA设备运行时API cudaMemsetAsync无法工作。

Question

我正在尝试从内核(即所谓的“动态并行性”)调用cudaMemsetAsync。但是无论我使用什么值，它总是将内存设置为0。

这是我的测试代码:

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "cuda_device_runtime_api.h"
#include <stdio.h>

const int size = 5;

__global__ void kernel(int *c)
{
    cudaMemsetAsync(c, 0x7FFFFFFF, size * 4, NULL);
}

int main()
{
    cudaError_t cudaStatus;
    int c[size] = { 12, 12, 12, 12, 12 };
    int *dev_c = 0;

    cudaStatus = cudaSetDevice(0);
    cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
    cudaStatus = cudaMemcpy(dev_c, c, size * sizeof(int), cudaMemcpyHostToDevice);
    kernel <<< 1, 1 >>>(dev_c);
    cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
    cudaFree(dev_c);
    cudaStatus = cudaDeviceReset();

    printf("%d\n", cudaStatus);
    printf("{%d,%d,%d,%d,%d}\n", c[0], c[1], c[2], c[3], c[4]);
    return 0;
}

如果我运行它，输出结果会像这样：

>nvcc -run kernel.cu -gencode=arch=compute_35,code=\"sm_35,compute_35\" -rdc=true -lcudadevrt
kernel.cu
   Creating library a.lib and object a.exp
0
{0,0,0,0,0}

当我调用内存设置函数时，我使用值0x7FFFFFFF。我期望看到非零数值，但它总是显示为零。

这是一个错误吗？还是我做错了什么？我正在使用CUDA 8.0。

- Xiang Zhang

---

我认为这是cudaMemsetAsync中的一个bug。我已在NVIDIA内部提交了一个bug。如果我有任何重要的发现，我会进行更新。在那之前，我的建议是使用另一种方法来初始化内存，比如使用代码循环。 - Robert Crovella

澄清一下Robert所说的：编写自己的设备端memset()函数，并调用它，而不是使用API函数。或者，编写自己的内核并使用动态并行性启动它。如果我是你，我非常希望避免内核启动。 - einpoklum

@RobertCrovella 谢谢 - Xiang Zhang

Answer 1

我可以确认在我测试的CUDA 8系统中，似乎无法正常工作。

如果您想让单个线程执行操作，您可以直接在设备代码中使用memset（就像memcpy一样，它已经被支持了很长时间）。内核将在内联循环中发出一个字节大小的循环，并且每个运行线程将处理该操作。

如果您想要动态并行风格的memset操作，则最简单的方法是自己制作。在您发布的代码中，一个微不足道（而且非常轻松测试）的实现可能如下所示：

#include <cstring>
#include <cstdio>

const int size = 5;

__global__ void myMemset_kernel(void* p, unsigned char val, size_t sz)
{
    size_t tid = threadIdx.x + blockDim.x * blockIdx.x;
    unsigned char* _p = (unsigned char*)p;
    for(; tid < sz; tid += blockDim.x * gridDim.x) {
       _p[tid] = val;
    }
}

__device__ void myMemset(void* p, unsigned int val, size_t sz, cudaStream_t s=NULL)
{
    const dim3 blocksz(256,1,1); 
    size_t nblocks = (sz + blocksz.x -1) / blocksz.x;

    unsigned charval = val & 0xff;
    myMemset_kernel<<< dim3(nblocks,1,1), blocksz, 0, s >>>(p, charval, sz); 
}

__global__ void kernel(int *c)
{
    cudaStream_t s;
    cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking);
    myMemset(c, 0x7FFFFFFF, size * 4, s);
    cudaDeviceSynchronize();
}

int main()
{
    int c[size];
    int *dev_c;

    memset(&c[0], 0xffffff0c, size * sizeof(int));
    printf("{%08x,%08x,%08x,%08x,%08x}\n", c[0], c[1], c[2], c[3], c[4]);

    cudaMalloc((void**)&dev_c, size * sizeof(int));
    cudaMemcpy(dev_c, c, size * sizeof(int), cudaMemcpyHostToDevice);
    kernel <<< 1, 1 >>>(dev_c);
    cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
    cudaFree(dev_c);

    printf("{%08x,%08x,%08x,%08x,%08x}\n", c[0], c[1], c[2], c[3], c[4]);
    return 0;
}

它编译并执行以下操作：

$ nvcc -rdc=true -arch=sm_52 -o memset memset.cu -lcudadevrt
$ ./memset 
{0c0c0c0c,0c0c0c0c,0c0c0c0c,0c0c0c0c,0c0c0c0c}
{ffffffff,ffffffff,ffffffff,ffffffff,ffffffff}

最后一点--请注意上面的值并阅读这个问题和答案。在您的代码中，不可能使用cudaMemset来应用值0x7FFFFFFF。虽然值参数是一个无符号整数，cudaMemset及其相关函数工作方式类似于常规的memset，设置字节值。仅使用32位参数的最低有效字节来设置值。如果您的目标是设置32位值，则无论如何都需要制作自己版本的memset。

---

- talonmies