Синхронизация потоков Cuda C с printf или другими функциями - PullRequest
Новая
       8

Синхронизация потоков Cuda C с printf или другими функциями

0 голосов
/ 28 октября 2019

У меня проблема с идентификатором потоков во время выполнения блока. enter image description here Мне бы хотелось, чтобы предложение звучало так: «Моя временная строка печатается через графический процессор!»как вы видите (на прикрепленной фотографии более ранней), предложение было отображено неправильно, и я не знаю, как это исправить.

Код: 

__global__ void Print(const char* const __string, const size_t* const loop_repeat)
{ 
    int id_x = threadIdx.x + blockIdx.x * blockDim.x;
    while (id_x < static_cast<int>(*loop_repeat))
    {
        printf("%c", __string[id_x]);
        __syncthreads();
        id_x += blockDim.x * gridDim.x;
    }
}

int main()
{
    const char* my_string = "My temporary string is printed via GPU!";
    size_t temp{};
    temp = Get_String_Length(my_string);    //get the string length
    //GPU MEMORY ALLOCATION
    size_t* my_string_length{};
    cudaMalloc((void**)&my_string_length, sizeof(size_t));
    //COPY VALUE FROM CPU(RAM) TO GPU
    cudaMemcpy(my_string_length, &temp, sizeof(size_t), HostToDevice);
    char* string_GPU{};
    cudaMalloc((void**)&string_GPU, (temp) * sizeof(char));
    //COPY VALUE FROM CPU(RAM) TO GPU
    cudaMemcpy(string_GPU, my_string, (temp) * sizeof(char), HostToDevice);
    dim3 grid_size(1);
    dim3 block_size((temp));
    Print <<< grid_size, temp >>> (string_GPU, my_string_length);
    cudaError_t final_error = cudaDeviceSynchronize(); //for synchronization e.g Hello_World then printf
    if (final_error == cudaSuccess)
    {
        printf("%cKernel executed successfully with code: %d !%\n", NEW_LINE, final_error);
    }
    else
    {
        printf("%cKernel executed with code error: %d !\n", NEW_LINE, final_error);
    }
    cudaFree(my_string_length);
    cudaFree(string_GPU);
    return 0;
}

Буду благодарен за любую помощь.

Ответы [ 2 ]

2 голосов
/ 28 октября 2019

Основная проблема заключается в том, что вы ожидаете, что порядок выполнения потока или деформации имеет некоторый предсказуемый порядок. На самом деле это не так. Использование __syncthreads() не устраняет и не решает эту проблему.

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

$ cat t1543.cu
#include <stdio.h>
#include <stdlib.h>

__global__ void Print(const char* const __string, const size_t* const loop_repeat)
{
    int id_x = threadIdx.x + blockIdx.x * blockDim.x;
    int warp_ID = threadIdx.x>>5;
    while (id_x < static_cast<int>(*loop_repeat))
    {
        if (warp_ID == 0)
          printf("%c", __string[id_x]);
        __syncthreads();
        if (warp_ID == 1)
          printf("%c", __string[id_x]);
        __syncthreads();
        id_x += blockDim.x * gridDim.x;
    }
}

int main()
{
    const char* my_string = "My temporary string is printed via GPU!";
    size_t temp;
    temp = 40;    //get the string length
    //GPU MEMORY ALLOCATION
    size_t* my_string_length;
    cudaMalloc((void**)&my_string_length, sizeof(size_t));
    //COPY VALUE FROM CPU(RAM) TO GPU
    cudaMemcpy(my_string_length, &temp, sizeof(size_t), cudaMemcpyHostToDevice);
    char* string_GPU;
    cudaMalloc((void**)&string_GPU, (temp) * sizeof(char));
    //COPY VALUE FROM CPU(RAM) TO GPU
    cudaMemcpy(string_GPU, my_string, (temp) * sizeof(char), cudaMemcpyHostToDevice);
    dim3 grid_size(1);
    dim3 block_size((temp));
    Print <<< grid_size, temp >>> (string_GPU, my_string_length);
    cudaError_t final_error = cudaDeviceSynchronize(); //for synchronization e.g Hello_World then printf
    if (final_error == cudaSuccess)
    {
        printf("\nKernel executed successfully with code: %d !%\n", final_error);
    }
    else
    {
        printf("\nKernel executed with code error: %d !\n", final_error);
    }
    cudaFree(my_string_length);
    cudaFree(string_GPU);
    return 0;
}
$ nvcc -o t1543 t1543.cu
$ cuda-memcheck ./t1543
========= CUDA-MEMCHECK
My temporary string is printed via GPU!
Kernel executed successfully with code: 0 !%
========= ERROR SUMMARY: 0 errors
$

Обратите внимание, что я не предполагаю, что вышесказанное является хорошим стилем кодирования. Это предусмотрено для понимания проблемы. Даже этот код основывается на идее, что потоки внутри деформации будут вызывать функцию printf в предсказуемом порядке, что не гарантируется моделью программирования CUDA. Так что код действительно все еще не работает.

0 голосов
/ 28 октября 2019

Это произошло потому, что Мультипроцессор создает, управляет, планирует и выполняет потоки в группах из 32 параллельных потоков, называемых warps , как вы можете видеть в Руководстве по программированию CUDA , поэтому первые 32Thread охватывает «Моя временная строка печатается v», а оставшаяся часть - «ia GPU!». Похоже, ядро ​​поместило последнюю обертку перед первой в порядке выполнения.

...