Блог

Главная - Вопросы по программированию - Параллельные пакетные малые матрицы в CUDA не работают с циклом for

Параллельные пакетные малые матрицы в CUDA не работают с циклом for

#c #cuda #gpu

#c #cuda #графический процессор

Вопрос:

У меня есть число (скажем, миллион) маленьких матриц 4 x 3. Я хотел бы выполнить с ними несколько простых операций, и я хотел бы, чтобы мое ядро CUDA распараллеливало только индекс матриц (а не операции со строками / столбцами). Позвольте мне объяснить лучше: я передаю в качестве входных данных в ядро моего графического процессора массив матриц A [MatrixNumb][строка] [col], и я хотел бы, чтобы распараллеливание операций было только для MatrixNumb (поэтому я хочу принудительно выполнить операцию в одном измерении. Для простоты приведенный ниже пример содержит только 3 матрицы. Он компилируется и запускается, однако он дает мне неправильные результаты. По сути, он возвращает те же матрицы, которые я даю ему в качестве входных данных. Я не могу понять, почему, и если я совершаю какую-либо ошибку, как я могу переписать / продумать код? Я написал код, используя также cudaMallocManaged, чтобы иметь общую память между хостом и устройством, однако он дает мне те же результаты, используя классический cudaMalloc и используя memcpy.

Source.cpp

 #include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <iostream>
#include <assert.h>
#include <chrono>
#include <random>
#include <time.h>
#include <math.h>

#include <cuda_runtime.h>
#include "device_launch_parameters.h"
#include <cuda.h>
#include <device_functions.h>

using namespace std;


__global__ void SVD(double*** a, const int m, const int n, const int numMatrices, double** w)
{
  int idx = blockIdx.x * blockDim.x   threadIdx.x;

  // I would like that each thread runs these loops independently
  for (int i = 0; i < m; i  ) {
    for (int j = 0; j < n; j  ) {
      a[idx][i][j] = (a[idx][i][j] * a[idx][i][j]) * 3.14;
    }
  }
  for (int j = 0; j < n; j  ) {
    w[idx][j] = 3.14 * a[idx][1][j]* a[idx][1][j];
  }

}


int main()
{
  const int n = 3;
  const int m = 4;
  const int lda = m;
  const int numMatrices = 3;

  random_device device;
  mt19937 generator(device());
  uniform_real_distribution<double> distribution(1., 5.);

  // create pointers
  double*** A = new double** [numMatrices];
  double** w = new double* [numMatrices];

  //ALLOCATE SHARED MEMORY
  for (int nm = 0; nm < numMatrices; nm  ) {
    A[nm] = new double* [lda];
    w[nm] = new double[n];

    for (int i = 0; i < lda; i  ) {
      A[nm][i] = new double[n];

      for (int j = 0; j < n; j  ) {
        cudaMallocManaged((void**)amp;A[nm][i][j], sizeof(double));
        cudaMallocManaged((void**)amp;w[nm][j], sizeof(double));
      }
    }
  }

  cout << " memory allocated" << endl;


  //FILL MATRICES INTO SHARED MEMORY
  for (int nm = 0; nm < numMatrices; nm  ) {
    A[nm] = new double* [lda];
    w[nm] = new double[n];                                   

    for (int i = 0; i < lda; i  ) {
      A[nm][i] = new double[n];

      for (int j = 0; j < n; j  ) {
        A[nm][i][j] = distribution(generator);
        w[nm][j] = 0.0;
      }
    }
  }
  cout << " matrix filled " << endl;


  // PRINT MATRICES BEFORE CUDA OPERATION
  for (int nm = 0; nm < numMatrices; nm  ) {
    for (int i = 0; i < lda; i  ) {
      for (int j = 0; j < n; j  ) {
        cout << A[nm][i][j] << " ";
      }
      cout << endl;
    }
    cout << endl;
  }

  //KERNEL ----------------------------------------------------------------------
  int NThreads = 3;   
  int NBlocks = int(numMatrices / NThreads   1);
 
  SVD << <NBlocks, NThreads >> > (A, n, m, numMatrices, w);
  cudaDeviceSynchronize();
  cout << " Kernel done " << endl << endl;

  cout << " --- GPU --- " << endl;
  cout << " NEW MATRIX: " << endl;
  for (int nm = 0; nm < numMatrices; nm  ) {
    for (int i = 0; i < lda; i  ) {
      for (int j = 0; j < n; j  ) {
        cout << A[nm][i][j] << " ";
      }
      cout << endl;
    }
    cout << endl;
  }

  cout << " NEW VECTOR RESULTS: " << endl;
  for (int nm = 0; nm < numMatrices; nm  ) {
    for (int i = 0; i < n; i  ) {
      cout << w[nm][i] << " ";
    }
    cout << endl;
  }

  cout << endl;

  //FREE THE DEVICE'S MEMORY -----------------------------------------------------
  cudaFree(A);
  cudaFree(w);
  cout << " Cuda free " << endl << endl;

  return 0;
}

(Неправильный) результат, который я получаю, выглядит следующим образом:

 memory allocated
 matrix filled
1.28689 3.76588 3.88649
1.52547 4.42371 2.62566
1.48002 3.33719 1.58413
3.78243 2.8394 3.0249

1.14322 1.70261 2.02784
2.86852 2.87918 3.2896
4.87268 3.52447 1.58414
3.52306 3.84931 3.18212

1.76397 1.41317 4.9765
1.63338 4.79316 2.64009
1.99873 1.72617 1.15974
1.18922 4.21513 1.6695

 Kernel done

 --- GPU ---
 NEW MATRIX:
1.28689 3.76588 3.88649
1.52547 4.42371 2.62566
1.48002 3.33719 1.58413
3.78243 2.8394 3.0249

1.14322 1.70261 2.02784
2.86852 2.87918 3.2896
4.87268 3.52447 1.58414
3.52306 3.84931 3.18212

1.76397 1.41317 4.9765
1.63338 4.79316 2.64009
1.99873 1.72617 1.15974
1.18922 4.21513 1.6695

 NEW VECTOR RESULTS:
0 0 0
0 0 0
0 0 0

 Cuda free

Я ожидал увидеть новые матрицы и векторы, измененные операциями:
a[idx] [i] [j] = (a [idx][i] [j] * a [idx] [i] [j]) * 3.14;
однако, похоже, что код не видит ядро или ядро не работает должным образом.

Комментарии:

1. каждый указатель, который может быть разыменован в коде устройства, должен использовать управляемый распределитель. Поэтому, поскольку вы передаете, например, A код устройства, они A должны быть выделены с помощью управляемого распределителя . Это не может работать: double*** A = new double** [numMatrices]; вы должны сделать это: double*** A; cudaMallocManaged(amp;A, sizeof(double**)*numMatrices); и вы должны следовать этому с каждым другим указателем в дереве для каждого указателя, который вы передаете ядру.

Ответ №1:

У вас было несколько проблем:

  1. При использовании управляемой памяти с двойным или тройным доступом к указателю каждый указатель в дереве должен быть выделен с помощью управляемого распределителя
  2. В ваших схемах распределения было слишком много уровней, и вы дважды выделяли некоторые указатели (утечка памяти).
  3. Порядок аргументов, которые вы передаете своему ядру, не соответствует порядку аргументов, ожидаемых вашим ядром ( n , m были обратными).
  4. Поскольку вы потенциально запускаете больше блоков / потоков, чем необходимо, вашему ядру требуется проверка потока (if-test).
  5. Ваш код должен быть в .cu файле, а не в .cpp файле.

В следующем коде устранены вышеуказанные проблемы, и, похоже, он выполняется без ошибок во время выполнения.

 $ cat t61.cu
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <iostream>
#include <assert.h>
#include <chrono>
#include <random>
#include <time.h>
#include <math.h>


using namespace std;


__global__ void SVD(double*** a, const int m, const int n, const int numMatrices, double** w)
{
  int idx = blockIdx.x * blockDim.x   threadIdx.x;
  if (idx < numMatrices){
  // I would like that each thread runs these loops independently
  for (int i = 0; i < m; i  ) {
    for (int j = 0; j < n; j  ) {
      a[idx][i][j] = (a[idx][i][j] * a[idx][i][j]) * 3.14;
    }
  }
  for (int j = 0; j < n; j  ) {
    w[idx][j] = 3.14 * a[idx][1][j]* a[idx][1][j];
  }
  }
}


int main()
{
  const int n = 3;
  const int m = 4;
  const int lda = m;
  const int numMatrices = 3;

  random_device device;
  mt19937 generator(device());
  uniform_real_distribution<double> distribution(1., 5.);

  // create pointers
  double*** A;
  cudaMallocManaged(amp;A, sizeof(double**)*numMatrices);
  double** w;
  cudaMallocManaged(amp;w, sizeof(double*)* numMatrices);

  //ALLOCATE SHARED MEMORY
  for (int nm = 0; nm < numMatrices; nm  ) {
    cudaMallocManaged(amp;(A[nm]), sizeof(double*)*lda);
    cudaMallocManaged(amp;(w[nm]), sizeof(double)*n);

    for (int i = 0; i < lda; i  ) {
      cudaMallocManaged(amp;(A[nm][i]), sizeof(double)*n);
      }
    }

  cout << " memory allocated" << endl;


  //FILL MATRICES INTO SHARED MEMORY
  for (int nm = 0; nm < numMatrices; nm  ) {
    for (int i = 0; i < lda; i  ) {
      for (int j = 0; j < n; j  ) {
        A[nm][i][j] = distribution(generator);
        w[nm][j] = 0.0;
      }
    }
  }
  cout << " matrix filled " << endl;


  // PRINT MATRICES BEFORE CUDA OPERATION
  for (int nm = 0; nm < numMatrices; nm  ) {
    for (int i = 0; i < lda; i  ) {
      for (int j = 0; j < n; j  ) {
        cout << A[nm][i][j] << " ";
      }
      cout << endl;
    }
    cout << endl;
  }

  //KERNEL ----------------------------------------------------------------------
  int NThreads = 3;
  int NBlocks = int(numMatrices / NThreads   1);

  SVD << <NBlocks, NThreads >> > (A, m, n, numMatrices, w);
  cudaDeviceSynchronize();
  cout << " Kernel done " << endl << endl;

  cout << " --- GPU --- " << endl;
  cout << " NEW MATRIX: " << endl;
  for (int nm = 0; nm < numMatrices; nm  ) {
    for (int i = 0; i < lda; i  ) {
      for (int j = 0; j < n; j  ) {
        cout << A[nm][i][j] << " ";
      }
      cout << endl;
    }
    cout << endl;
  }

  cout << " NEW VECTOR RESULTS: " << endl;
  for (int nm = 0; nm < numMatrices; nm  ) {
    for (int i = 0; i < n; i  ) {
      cout << w[nm][i] << " ";
    }
    cout << endl;
  }

  cout << endl;

  //FREE THE DEVICE'S MEMORY -----------------------------------------------------
  cudaFree(A);
  cudaFree(w);
  cout << " Cuda free " << endl << endl;

  return 0;
}
$ nvcc -o t61 t61.cu
$ cuda-memcheck ./t61
========= CUDA-MEMCHECK
 memory allocated
 matrix filled
3.73406 3.51919 3.249
1.52374 2.678 2.50944
3.67358 1.15831 3.26327
2.58468 1.49937 2.67133

1.72144 2.99183 3.11156
1.06247 3.34983 4.23568
3.49749 3.07641 3.42827
4.09607 2.00557 2.12049

3.65427 3.98966 4.73428
1.68397 4.3746 2.95533
2.1914 4.96086 1.7165
3.10095 2.61781 4.52626

 Kernel done

 --- GPU ---
 NEW MATRIX:
43.7816 38.888 33.1458
7.29041 22.5191 19.7735
42.375 4.2129 33.4376
20.977 7.05908 22.407

9.30494 28.1062 30.4008
3.54453 35.2351 56.3348
38.41 29.7179 36.9045
52.6821 12.6301 14.1189

41.9306 49.9807 70.3782
8.90432 60.0905 27.4247
15.079 77.2757 9.25165
30.1939 21.5182 64.3294

 NEW VECTOR RESULTS:
166.891 1592.32 1227.71
39.4501 3898.35 9965.14
248.961 11338.1 2361.64

 Cuda free

========= ERROR SUMMARY: 0 errors
$