-ta=tesla:managed:cuda8, ​​но cuMemAllocManaged вернул ошибку 2: Недостаточно памяти

Я считаю, что проблема здесь:

size_t n = GB*1024*1024*1024/sizeof(float);

когда я компилирую эту строку кода с помощью g++, я получаю предупреждение о целочисленном переполнении. По какой-то причине компилятор PGI не предупреждает, но под капотом происходит то же самое. После объявлений s и n, если я добавлю такую ​​распечатку:

  size_t n = GB*1024*1024*1024/sizeof(float);
  size_t s = n * sizeof(float);
  printf("n = %lu, s = %lu\n", n, s);  // add this line

и скомпилировать с PGI 17.04 и запустить (на P100, с 16 ГБ), я получаю такой вывод:

$ pgcc -acc -fast -ta=tesla:managed:cuda8 -Minfo m1.c
main:
     16, Loop not fused: function call before adjacent loop
         Generated vector simd code for the loop
     22, Loop not fused: function call before adjacent loop
         Generated vector simd code for the loop containing reductions
         Generated a prefetch instruction for the loop
$ ./a.out
n = 4611686017890516992, s = 18446744071562067968
malloc: call to cuMemAllocManaged returned error 2: Out of memory
Aborted
$

так что очевидно, что n и s не то, что вы хотели.

Мы можем исправить это, пометив все эти константы ULL, и тогда у меня все будет работать правильно:

$ cat m1.c
#include <stdio.h>
#include <openacc.h>
#include <stdlib.h>

#define GB 20ULL

int main()
{
  float *a;
  size_t n = GB*1024ULL*1024ULL*1024ULL/sizeof(float);
  size_t s = n * sizeof(float);
  printf("n = %lu, s = %lu\n", n, s);
  a = (float *)malloc(s);
  if (!a) { printf("Failed to malloc.\n"); return 1; }
  printf("Initializing ... ");
  for (int i = 0; i < n; ++i) {
    a[i] = 0.1f;
  }
  printf("done\n");
  double sum=0.0;
  #pragma acc loop reduction (+:sum)
  for (int i = 0; i < n; ++i) {
    sum+=a[i];
  }
  printf("Sum is %f\n", sum);
  free(a);
  return 0;
}
$ pgcc -acc -fast -ta=tesla:managed:cuda8 -Minfo m1.c
main:
     16, Loop not fused: function call before adjacent loop
         Generated vector simd code for the loop
     22, Loop not fused: function call before adjacent loop
         Generated vector simd code for the loop containing reductions
         Generated a prefetch instruction for the loop
$ ./a.out
n = 5368709120, s = 21474836480
Initializing ... done
Sum is 536870920.000000
$

Обратите внимание, что я сделал еще одно изменение выше. Я изменил переменную накопления sum с float на double. Это необходимо, чтобы сохранить в некоторой степени «разумные» результаты при выполнении очень большого сокращения для очень малых количеств.

И, как указал @MatColgrove в своем ответе, я пропустил еще несколько вещей.

Помимо того, что упомянул Боб, я сделал еще несколько исправлений.

Во-первых, вы на самом деле не создаете вычислительную область OpenACC, поскольку у вас есть только директива «#pragma acc loop». Это должно быть "параллельный цикл #pragma acc". Вы можете увидеть это в сообщениях обратной связи компилятора, где он показывает только оптимизацию кода хоста.

Во-вторых, индекс «i» должен быть объявлен как «длинный». В противном случае вы переполните индекс.

Наконец, вам нужно добавить «cc60» к вашим параметрам целевого ускорителя, чтобы указать компилятору ориентироваться на графический процессор на основе Pascal.

% cat mi.c  
#include <stdio.h>
#include <openacc.h>
#include <stdlib.h>

#define GB 20ULL

int main()
{
  float *a;
  size_t n = GB*1024ULL*1024ULL*1024ULL/sizeof(float);
  size_t s = n * sizeof(float);
  printf("n = %lu, s = %lu\n", n, s);
  a = (float *)malloc(s);
  if (!a) { printf("Failed to malloc.\n"); return 1; }
  printf("Initializing ... ");
  for (int i = 0; i < n; ++i) {
    a[i] = 0.1f;
  }
  printf("done\n");
  double sum=0.0;
  #pragma acc parallel loop reduction (+:sum)
  for (long i = 0; i < n; ++i) {
    sum+=a[i];
  }
  printf("Sum is %f\n", sum);
  free(a);
  return 0;
}

% pgcc -fast -acc -ta=tesla:managed,cuda8.0,cc60 -Minfo=accel mi.c
main:
     21, Accelerator kernel generated
         Generating Tesla code
         21, Generating reduction(+:sum)
         22, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
     21, Generating implicit copyin(a[:5368709120])
% ./a.out
n = 5368709120, s = 21474836480
Initializing ... done
Sum is 536870920.000000