Api cudaMemsetAsync во время выполнения устройства CUDA не работает

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

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. Хотя аргумент value является целым числом без знака, cudaMemset и его родственники работают как обычные memset и установить байтовые значения. Только младший байт 32-битного аргумента используется для установки значений. Если ваша цель состоит в том, чтобы установить 32-битные значения, то вам все равно нужно будет создать собственную версию memset для этой цели.

---