cuda异常后的内存数据状态

Question

在CUDA应用程序抛出异常后,CUDA文档不清楚内存数据如何变化.

例如,内核启动(动态)遇到异常(例如Warp超出范围的地址),将停止当前的内核启动.在此之后,设备上的数据(例如__device__变量)是否仍然保留,或者它们是否与例外一起被删除？

一个具体的例子是这样的:

1. CPU启动内核
2. 内核将__device__ variableA的值更新为5,然后崩溃
3. CPU memcpy从设备到主机的variableA的值,在这种情况下CPU获得的值是5,还是其他什么？

有人能说明这背后的理由吗？

Answer 1

如果发生CUDA错误会破坏CUDA上下文,则行为未定义.

这种类型的错误是显而易见的,因为它是"粘性的",这意味着一旦发生,每个CUDA API调用都将返回该错误,直到上下文被销毁.

非粘性错误在cuda API调用返回后自动清除(除了cudaPeekAtLastError).任何"崩溃的内核"类型错误(无效访问,未指定的启动失败等)都将是一个棘手的错误.在您的示例中,步骤3将(始终)对cudaMemcpy从设备到主机的传递variableA 的调用结果返回API错误,因此cudaMemcpy操作的结果是未定义且不可靠的 - 就好像cudaMemcpy某些操作在某些操作中也失败了未指明的方式.

由于损坏的CUDA上下文的行为未定义,因此没有任何分配内容的定义,或者通常是在发生此类错误之后的机器状态.

非粘性错误的示例可能是尝试获得cudaMalloc比设备内存中可用数据更多的数据.此类操作将返回内存不足错误,但该错误将在返回后被清除,并且后续(有效)cuda API调用可以成功完成,而不会返回错误.非粘性错误不会破坏CUDA上下文,并且cuda上下文的行为与从未请求过无效操作的行为完全相同.

在许多记录的错误代码描述中都会调出粘性和非粘性错误之间的区别,例如:

非粘性,非cuda-context-corrupting:

cudaErrorMemoryAllocation = 2 API调用失败,因为它无法分配足够的内存来执行请求的操作.

粘性,cuda-context-corrupting:

cudaErrorMisalignedAddress = 74设备在未对齐的内存地址上遇到加载或存储指令.不能使用上下文,因此必须销毁它(并且应该创建一个新的上下文).来自此上下文的所有现有设备内存分配均无效,如果程序要继续使用CUDA,则必须重建.