CUDA 设备运行时 api cudaMemsetAsync 不起作用

Question

我正在尝试从内核调用 cudaMemsetAsync（所谓的 "dynamic parallelism"）。但无论我使用什么值，它总是将内存设置为 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

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}

最后一点 -- 注意上面的值并阅读 this question and answer。在您的代码中，不可能使用 cudaMemset 来应用值 0x7FFFFFFF。虽然 value 参数是一个无符号整数，但 cudaMemset 及其亲戚像常规 memset 一样工作并设置 byte 值。只有 32 位参数的最低有效字节用于设置值。如果您的 objective 是设置 32 位值，那么无论如何您都需要为此目的制作自己的 memset 版本。

CUDA 设备运行时 api cudaMemsetAsync 不起作用

CUDA device runtime api cudaMemsetAsync doesn't work

cuda

dynamic-parallelism