Как обмениваться переменными через память хоста, заблокированную на странице - PullRequest
Новая
       57

Как обмениваться переменными через память хоста, заблокированную на странице

0 голосов
/ 26 июня 2018

Я хочу использовать Page-locked Host memory в CUDA для обмена сообщениями между ХОСТОМ и УСТРОЙСТВОМ. Позвольте мне выразить свои идеи на следующих примерах. Я не уверен, разумно ли это.

Окружение моей машины: 

 - Ubuntu 14.04.5 LTS
 - gcc (Ubuntu 4.8.4-2ubuntu1~14.04.3) 4.8.4
 - CUDA 9.1

Я разделил свою программу на четыре шага, как показано ниже:

  1. Предположим, что есть два блока, и для первого блока он выполняет некоторые вычисления, и в конце первого генерируется сигнал. блок;
  2. Когда первый блок завершает функцию, он информирует терминал ЦП, а затем соответствующие данные организуются в ЦП;
  3. Затем скопируйте данные в GPU и подайте GPU сигнал, когда копирование данных будет завершено;
  4. Второй блок в gpu запускается на основе сигнала на шаге 3.

После того, как вы узнали о том, что я хочу сделать, У меня проблема в том, что когда я изменяю данные в памяти, заблокированной на странице (в моей программе есть сигнал), кажется, что они не могут быть идентифицированы противоположное устройство.

Для этого вопроса я попробовал следующее

  • Я обнаружил, что компилятор CUDA может оптимизировать значение и сохранить значение в регистре, поэтому я не могу получить самое новое значение в ядре, поэтому я заметил PTX.
  • Я попытался использовать PTX для предотвращения оптимизации компилятором части кода, и я успешно получил сигнал в ядре, , но не смог передать устройство с формой сигнала на хост , что сбило с толку мне очень.

Часть кода моего проекта показана ниже: 

__global__ void pipeline(int *flag_a, int*flag_b, int*Input, int*Out){
    int idx = threadIdx.x;
    if (blockIdx.x == 0) {
        if (0 == idx) {
            flag_a[0] = 1;    //to generate signal in the step one 
                              //why the host cannot get the flag_a[0]==1?
        }
    }

    if (blockIdx.x == 1) {
        if (0 == idx) {
            int value = 0;
            do{
                asm volatile("ld.global.cg.u32 %0, [%1];" :"=r"(value) : "l"(&flag_b[0]));
                //receipt signal form the host generate in step 3
                //and the asm volatile to make sure I can get the newest flag_b[0]
            } while (value != 1);
        }
        __syncthreads();
        Out[idx] = Input[idx] + idx;
    }
}

int main()
{
    /*1*/
    int *flag_a, *flag_b;
    cudaHostAlloc((void**)&flag_a, sizeof(int), cudaHostAllocMapped);
    cudaHostAlloc((void**)&flag_b, sizeof(int), cudaHostAllocMapped);
    flag_a[0] = 0;
    flag_b[0] = 0;
    /*2*/
    int*Input, *Out;
    int *d_Input, *d_Out;
    int*d_float_a, *d_float_b;
    Input = (int*)malloc(sizeof(int) * 32);
    Out = (int*)malloc(sizeof(int) * 32);
    for (int i = 0; i<32; i++) {
        Input[i] = i;
    }
    memset(Out, 0, sizeof(int) * 32);

    cudaMalloc((void**)&d_Input, sizeof(int) * 32);
    cudaMemset(d_Input, 0, sizeof(int) * 32);
    cudaMalloc((void**)&d_Out, sizeof(int) * 32);
    cudaMemset(d_Out, 0, sizeof(int) * 32);

    cudaHostGetDevicePointer((void **)&d_float_a, (void *)flag_a, 0);
    cudaHostGetDevicePointer((void **)&d_float_b, (void *)flag_b, 0);

    cudaStream_t stream_kernel, stream_datacopy;
    cudaStreamCreate(&stream_kernel);
    cudaStreamCreate(&stream_datacopy);

    pipeline <<< 2, 32, 0, stream_kernel >>> (d_float_a, d_float_b, d_Input, d_Out);
    int count = 0;
    do{
        if (flag_a[0]==1){
            cudaMemcpyAsync(d_Input, Input, sizeof(int) * 32, cudaMemcpyHostToDevice, stream_datacopy);
            cudaStreamSynchronize(stream_datacopy);
            flag_b[0] = 1;  //step 3；
            break;
        }
        if (count==10)
            break;
    } while (1 != flag_a[0]);

    cudaStreamSynchronize(stream_kernel);
    cudaMemcpy(Out, d_Out, sizeof(int) * 32, cudaMemcpyDeviceToHost);
    for (int i = 0; i<32; i++) {
        printf("%d:%d\n", i, Out[i]);
    }
    // free()
    return 0;
}

Я не очень хорош в программировании на CUDA, и я не уверен, что это правильный путь для переключения сигнала между хостом и устройством, все, что я сделал, это просто попытка, и если кто-то может дать мне совет, я буду признателен , Заранее спасибо :)

1 Ответ

0 голосов
/ 02 июля 2018

В конце я удалил код части PTX и ввел код Tesla P100-PCIE (TCC mode), что может правильно запустить программу, которую я ожидал. Спасибо Роберту Кровелле, подсказку дал в комментарии.

Вот обновленный код и результаты. 

__global__ void pipeline(volatile float *flag_a, volatile float*flag_b, int*Input, int*Out)
{
    int idx = threadIdx.x;
    if (blockIdx.x == 0) {
        if (0 == idx) {
            flag_a[idx] = 1;    
        }
    }

    if (blockIdx.x == 1) {
        if (0 == idx) {
            while (!(1 == flag_b[0])) {
            }
        }
        __syncthreads();
        Out[idx] = Input[idx] + idx;
    }
}

В основной функции можно получить сигнал от ядра. 

int main()
{
    //Data definition
    pipeline << < 2, 32, 0, stream_kernel >> > (flag_a, flag_b, d_Input, d_Out);
    while (flag_a[0] == 0);
    if (flag_a[0] == 1)
    {
        std::cout << "get the flag_a[0]==1" << std::endl;
        cudaMemcpyAsync(d_Input, Input, sizeof(int) * 32, cudaMemcpyHostToDevice, stream_datacopy);
        cudaStreamSynchronize(stream_datacopy);
        flag_b[0] = 1;
        std::cout << "data transfer has finished" << std::endl;
    }

    cudaStreamSynchronize(stream_kernel);
    cudaMemcpy(Out, d_Out, sizeof(int) * 32, cudaMemcpyDeviceToHost);
    for (int i = 0; i < 32; i++) 
    {
        printf("%d:%d\n", i, Out[i]);
    }
    //free the memory;
    return 0;
}

Вот результат .

...