Есть ли способ, которым я могу иметь барьер в коде устройства, который контролируется Host?

Question

Например, мой код что-то вроде этого (но это не работает, и ядро ​​киоски):

__device__ __managed__ int x; 

__global__ void kernel() { 

    // do something 

    while(x == 1); // a barrier 

    // do the rest 
} 

int main() { 
    x = 1; 
    kernel<<< 1, 1 >>>(); 
    x = 0; 

    //... 
} 

есть в любом случае я могу это сделать?

Answer 1

Вы не можете сделать это с текущей реализацией управляемой памяти, так как для управляемой памяти требуется эксклюзивный доступ к управляемым данным на устройстве, когда ядра запущены. Хост-доступ к управляемым данным в течение времени, когда ядра работают, приведет к undefined behavior, как правило, к сбою.

Это должно быть возможно с использованием методов нулевой копии, включая рекомендацию от 44C.

Вот обработанный пример:

$ cat t736.cu 
#include <stdio.h> 
#include <unistd.h> 

__global__ void mykernel(volatile int *idata, volatile int *odata){ 

    *odata = *idata; 
    while (*idata == 1); 
    *odata = *idata+5; 
} 

int main(){ 

    int *idata, *odata; 

    cudaHostAlloc(&idata, sizeof(int), cudaHostAllocMapped); 
    cudaHostAlloc(&odata, sizeof(int), cudaHostAllocMapped); 

    *odata = 0; 
    *idata = 1; // set barrier 
    mykernel<<<1,1>>>(idata, odata); 
    sleep(1); 
    printf("odata = %d\n", *odata); // expect this to be 1 
    *idata = 0; // release barrier 
    sleep(1); 
    printf("odata = %d\n", *odata); // expect this to be 5 
    cudaDeviceSynchronize(); // if kernel is hung, we will hang 
    return 0; 
} 


$ nvcc -o t736 t736.cu 
$ cuda-memcheck ./t736 
========= CUDA-MEMCHECK 
odata = 1 
odata = 5 
========= ERROR SUMMARY: 0 errors 
$ 

выше предполагает немного среды Linux 64.