有什么办法可以在主机控制的设备代码中设置障碍？

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，通常是段错误。

这应该可以使用零拷贝技术，但是，包括来自@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 位环境。

有什么办法可以在主机控制的设备代码中设置障碍？

Is there any way I can have a barrier within Device code that is controlled by Host?

cuda

barrier