Асинхронное размещение объектов текстуры в коде с несколькими графическими процессорами

Question

У меня есть некоторый код для размещения объектов текстуры и копирования с хоста на устройство.Это просто модификация ответа здесь .Я явно не использую потоки, просто cudaSetDevice()

. Этот код работает нормально, однако, когда я запускаю Visual Profiler, я вижу, что копии памяти из Host to Array не асинхронны.Каждому из них назначается собственный поток устройств, но второй не запускается до тех пор, пока не завершится первый (работающий на 2 графических процессорах).Я пробовал это с большими изображениями, поэтому я уверен, что это не накладные расходы процессора.

Я предполагаю, что в коде есть что-то, что требует синхронности, поэтому останавливает процессор, но я не знаю, что.Что я могу сделать, чтобы сделать этот цикл асинхронным?

MCVE:

    void CreateTexture(int num_devices,float* imagedata, int nVoxelX, int nVoxelY, int nVoxelZ ,cudaArray** d_cuArrTex, cudaTextureObject_t *texImage);

int main(void)
{

int deviceCount =0 ;
cudaGetDeviceCount(&deviceCount);

int nVoxelX=512;
int nVoxelY=512;
int nVoxelZ=512;
float* image=(float*)malloc(nVoxelX*nVoxelY*nVoxelZ*sizeof(float));

cudaTextureObject_t *texImg =new cudaTextureObject_t[deviceCount];
cudaArray **d_cuArrTex = new cudaArray*[deviceCount];

CreateTexture(deviceCount,image, nVoxelX,nVoxelY, nVoxelZ,d_cuArrTex,texImg);


}

Фактическая функция:

void CreateTexture(int num_devices, float* imagedata, int nVoxelX, int nVoxelY, int nVoxelZ ,cudaArray** d_cuArrTex, cudaTextureObject_t *texImage)
{
    //size_t size_image=nVoxelX*nVoxelY*nVoxelZ;
    for (unsigned int i = 0; i < num_devices; i++){
        cudaSetDevice(i);

        //cudaArray Descriptor
        const cudaExtent extent = make_cudaExtent(nVoxelX, nVoxelY, nVoxelZ);
        cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
        //cuda Array
        cudaMalloc3DArray(&d_cuArrTex[i], &channelDesc, extent);
        //cudaCheckErrors("Texture memory allocation fail");
        cudaMemcpy3DParms copyParams = {0};


        //Array creation
        copyParams.srcPtr   = make_cudaPitchedPtr((void *)imagedata, extent.width*sizeof(float), extent.width, extent.height);
        copyParams.dstArray = d_cuArrTex[i];
        copyParams.extent   = extent;
        copyParams.kind     = cudaMemcpyHostToDevice;
        cudaMemcpy3DAsync(&copyParams);
        //cudaCheckErrors("Texture memory data copy fail");


        //Array creation End
        cudaResourceDesc    texRes;
        memset(&texRes, 0, sizeof(cudaResourceDesc));
        texRes.resType = cudaResourceTypeArray;
        texRes.res.array.array  = d_cuArrTex[i];
        cudaTextureDesc     texDescr;
        memset(&texDescr, 0, sizeof(cudaTextureDesc));
        texDescr.normalizedCoords = false;
        texDescr.filterMode = cudaFilterModePoint;
        texDescr.addressMode[0] = cudaAddressModeBorder;
        texDescr.addressMode[1] = cudaAddressModeBorder;
        texDescr.addressMode[2] = cudaAddressModeBorder;
        texDescr.readMode = cudaReadModeElementType;
        cudaCreateTextureObject(&texImage[i], &texRes, &texDescr, NULL);
        //cudaCheckErrors("Texture object creation fail");
    }
}

Answer 1

Две основные проблемы, которые я вижу в коде:

1. Ваше размещение хоста - это размещение с возможностью просмотра страниц.Асинхронность операций копирования в CUDA, где одной из целей является память хоста, требует закрепленного выделения памяти хоста.

2. В цикле создания текстур есть другие операции синхронизации.Операции выделения устройств (cudaMalloc3DArray в данном случае) синхронизируются, по моему опыту.Я не запускал тесты, чтобы определить, синхронизируется ли cudaCreateTextureObject, но я не удивлюсь, если бы это было.Поэтому моя рекомендация для асинхронности в целом состоит в том, чтобы вывести синхронизирующие операции из цикла.

В вашем случае мы можем реорганизовать ваш код следующим образом, что, по-видимому, позволяет перекрывать операции изперспектива nvprof:

