Возможно ли в CUDA запустить ядро ​​в потоке A между двумя событиями в потоке B?

Question

У меня есть два потока A и B в разных контекстах (разные устройства, разные потоки хоста, созданные автоматически). Мне нужно синхронизировать выполнение ядра K в потоке A с 2 различными событиями в потоке B, и поэтому K начинается сразу после запуска события 1, но не события 2. Возможно ли это в принципе? (используя cudaStreamWaitEvent?)

Мой поток В содержит цикл, так что у меня есть последовательность событий, таких как 1,2,1,2,1,2,1,2 и т. д. Но мое ядро ​​K в отдельном потоке должно начинаться только между 1 и 2, а не после того, как 2 и до 1.

Пример:

Host thread X: 
Loop: 
    kernel1(userStreamA) 
    ... 
    kernelK(userStreamA) <-- this should start only between E1 and E2 
    ... 
    kernelN(userStreamA) 

Host thread Y: 
Loop: 
    kernel1(userStreamB) 
    record(E1) 
    kernel2(userStreamB) 
    some other kernels in userStreamB 
    record(E2) 
    kernel3(userStreamB) 
    ... 
    kernelN(userStreamB) 

Answer 1

Из того, что я знаю, CUDA не может быть лучшим вариантом для использования хоста параллельные вызовы для синхронизации потоков/процессов так, как вам нужно (т. е. с помощью такого мелкозернистого управления конкретными ситуациями хоста). Если ваш вопрос ограничивался только использованием вызовов CUDA, вы можете игнорировать альтернативное решение, данное впереди.

Вы можете легко получить то, что вам требуется, с помощью OpenMP или MPI, что даст вам инструменты высокого уровня для управления потоком синхронизации и выполнения процессора потоков и процессов без ущерба для ядер GPU.

Подсказка: рассмотрите барьеры синхронизации OpenMP или сообщения ожидания/отправки/получения MPI.

Answer 2

Основная проблема, которую я вижу (что вы, вероятно, получили в своем последнем вопросе), заключается в том, что cudaStreamWaitEvent не полезен, если событие еще не «записано».

Поэтому, чтобы исправить это, я предлагаю добавить общий флаг или семафор между двумя нитями, которые будут гарантировать, что cudaStreamWaitEvent в streamA не получает выданные до соответствующего cudaEventRecord в streamB, и наоборот (мы можем использовать один и тот же флаг в этом случае).

Следующий код реализует это с помощью Pthreads:

$ cat t884.cu 
#include <pthread.h> 
#include <stdio.h> 
#include <stdlib.h> 
#include <iostream> 
#include <unistd.h> 

#define N_LOOPS 3 
#define MAX_EVENT 4 
#define MAX_STREAM 4 
#define PTHREADS 2 

#define cudaCheckErrors(msg) \ 
    do { \ 
     cudaError_t __err = cudaGetLastError(); \ 
     if (__err != cudaSuccess) { \ 
      fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \ 
       msg, cudaGetErrorString(__err), \ 
       __FILE__, __LINE__); \ 
      fprintf(stderr, "*** FAILED - ABORTING\n"); \ 
      exit(1); \ 
     } \ 
    } while (0) 

#include <time.h> 
#include <sys/time.h> 
#define USECPSEC 1000000ULL 

long long dtime_usec(unsigned long long start){ 

    timeval tv; 
    gettimeofday(&tv, 0); 
    return ((tv.tv_sec*USECPSEC)+tv.tv_usec)-start; 
} 

#define DELAY_T 1000000000ULL 
template <int type> 
__global__ void delay_kern(int i){ 

    unsigned long long time = clock64(); 
#ifdef DEBUG 
    printf("hello %d\n", type); 
#endif 
    while (clock64() < time+(i*DELAY_T)); 
} 

static int flag; 

// The thread configuration structure. 
typedef struct 
{ 
    int device; 
    int my_thread_ordinal; 
    pthread_t thread; 
    cudaError_t status; 
    cudaStream_t *streams[MAX_STREAM]; 
    cudaEvent_t *events[MAX_EVENT]; 
} 
config_t; 


