Как читать с другого устройства GPU?

Question

Я хочу запустить следующий простой код одновременно на двух графических процессорах. Здесь у меня есть переменная A [i] = [0 1 2 3 4 5 6 7 8 9], и я хочу вычислить C [i] = A [i + 1] + A [i] + A [i-1]. Вот ответ: C [i] = [ 1 3 6 9 7 11 18 21 24 17 ]. Жирные числа неверны. Для двух устройств C [4] с устройства = 1 должен получить доступ к A [5] с устройства = 2. Как я могу сделать это самым простым способом?

Мой опыт не связан с программированием, и я предполагаю использовать multiGPU для решения уравнения PDE. Итак, я очень признателен за любую помощь в изменении этого кода для решения моей текущей проблемы.

Спасибо.

#include <stdio.h>
#include <assert.h>
#include <cuda_runtime.h>
#include <stdlib.h>
#include<time.h>

__global__ void iKernel(float *A, float *C, const int N)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < N) C[i] = A[i-1] + A[i] + A[i+1];
}


int main(int argc, char **argv)
{
    int ngpus;

    printf("> starting %s", argv[0]);

    cudaGetDeviceCount(&ngpus);
    printf(" CUDA-capable devices: %i\n", ngpus);

    ngpus = 2;

    int size = 10;

    int    iSize = size / ngpus;
    size_t iBytes = iSize * sizeof(float);

    printf("> total array size %d M, using %d devices with each device "
        "handling %d M\n", size / 1024 / 1024, ngpus, iSize / 1024 / 1024);


    // allocate device memory
    float **d_A = (float **)malloc(sizeof(float *) * ngpus);
    float **d_C = (float **)malloc(sizeof(float *) * ngpus);

    float **h_A = (float **)malloc(sizeof(float *) * ngpus);
    float **gpuRef = (float **)malloc(sizeof(float *) * ngpus);
    cudaStream_t *stream = (cudaStream_t *)malloc(sizeof(cudaStream_t) * ngpus);

    for (int i = 0; i < ngpus; i++){
        // set current device
        cudaSetDevice(i);

        // allocate device memory
        cudaMalloc((void **)&d_A[i], iBytes);
        cudaMalloc((void **)&d_C[i], iBytes);

        // allocate page locked host memory for asynchronous data transfer
        cudaMallocHost((void **)&h_A[i], iBytes);
        cudaMallocHost((void **)&gpuRef[i], iBytes);

        // create streams for timing and synchronizing
        cudaStreamCreate(&stream[i]);
    }

    dim3 block(512);
    dim3 grid((iSize + block.x - 1) / block.x);

    //h_A[ngpus][index]
    for (int i = 0; i < ngpus; i++){
        cudaSetDevice(i);
        for (int j = 0; j < iSize; j++){
            h_A[i][j] = j + i*iSize;
            printf("%d %d %d %0.8f \n", i,j,iSize, h_A[i][j]);
        }
    }
    // record start time
    double iStart = clock();

    // distributing the workload across multiple devices
    for (int i = 0; i < ngpus; i++){
        cudaSetDevice(i);

        cudaMemcpyAsync(d_A[i], h_A[i], iBytes, cudaMemcpyHostToDevice, stream[i]);

        iKernel << <grid, block, 0, stream[i] >> >(d_A[i], d_C[i], iSize);

        cudaMemcpyAsync(gpuRef[i], d_C[i], iBytes, cudaMemcpyDeviceToHost,
            stream[i]);
    }

    // synchronize streams
    for (int i = 0; i < ngpus; i++){
        cudaSetDevice(i);
        cudaStreamSynchronize(stream[i]);
    }

    for (int i = 0; i < ngpus; i++){
        for (int j = 0; j < iSize; j++){
            printf("%d %d %0.8f \n", i,j,gpuRef[i][j]);
        }
    }
    return EXIT_SUCCESS;
}

Answer 1

Вы должны загрузить области перекрытия на оба устройства. Вы не можете (легко) прочитать значения с другого устройства, поэтому вам нужно продублировать и дополнить хотя бы некоторые входные значения по мере необходимости. iSize явно недостаточен входного размера при доступе к iSize + 2 различным входным значениям.

Если бы это был алгоритм с несколькими проходами, вам нужно было бы явно выполнять копию соответствующих областей между проходами.

Попробуйте смоделировать зависимости данных формально на бумаге при попытке настроить таргетинг на системы с несколькими GPU.

Оба графических процессора могут получить доступ к памяти, выделенной с помощью cudaMallocHost, но обычно не рекомендуется использовать этот тип памяти в качестве производительности по сравнению с P * Шина 1012 *Ie довольно плохая по сравнению с локальной памятью устройства. Также имеется память, управляемая драйверами, но она также не подходит для двух графических процессоров с одним и тем же активным рабочим набором.