CUDA interthread communication

Question

Я новичок в программировании CUDA, и у меня есть проблема. Я пытался написать программу, которая требует перекрестного общения; Я пробовал это всеми возможными способами, которые я нашел, и он все еще не работает. Как вы думаете, чего мне не хватает?

Ниже приведен фрагмент кода - моя вся программа. Он запускает 2 потока в том же самом блоке. Они получают вход, и выходной массив, и другую глобальную переменную для связи. Значение 0 означает, что переменная пуста и поэтому доступна для записи. В основном первый считывает элемент со входа, передает значение второму, который записывает его в выходной массив. Позже он должен быть трубопровод, с большим количеством нитей между A и B.

#include <cuda.h> 
#include <cuda_runtime.h> 

#include <stdio.h> 

#define N 1 

__global__ void link(int *in, int *out, int *pipe){ 
    int id = threadIdx.y*blockDim.x + threadIdx.x; //compute index  

    if(id == 0){  //writer thread 

     for(int index = 0;index<N;){     
      if(pipe[0]==0){    
       atomicExch(pipe, in[index++]);    
      }   
     } 
    } 
    else if(id == 1){ // reader thread  

     for(int index=0;index<N;) { 
      if(pipe[0]!=0){ 
       out[index++] = atomicExch(pipe, 0); //read and make it empty  
      }   
     }   
    } 
} 

int main(){ 
    int input[] = {8,7}; 
    int *dev_input; 
    int *dev_output; 
    int *dev_pipe; 
    int *output = (int*) malloc (N*sizeof(int)); 

    cudaMalloc((void**) &dev_input, N*sizeof(int)); 
    cudaMalloc((void**) &dev_output, N*sizeof(int)); 
    cudaMalloc((void**) &dev_pipe, 1*sizeof(int)); 
    cudaMemset(dev_pipe, 0, 1); 
    cudaMemcpy(dev_input, &input[0], N*sizeof(int), cudaMemcpyHostToDevice); 

    link<<<1, 2>>>(dev_input, dev_output, dev_pipe); 

    cudaMemcpy(output, dev_output, N*sizeof(int), cudaMemcpyDeviceToHost); 

    printf("[%d", output[0]); 
    for(int i = 1;i<N;i++) 
     printf(", %d", output[i]); 
    printf("]\n"); 
    int d = 0; 
    scanf("\n", &d); 

} 

Если читатель видеть, что труба 0 (пусто), он ставит первый элемент на нем, но писатель не может см. любые изменения, и программа переходит в тупик. Я попытался добавить __threadfence и __syncthreads, но это не помогло. Я также пытался использовать измененную память, но это тоже не сработало. Пожалуйста, помогите мне, если можете, потому что я понятия не имею, что с этим не так.

Answer 1

Опасайтесь, Нити CUDA сильно отличаются от потоков POSIX. Они работают в соответствии с парадигмой Single Instruction Multiple Threads (SIMT, см. this interesting discussion): при каждом такте каждого потока (в одном «обтекателе») выполняется одна и та же (низкоуровневая) команда.

В вашем коде поток писателя будет работать, пока поток читателя выполнит NOP s, а затем второй поток будет работать, пока первый выполнит NOP s, но они никогда не будут работать одновременно, так что вы не выиграете от массового параллельная структура графических процессоров.

Во всяком случае, чтобы ответить на ваш вопрос, ваших for петли

for(int index=0;index<N;) 

не увеличивает индекс; поэтому они представляют собой бесконечные петли. Заменить

for(int index=0;index<N;index++) 

И разделять память, что является общим для всех потоков в пределах одного блока, будет гораздо быстрее, чем глобальная память между потоками связи.