我的问题的简短描述如下:
我开发了一个调用 CUDA 内核的函数。我的函数接收到一个指向主机数据缓冲区(内核的输入和输出)的指针,并且无法控制这些缓冲区的分配。
--> 主机数据可能是使用 malloc 或 cudaHostAlloc 分配的。我的函数没有具体告诉我使用了哪种分配方法。
问题是:我的函数确定主机缓冲区是固定/页面锁定(cudaHostAlloc)还是非(常规 malloc)的可行方法是什么?
我问的原因是,如果它们不是页面锁定的,我想使用 cudaHostRegister() 使它们(缓冲区)如此,使它们适合流。
我试过三种方法都失败了: 1- 始终应用 cudaHostRegister():如果主机缓冲区已经固定,则这种方式不好 2- 运行 cudaPointerGetAttributes(),如果返回错误为 cudaSuccess,则缓冲区已被固定,无需执行任何操作; else if cudaErrorInvalidValue, apply cudaHostRegister :出于某种原因,这种方式导致内核执行返回错误 3- 运行 cudaHostGetFlags(),如果返回不成功,则应用 cudaHostRegister:与 2- 相同的行为。
在2-和3-的情况下,错误是“invalid argumentn”
请注意,我的代码当前未使用流,而是始终为整个主机缓冲区调用 cudaMemcpy()。如果我不使用上述三种方式中的任何一种,无论主机缓冲区是否固定,我的代码都会运行完成。
有什么建议吗?非常感谢。
最佳答案
您的方法 2 应该有效(我认为方法 3 也应该有效)。您可能对如何在这种情况下进行正确的 CUDA 错误检查感到困惑。
由于您有一个失败的运行时 API 调用,如果您在内核调用之后执行类似 cudaGetLastError
的操作,它将显示 之前 发生的运行时 API 失败cudaPointerGetAttributes()
调用。就您而言,这不一定是灾难性的。你想要做的是清除那个错误,因为你知道它发生了并且已经正确地处理了它。您可以通过额外调用 cudaGetLastError
来做到这一点(对于这种类型的“非粘性”API 错误,即不表示 CUDA 上下文已损坏的 API 错误)。
这是一个完整的例子:
$ cat t642.cu
#include <stdio.h>
#include <stdlib.h>
#define DSIZE 10
#define nTPB 256
#define cudaCheckErrors(msg) \
do { \
cudaError_t __err = cudaGetLastError(); \
if (__err != cudaSuccess) { \
fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
msg, cudaGetErrorString(__err), \
__FILE__, __LINE__); \
fprintf(stderr, "*** FAILED - ABORTING\n"); \
exit(1); \
} \
} while (0)
__global__ void mykernel(int *data, int n){
int idx = threadIdx.x+blockDim.x*blockIdx.x;
if (idx < n) data[idx] = idx;
}
int my_func(int *data, int n){
cudaPointerAttributes my_attr;
if (cudaPointerGetAttributes(&my_attr, data) == cudaErrorInvalidValue) {
cudaGetLastError(); // clear out the previous API error
cudaHostRegister(data, n*sizeof(int), cudaHostRegisterPortable);
cudaCheckErrors("cudaHostRegister fail");
}
int *d_data;
cudaMalloc(&d_data, n*sizeof(int));
cudaCheckErrors("cudaMalloc fail");
cudaMemset(d_data, 0, n*sizeof(int));
cudaCheckErrors("cudaMemset fail");
mykernel<<<(n+nTPB-1)/nTPB, nTPB>>>(d_data, n);
cudaDeviceSynchronize();
cudaCheckErrors("kernel fail");
cudaMemcpy(data, d_data, n*sizeof(int), cudaMemcpyDeviceToHost);
cudaCheckErrors("cudaMemcpy fail");
int result = 1;
for (int i = 0; i < n; i++) if (data[i] != i) result = 0;
return result;
}
int main(int argc, char *argv[]){
int *h_data;
int mysize = DSIZE*sizeof(int);
int use_pinned = 0;
if (argc > 1) if (atoi(argv[1]) == 1) use_pinned = 1;
if (!use_pinned) h_data = (int *)malloc(mysize);
else {
cudaHostAlloc(&h_data, mysize, cudaHostAllocDefault);
cudaCheckErrors("cudaHostAlloc fail");}
if (!my_func(h_data, DSIZE)) {printf("fail!\n"); return 1;}
printf("success!\n");
return 0;
}
$ nvcc -o t642 t642.cu
$ ./t642
success!
$ ./t642 1
success!
$
在您的情况下,我相信您没有像我在放置评论的那一行所做的那样正确处理 API 错误:
// clear out the previous API error
如果您省略此步骤(您可以尝试将其注释掉),那么当您在案例 0 中运行代码时(即在函数调用之前不使用固定内存),您将获得一个“下一个错误检查步骤(在我的情况下是下一个 API 调用,但在您的情况下可能是在内核调用之后)出现虚假”错误。
关于c++ - CUDA:查明主机缓冲区是否已固定(页面锁定),我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/28861161/