// The function executed by each thread assigned with CUDA device. 
void *thread_func(void *arg) 
{ 
    // Unpack the config structure. 
    config_t *config = (config_t *)arg; 

    int device = config->device; 
    int my_thread=config->my_thread_ordinal; 
    cudaError_t cuda_status = cudaSuccess; 
    cuda_status = cudaSetDevice(device); 
    if (cuda_status != cudaSuccess) { 
     fprintf(stderr, "Cannot set focus to device %d, status = %d\n", 
      device, cuda_status); 
     config->status = cuda_status; 
     pthread_exit(NULL); 
    } 


    printf("thread %d initialized\n", my_thread); 

    switch(config->my_thread_ordinal){ 
     case 0: 
     //master thread - thread Y 
     for (int i = 0; i < N_LOOPS; i++){ 
      delay_kern<0><<<1,1,0,*(config->streams[0])>>>(1); 
      cudaEventRecord(*(config->events[0]), *(config->streams[0])); 
      flag = 1; 
      delay_kern<1><<<1,1,0,*(config->streams[0])>>>(1); 
      while (flag == 1){}; 
      cudaStreamWaitEvent(*(config->streams[0]), *(config->events[2]),0); 
      cudaEventRecord(*(config->events[1]), *(config->streams[0])); 
      delay_kern<2><<<1,1,0,*(config->streams[0])>>>(1); 
      } 
     break; 
     default: 
     //slave thread - thread X 
     for (int i = 0; i < N_LOOPS; i++){ 
      delay_kern<3><<<1,1,0,*(config->streams[1])>>>(1); 
      while (flag == 0){}; 
      cudaStreamWaitEvent(*(config->streams[1]), *(config->events[0]),0); 
      delay_kern<4><<<1,1,0,*(config->streams[1])>>>(1); 
      cudaEventRecord(*(config->events[2]), *(config->streams[1])); 
      flag = 0; 
      delay_kern<5><<<1,1,0,*(config->streams[1])>>>(1); 
      } 
     break; 
    } 
    cudaDeviceSynchronize(); 
    cudaCheckErrors("thread CUDA error"); 
    printf("thread %d complete\n", my_thread); 
    config->status = cudaSuccess; 
    return NULL; 
} 

int main(int argc, char* argv[]) 
{ 
    flag = 0; 
    const int nthreads = PTHREADS; 

    // Create workers configs. Its data will be passed as 
    // argument to thread_func. 
    config_t* configs = (config_t*)malloc(sizeof(config_t) * nthreads); 
    cudaStream_t s[MAX_STREAM]; 
    cudaEvent_t e[MAX_EVENT]; 
    cudaSetDevice(0); 
    cudaStreamCreate(s+0); 
    cudaEventCreate(e+0); 
    cudaEventCreate(e+1); 
    cudaSetDevice(1); 
    cudaStreamCreate(s+1); 
    cudaEventCreate(e+2); 
    // create a separate thread 
    // and execute the thread_func. 
    for (int i = 0; i < nthreads; i++) { 
     config_t *config = configs + i; 
     config->device = i; 
     config->my_thread_ordinal = i; 
     for (int j = 0; j < PTHREADS; j++) config->streams[j] = s+j; 
     for (int j = 0; j < PTHREADS+1; j++) config->events[j] = e+j; 
     int status = pthread_create(&config->thread, NULL, thread_func, config); 
     if (status) { 
      fprintf(stderr, "Cannot create thread for device %d, status = %d\n", 
       i, status); 
     } 
    } 
    // Wait for device threads completion. 
    // Check error status. 
    int status = 0; 
    for (int i = 0; i < nthreads; i++) { 
     pthread_join(configs[i].thread, NULL); 
     status += configs[i].status; 
    } 
    if (status) 
     return status; 

    free(configs); 

    return 0; 
} 
$ nvcc -o t884 t884.cu -lpthread 
$ time ./t884 
thread 1 initialized 
thread 0 initialized 
thread 1 complete 
thread 0 complete 

real 0m9.738s 
user 0m12.102s 
sys  0m6.235s 
$ 

Ядра шаблонный, так что мы можем различать вещи более легко в профилировщике. Само ядро ​​предназначено просто для реализации задержки ~ 1 с (это будет видно на выходе профилировщика ниже). Поскольку каждое устройство запускает в общей сложности 9 ядер, мы можем грубо наблюдать хорошее совпадение/параллелизм просто, наблюдая, что общее время выполнения составляет ~ 9 с. Я реализовал «третье» событие, идущее «по-другому», чтобы гарантировать, что запись события E2 в потоке B не произойдет до завершения ядра K. Но свидетель конкретной синхронности вы ищете между ядром K в потоке A и событие, записанное в потоке B, мы должны смотреть на это с профилировщиком:

