CUDA：查明主机缓冲区是否已固定（页面锁定）

Question

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

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

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

问题是：我的函数确定主机缓冲区是 pinned/page-locked（cudaHostAlloc）还是不是（常规 malloc）的可行方法是什么？

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

我试过三种方法都失败了： 1- 始终应用 cudaHostRegister()：如果主机缓冲区已经固定，则这种方式不好 2- 运行 cudaPointerGetAttributes()，如果 return 错误是 cudaSuccess，则缓冲区已经固定，无需执行任何操作； else if cudaErrorInvalidValue, apply cudaHostRegister : 由于某种原因，这种方式导致内核执行 returning 错误 3- 运行 cudaHostGetFlags()，如果 return 不成功，则应用 cudaHostRegister ：与 2-.

相同的行为

在2-和3-的情况下，错误是"invalid argumentn"

请注意，我的代码当前未使用流，而是始终为整个主机缓冲区调用 cudaMemcpy()。如果我不使用上述三种方式中的任何一种，我的代码将运行完成，无论主机缓冲区是否固定。

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

Answer 1

您的方法 2 应该有效（我认为方法 3 也应该有效）。您可能对如何在这种情况下进行正确的 CUDA 错误检查感到困惑。

由于您有一个失败的运行time API 调用，如果您在内核调用后执行类似 cudaGetLastError 的操作，它将显示运行time API 之前在 cudaPointerGetAttributes() 调用中发生的故障。就您而言，这不一定是灾难性的。您要做的是清除那个错误，因为您知道它发生了并且已经正确处理了它。您可以通过额外调用 cudaGetLastError 来做到这一点（对于这种类型的 "non-sticky" API 错误，即 API 错误并不意味着 CUDA 上下文已损坏）。

这是一个完整的示例：

$ 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 中的代码时（即在函数调用之前不使用固定内存），那么您将在下一个错误检查步骤中出现 "spurious" 错误（在我的情况下是下一个 API 调用，但在您的情况下可能是在内核调用之后）。

CUDA：查明主机缓冲区是否已固定（页面锁定）

CUDA: Find out if host buffer is pinned (page-locked)

c++

memory

cuda

gpu