例如，我的代码是这样的（但是它不起作用并且内核停止运行）：
__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;
//...
}
无论如何，我可以做到这一点吗？
您无法使用托管内存的当前实现来执行此操作，因为在内核运行时，托管内存要求设备以独占方式访问托管数据。在内核运行期间，主机对托管数据的访问将导致不确定的行为，通常是段错误。
但是，使用零拷贝技术（包括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位环境。
本文收集自互联网，转载请注明来源。
如有侵权，请联系[email protected] 删除。
我来说两句