CUDA: выяснить, закреплен ли буфер хоста (заблокирован ли он страницей)

Question

Краткое описание моей проблемы таково:

Я разработал функцию, которая вызывает ядро ​​CUDA. Моя функция получает указатель на буферы данных хоста (вход и выход ядра) и не контролирует распределение этих буферов.

-> Возможно, данные хоста были выделены с помощью malloc или cudaHostAlloc. В моей функции конкретно не указано, какой метод размещения использовался.

Вопрос заключается в следующем: каков способ для моей функции выяснить, закреплены ли буферы хоста / заблокированы ли страницы (cudaHostAlloc) или нет (обычный malloc)?

Причина, по которой я спрашиваю, состоит в том, что, если они не заблокированы на странице, я хотел бы использовать cudaHostRegister (), чтобы сделать их (буферы) так, чтобы сделать их пригодными для потоков.

Я пробовал три способа, которые потерпели неудачу: 1- Всегда применять cudaHostRegister (): этот способ не годится, если буферы хоста уже закреплены. 2- Запустите cudaPointerGetAttributes(), и если ошибка возврата - cudaSuccess, то буферы уже закреплены., нечего делать; иначе, если cudaErrorInvalidValue, примените cudaHostRegister: по какой-то причине этот путь приводит к тому, что выполнение ядра возвращает ошибку 3- Запустите cudaHostGetFlags(), а если return не удастся, тогда примените cudaHostRegister: то же поведение, что и 2-.

В случае 2- и 3- ошибкой является "неверный аргумент n"

Обратите внимание, что мой код в настоящее время не использует потоки, скорее, всегда вызывает cudaMemcpy() для всех буферов хоста. Если я не использую ни один из трех вышеперечисленных способов, мой код выполняется до завершения независимо от того, закреплен ли буфер хоста или нет.

Любой совет? Спасибо заранее.

Answer 1

Ваш метод 2 должен работать (я думаю, что метод 3 должен работать также). Вероятно, вас смущает, как правильно выполнить проверку ошибок CUDA в этом сценарии.

Поскольку у вас есть вызов API времени выполнения, который не работает, если вы делаете что-то вроде cudaGetLastError после вызова ядра, он покажет сбой API времени выполнения, который произошел ранее на cudaPointerGetAttributes() вызов. Это не обязательно катастрофично, в вашем случае. Что вы хотите сделать, это очистить эту ошибку, поскольку вы знаете, что она произошла, и правильно ее обработали. Вы можете сделать это с дополнительным вызовом cudaGetLastError (для этого типа "нелипкой" ошибки 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 (т.е. не используете закрепленную память до вызова функции), то вы увидите, что вы получите "ложную" ошибку на следующем шаге проверки ошибок (следующий вызов API в моем случае, но может быть после вызова ядра в вашем случае).