$ nvprof --print-gpu-trace ./t884 
==14914== NVPROF is profiling process 14914, command: ./t884 
thread 0 initialized 
thread 1 initialized 
thread 1 complete 
thread 0 complete 
==14914== Profiling application: ./t884 
==14914== Profiling result: 
    Start Duration   Grid Size  Block Size  Regs* SSMem* DSMem*  Size Throughput   Device Context Stream Name 
887.60ms 974.65ms    (1 1 1)   (1 1 1)   7  0B  0B   -   - Quadro 5000 (0)   1  13 void delay_kern<int=0>(int) [188] 
887.64ms 956.79ms    (1 1 1)   (1 1 1)   8  0B  0B   -   - GeForce GT 640   2  21 void delay_kern<int=3>(int) [192] 
1.86225s 974.65ms    (1 1 1)   (1 1 1)   7  0B  0B   -   - Quadro 5000 (0)   1  13 void delay_kern<int=1>(int) [195] 
1.86225s 956.79ms    (1 1 1)   (1 1 1)   8  0B  0B   -   - GeForce GT 640   2  21 void delay_kern<int=4>(int) [199] 
2.81905s 956.79ms    (1 1 1)   (1 1 1)   8  0B  0B   -   - GeForce GT 640   2  21 void delay_kern<int=5>(int) [204] 
2.83690s 974.65ms    (1 1 1)   (1 1 1)   7  0B  0B   -   - Quadro 5000 (0)   1  13 void delay_kern<int=2>(int) [208] 
3.77584s 956.79ms    (1 1 1)   (1 1 1)   8  0B  0B   -   - GeForce GT 640   2  21 void delay_kern<int=3>(int) [212] 
3.81155s 974.65ms    (1 1 1)   (1 1 1)   7  0B  0B   -   - Quadro 5000 (0)   1  13 void delay_kern<int=0>(int) [214] 
4.78619s 974.65ms    (1 1 1)   (1 1 1)   7  0B  0B   -   - Quadro 5000 (0)   1  13 void delay_kern<int=1>(int) [219] 
4.78620s 956.79ms    (1 1 1)   (1 1 1)   8  0B  0B   -   - GeForce GT 640   2  21 void delay_kern<int=4>(int) [222] 
5.74300s 956.79ms    (1 1 1)   (1 1 1)   8  0B  0B   -   - GeForce GT 640   2  21 void delay_kern<int=5>(int) [227] 
5.76084s 974.65ms    (1 1 1)   (1 1 1)   7  0B  0B   -   - Quadro 5000 (0)   1  13 void delay_kern<int=2>(int) [231] 
6.69979s 956.79ms    (1 1 1)   (1 1 1)   8  0B  0B   -   - GeForce GT 640   2  21 void delay_kern<int=3>(int) [235] 
6.73549s 974.65ms    (1 1 1)   (1 1 1)   7  0B  0B   -   - Quadro 5000 (0)   1  13 void delay_kern<int=0>(int) [237] 
7.71014s 974.65ms    (1 1 1)   (1 1 1)   7  0B  0B   -   - Quadro 5000 (0)   1  13 void delay_kern<int=1>(int) [242] 
7.71015s 956.79ms    (1 1 1)   (1 1 1)   8  0B  0B   -   - GeForce GT 640   2  21 void delay_kern<int=4>(int) [245] 
8.66694s 956.79ms    (1 1 1)   (1 1 1)   8  0B  0B   -   - GeForce GT 640   2  21 void delay_kern<int=5>(int) [250] 
8.68479s 974.65ms    (1 1 1)   (1 1 1)   7  0B  0B   -   - Quadro 5000 (0)   1  13 void delay_kern<int=2>(int) [254] 

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. 
$ 

То, что мы наблюдаем во всех 3-х итераций что delay_kern<int=1> на одном устройстве и delay_kern<int=4> (фактически ваше Kernel K) на другом устройстве имеют почти точно такое же время начала. Это дает мне довольно уверенную уверенность в том, что поведение события приводит к желаемому поведению.

Синхронизация потоков, наложенная в этом предложении, имеет тот недостаток, что мы больше не можем иметь длинный взрыв асинхронной активности, выдаваемый хост-потоком (хотя мы все же достигаем параллелизма обработки на устройствах с требуемой синхронизацией потока). Тем не менее, я действительно не вижу способа обеспечить требуемое поведение без какого-либо элемента синхронизации потоков.