有什麼辦法可以在由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，通常是seg故障。

這應該可以使用零拷貝技術，但是，包括來自@Cicada的volatile推薦。

這裏有一個工作例如：

$ 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位環境。