CUDA:查明主机缓冲区是否被固定(页面锁定)

Question

我的问题的简短描述如下：

我开发了一个调用CUDA内核的函数。我的函数接收一个指向主机数据缓冲区（内核的输入和输出）的指针，并且不能控制这些缓冲区的分配。

-->主机数据可能是用malloc或cudaHostAlloc分配的。我的函数没有具体说明使用了哪种分配方法。

问题是：对于我的函数来说，什么是可行的方法来判断主机缓冲区是否被固定/页面锁定（cudaHostAlloc）（常规malloc）？

我问的原因是，如果它们没有页面锁定，我想使用cudaHostRegister（）使它们（缓冲区）如此，使它们适合流。

我尝试了三种失败的方法：1-始终应用cudaHostRegister（）：如果主机缓冲区已经固定，这种方式是不好的2-运行cudaPointerGetAttributes（），如果返回错误为cudaSuccess，则缓冲区已被固定，无需执行任何操作；否则，如果cudaErrorInvalidValue，则应用cudaHostRegister：由于某种原因，这种方式会导致内核执行返回错误3-运行cudaHostGetFlags（），如果返回不成功，则应用cudaHostRegister:与2-相同的行为。

在2-和3-的情况下，错误为"无效论证n"

请注意，我的代码目前不使用流，而是总是为整个主机缓冲区调用cudaMemcpy（）。如果我不使用上面三种方法中的任何一种，那么无论主机缓冲区是否固定，我的代码都会运行到完成。

有什么建议吗？非常感谢。

Answer 1

您的方法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 - ABORTINGn"); 
            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调用，但在您的情况下可能是在内核调用之后），您似乎会得到一个"虚假"错误。