memset cuArray для поверхностной памяти

Question

Скажем, у вас есть cuArray для привязки поверхностного объекта.

Что-то в форме:

// These are inputs to a function really.
cudaArray* d_cuArrSurf
cudaSurfaceObject_t * surfImage;

const cudaExtent extent = make_cudaExtent(width, height, depth);
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
cudaMalloc3DArray(&d_cuArrSurf, &channelDesc, extent);

// Bind to Surface
cudaResourceDesc    surfRes;
memset(&surfRes, 0, sizeof(cudaResourceDesc));
surfRes.resType = cudaResourceTypeArray;
surfRes.res.array.array  = d_cuArrSurf;

cudaCreateSurfaceObject(surfImage, &surfRes);

Теперь я хочу инициализировать это cuArray нулем. По-видимому, для объектов типа cuArray не memset. Каков был бы лучший способ сделать это? Возможно несколько вариантов, и некоторые из них могут иметь лучшие или худшие функции. Какие это варианты?

Я могу подумать о

1. выделить и обнулить память хоста и скопировать ее, используя cudaMemcpy3D().

2. создайте ядро ​​инициализации и запишите его с surf3Dwrite()

Answer 1

Можно ли показать пример этих строк?

Вот примерный пример, грубо расширяющий предыдущий пример:

$ cat t1648.cu
// Includes, system
#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <math.h>


__device__ float my_common(float *d, int width, unsigned int x, unsigned int y){

// 200 lines of common code...
  return d[y *width +x];
}




////////////////////////////////////////////////////////////////////////////////
// Kernels
////////////////////////////////////////////////////////////////////////////////
//! Write to a cuArray using surface writes
//! @param gIData input data in global memory
////////////////////////////////////////////////////////////////////////////////
__global__ void WriteKernel(float *gIData, int width, int height,
                                       cudaSurfaceObject_t outputSurface)
{
    // calculate surface coordinates
    unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
    unsigned int z = blockIdx.z*blockDim.z + threadIdx.z;
    // read from global memory and write to cuarray (via surface reference)
    surf3Dwrite(my_common(gIData, width, x, y),
                outputSurface, x*4, y, z, cudaBoundaryModeTrap);
}

__global__ void WriteKernel(float *gIData, int width, int height,
                                       float *out)
{
    // calculate coordinates
    unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;

    // read from global memory and write to global memory
    out[y*width+x] = my_common(gIData, width, x, y);
}

__global__ void ReadKernel(float tval, cudaSurfaceObject_t outputSurface)
{
    // calculate surface coordinates
    unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
    unsigned int z = blockIdx.z*blockDim.z + threadIdx.z;;
    // read from global memory and write to cuarray (via surface reference)
    float val;
    surf3Dread(&val,
                outputSurface, x*4, y, z, cudaBoundaryModeTrap);
    if (val != tval) printf("oops\n");
}


////////////////////////////////////////////////////////////////////////////////
// Program main
////////////////////////////////////////////////////////////////////////////////
int main(int argc, char **argv)
{
    printf("starting...\n");


    unsigned width = 256;
    unsigned height = 256;
    unsigned depth = 256;
    unsigned int size = depth*width * height * sizeof(float);

    // Allocate device memory for result
    float *dData = NULL;
    cudaMalloc((void **) &dData, size);

    // Allocate array and copy image data
    float *out, *h_out;
    h_out = new float[height*width*depth];
    float tval = 1.0f;
    for (int i = 0; i < height*width*depth; i++) h_out[i] = tval;
    cudaArray* d_cuArrSurf;
    cudaSurfaceObject_t  surfImage;

    const cudaExtent extent = make_cudaExtent(width, height, depth);
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
    cudaMalloc3DArray(&d_cuArrSurf, &channelDesc, extent);

    // Bind to Surface
    cudaResourceDesc    surfRes;
    memset(&surfRes, 0, sizeof(cudaResourceDesc));
    surfRes.resType = cudaResourceTypeArray;
    surfRes.res.array.array  = d_cuArrSurf;

    cudaCreateSurfaceObject(&surfImage, &surfRes);

    cudaMalloc(&out, size);
    cudaMemcpy(out, h_out, size, cudaMemcpyHostToDevice);
    dim3 dimBlock(8, 8, 8);
    dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1);
    // initialize array
    cudaMemcpy3DParms p = {0};
    p.srcPtr = make_cudaPitchedPtr(out, width*sizeof(out[0]), width, height);
    p.srcPos = make_cudaPos(0,0,0);
    p.dstArray = d_cuArrSurf;
    p.dstPos = make_cudaPos(0,0,0);
    p.extent = make_cudaExtent(width, height, 1);
    p.kind   = cudaMemcpyDefault;
    for (int i = 0; i < depth; i++){
      cudaMemcpy3D(&p);
      p.dstPos = make_cudaPos(0,0, i+1);}

    ReadKernel<<<dimGrid, dimBlock>>>(tval, surfImage);
    WriteKernel<<<dimGrid, dimBlock>>>(dData, width, height, surfImage);
    WriteKernel<<<dimGrid, dimBlock>>>(dData, width, height, out);
    cudaDeviceSynchronize();
}
$ nvcc -o t1648 t1648.cu
$ cuda-memcheck ./t1648
========= CUDA-MEMCHECK
starting...
========= ERROR SUMMARY: 0 errors
$

(Общий) экстент выше 256x256x256. Поэтому я решил выполнить передачу 256x256 (экстент на каждую передачу) (в основном каждый z-фрагмент) в течение 256 итераций cudaMemcpy3D. Похоже, что он прошел тест sniff.

Я использовал 1 в качестве начального значения для памяти устройства здесь "просто потому что". Если вы хотите сделать это быстрее и инициализировать до нуля, пропустите копию host-> device и просто используйте cudaMemset, чтобы инициализировать линейную память (источник для передачи 3D) в ноль.