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

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

您不能使用托管内存的当前实现来执行此操作，因为托管内存需要在内核运行时由设备独占访问托管数据。在内核运行期间主机访问托管数据将导致未定义的行为，通常是段错误。

然而，这应该可以使用零拷贝技术，包括volatile来自@Cicada 的建议。

这是一个工作示例：

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