$ cat t399.cu
void CreateTexture(int num_devices, float* imagedata, int nVoxelX, int nVoxelY, int nVoxelZ ,cudaArray** d_cuArrTex, cudaTextureObject_t *texImage)
{
    //size_t size_image=nVoxelX*nVoxelY*nVoxelZ;

    const cudaExtent extent = make_cudaExtent(nVoxelX, nVoxelY, nVoxelZ);
    for (unsigned int i = 0; i < num_devices; i++){
        cudaSetDevice(i);

        //cudaArray Descriptor
        cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
        //cuda Array
        cudaMalloc3DArray(&d_cuArrTex[i], &channelDesc, extent);
        //cudaCheckErrors("Texture memory allocation fail");
        }
    for (unsigned int i = 0; i < num_devices; i++){
        cudaSetDevice(i);
        cudaMemcpy3DParms copyParams = {0};
        //Array creation
        copyParams.srcPtr   = make_cudaPitchedPtr((void *)imagedata, extent.width*sizeof(float), extent.width, extent.height);
        copyParams.dstArray = d_cuArrTex[i];
        copyParams.extent   = extent;
        copyParams.kind     = cudaMemcpyHostToDevice;
        cudaMemcpy3DAsync(&copyParams);
        //cudaCheckErrors("Texture memory data copy fail");
        }
    for (unsigned int i = 0; i < num_devices; i++){
        cudaSetDevice(i);
        //Array creation End
        cudaResourceDesc    texRes;
        memset(&texRes, 0, sizeof(cudaResourceDesc));
        texRes.resType = cudaResourceTypeArray;
        texRes.res.array.array  = d_cuArrTex[i];
        cudaTextureDesc     texDescr;
        memset(&texDescr, 0, sizeof(cudaTextureDesc));
        texDescr.normalizedCoords = false;
        texDescr.filterMode = cudaFilterModePoint;
        texDescr.addressMode[0] = cudaAddressModeBorder;
        texDescr.addressMode[1] = cudaAddressModeBorder;
        texDescr.addressMode[2] = cudaAddressModeBorder;
        texDescr.readMode = cudaReadModeElementType;
        cudaCreateTextureObject(&texImage[i], &texRes, &texDescr, NULL);
        //cudaCheckErrors("Texture object creation fail");
    }
    for (unsigned int i = 0; i < num_devices; i++){
        cudaSetDevice(i);
        cudaDeviceSynchronize();
    }
}

int main(void)
{
  int deviceCount =0 ;
  cudaGetDeviceCount(&deviceCount);

  int nVoxelX=512;
  int nVoxelY=512;
  int nVoxelZ=512;
  float* image;

  cudaHostAlloc(&image, nVoxelX*nVoxelY*nVoxelZ*sizeof(float), cudaHostAllocDefault);

  cudaTextureObject_t *texImg =new cudaTextureObject_t[deviceCount];
  cudaArray **d_cuArrTex = new cudaArray*[deviceCount];

  CreateTexture(deviceCount,image, nVoxelX,nVoxelY, nVoxelZ,d_cuArrTex,texImg);
}


$ nvcc -o t399 t399.cu
$ cuda-memcheck ./t399
========= CUDA-MEMCHECK
========= ERROR SUMMARY: 0 errors
$ nvprof --print-gpu-trace ./t399
==19953== NVPROF is profiling process 19953, command: ./t399
==19953== Profiling application: ./t399
==19953== Profiling result:
   Start  Duration            Grid Size      Block Size     Regs*    SSMem*    DSMem*      Size  Throughput  SrcMemType  DstMemType           Device   Context    Stream  Name
1.55311s  90.735ms                    -               -         -         -         -  512.00MB  5.5106GB/s      Pinned       Array  Tesla P100-PCIE         1         7  [CUDA memcpy HtoA]
1.55316s  90.640ms                    -               -         -         -         -  512.00MB  5.5163GB/s      Pinned       Array   Tesla K40m (1)         2        18  [CUDA memcpy HtoA]
1.55318s  85.962ms                    -               -         -         -         -  512.00MB  5.8165GB/s      Pinned       Array  Tesla K20Xm (2)         3        29  [CUDA memcpy HtoA]
1.55320s  89.908ms                    -               -         -         -         -  512.00MB  5.5612GB/s      Pinned       Array  Tesla K20Xm (3)         4        40  [CUDA memcpy HtoA]

Regs: Number of registers used per CUDA thread. This number includes registers used internally by the CUDA driver and/or tools and can be more than what the compiler shows.
SSMem: Static shared memory allocated per CUDA block.
DSMem: Dynamic shared memory allocated per CUDA block.
SrcMemType: The type of source memory accessed by memory operation/copy
DstMemType: The type of destination memory accessed by memory operation/copy
$

моя система здесь представляет собой систему с 4 графическими процессорами с двумя графическими процессорами, висящими на каждом из 2 корневых портов.Поэтому полоса пропускания закрепленной передачи Host-> Device на PCIE Gen3 около 10 ГБ / с разделяется с точки зрения профилировщика между 2 графическими процессорами на каждом порту, но тщательное изучение времени запуска и продолжительности профилировщика для передач указывает на то, что все 4 перекрываютсяс точки зрения профилировщика.