умножение куда

Question

Серийный код выглядит следующим образом:

int i, j;
for(j=0; j<ny; j++)
{
    for(i=0; i<nx; i++)
    {
        x[i + j*nx] *= y[i];
    }
}

Я преобразовал это в CUDA, используя это ядро:

int tid = blockIdx.x * blockDim.x + threadIdx.x;
int i,j;
for(tid = 0; tid <nx*ny; tid++)
{
    j = tid/nx;
    i = tid - j*nx;
    x[tid] *= y[i];
}

Однако ядро ​​графического процессора не дает улучшения ускорения? Любые предложения по лучшему решению? Заранее спасибо

Answer 1

Если это серийный код:

  int i, j;
  for(j=0; j<ny; j++)
  {
      for(i=0; i<nx; i++)
      {
          x[i + j*nx] *= y[i];
      }
  }

, то вы должны делать это:

  __global__ void fn(float *x, int nx)
  {
     int tid = blockIdx.x * blockDim.x + threadIdx.x;
     int j = tid/nx, i = tid - j * nx;
     x[tid] *= y[i];
  }

  fn<<<nx*ny/B, B>>>(x, nx); // with B = 256, 512, etc.

То, что вы делаете, довольно странно: вы инструктируете каждый потокядра CUDA для итерации по всем значениям tid от 0 до nx * ny и вычислению той же функции, что и версия вашего процессора!Более того, вместо того, чтобы просто перебирать индексы, вы фактически делаете цикл менее эффективнее, чем для версии CPU;другими словами, вы делаете то же самое в каждом потоке, но менее эффективно, чем в 1 потоке на процессоре.Неудивительно, что это медленнее;это должно быть намного, намного медленнее.Ваше ядро ​​CUDA:

  int **tid** = blockIdx.x * blockDim.x + threadIdx.x;
  int i,j;
  for(**tid** = 0; **tid** <nx*ny; **tid**++)
  {
      j = tid/nx;
      i = tid - j*nx;
      x[tid] *= y[i];
  }

Это делает nx * ny итераций, так же, как ваш код хоста, для каждого потока;вы теряете все преимущества параллелизма, поскольку каждый поток делает одно и то же;вы получите ту же производительность при использовании одного потока в графическом процессоре и тот же результат!

Если это дословный код из вашего исходного файла CUDA, вам нужно изменить его и повторить сравнение;если это код, который вы написали, чтобы помочь объяснить, что ваш код делает для непрофессиональной аудитории, не относящейся к CUDA, то вам нужно представить свой фактический код CUDA, чтобы мы могли видеть, что происходит ... как есть, анализ производительностиЯ сделал - тривиальный - это все, что вы можете ожидать.

Answer 2

учитывая ваш комментарий к этому ответу :

the nx * ny = 2205;поэтому я использовал нет.of blocks = (nx * ny + (threads-1)) / threads и threads = 64.

подразумевает, что вы собираетесь запускать один поток на каждое вычисление, правильная реализация CUDA будет выглядеть так:

int tid = blockIdx.x * blockDim.x + threadIdx.x;
int j = tid/nx;
int i = tid - j*nx;

if (tid < (nx*ny))
    x[tid] *= y[i];

Если вы намеревались для каждого потока вычислить более одного вычисления за запуск ядра, то вы бы измеряли размер сетки, чтобы "заполнить" каждый из SM на целевом GPU, а не использовали бы одинаковое количествопотоками в качестве входного размера, а затем выполните что-то вроде:

int tid = blockIdx.x * blockDim.x + threadIdx.x;
int gsize = blockDim.x * gridDim.x;
int i,j;

for(; tid <nx*ny; tid+=gsize)
{
    j = tid/nx;
    i = tid - j*nx;
    x[tid] *= y[i];
}

Это позволит вам по крайней мере объединить чтение и запись в x и удалить огромное количество избыточных вычислений в опубликованной версии.Можно провести ряд дальнейших оптимизаций, но для этого потребуется больше информации о проблеме, чем было предоставлено в вопросе и последующих комментариях.Ваша схема индексации содержит целочисленное деление, а затем целочисленное умножение-сложение на расчёт.Это много накладных расходов для одного FLOP на входное значение.Тем не менее, учитывая все вышесказанное, если размер проблемы, который я привел, является фактическим размером проблемы, который вас интересует, графический процессор никогда не будет быстрее, чем даже скромный хост-процессор.Чтобы реализовать полезное ускорение с помощью графического процессора для операций с такой низкой арифметической интенсивностью, потребовались бы задачи на много порядков большего размера.

Answer 3

Попробуйте, используя общую память. Одна из лучших реализаций вокруг:

// Matrices are stored in row-major order:
// M(row, col) = *(M.elements + row * M.stride + col)
typedef struct {
   int width;
   int height;
   int stride; // In number of elements
   float *elements;
} Matrix;

// Thread block size
#define BLOCK_SIZE 16

// Get a matrix element
__device__ float GetElement(const Matrix A, int row, int col)
{
   return A.elements[row * A.stride + col];
}

// Set a matrix element
__device__ void SetElement(Matrix A, int row, int col, float value)
{
   A.elements[row * A.stride + col] = value;
}
// Get the BLOCK_SIZExBLOCK_SIZE sub-matrix Asub of A that is
// located col sub-matrices to the right and row sub-matrices down
// from the upper-left corner of A
__device__ Matrix GetSubMatrix(Matrix A, int row, int col)
{
   Matrix Asub;
   Asub.width = BLOCK_SIZE; Asub.height = BLOCK_SIZE;
   Asub.stride = A.stride;
   Asub.elements = &A.elements[A.stride * BLOCK_SIZE * row + 
                               BLOCK_SIZE * col];
   return Asub;
}

// Forward declaration of the matrix multiplication kernel
__global__ void MatMulKernel(const Matrix, const Matrix, Matrix);

// Matrix multiplication - Host code
// Matrix dimensions are assumed to be multiples of BLOCK_SIZE
void MatMul(const Matrix A, const Matrix B, Matrix C)
{
   // Same as in previous example, except the followings:
   // d_A.width = d_A.stride = A.width;
   // d_B.width = d_B.stride = B.width;
   // d_C.width = d_C.stride = C.width;
}
// Matrix multiplication kernel called by MatMul()
__global__ void MatMulKernel(Matrix A, Matrix B, Matrix C)
{
   // Block row and column
   int blockRow = blockIdx.y;
   int blockCol = blockIdx.x;

   // Each thread block computes one sub-matrix Csub of C
   Matrix Csub = GetSubMatrix(C, blockRow, blockCol);

   // Each thread computes one element of Csub
   // by accumulating results into Cvalue
   float Cvalue = 0;

   // Thread row and column within Csub
   int row = threadIdx.y;
   int col = threadIdx.x;
// Loop over all the sub-matrices of A and B that are
   // required to compute Csub
   // Multiply each pair of sub-matrices together
   // and accumulate the results
   for (int m = 0; m < (A.width / BLOCK_SIZE); ++m) 
   {
      // Get sub-matrix Asub of A and Bsub of B
      Matrix Asub = GetSubMatrix(A, blockRow, m);
      Matrix Bsub = GetSubMatrix(B, m, blockCol);

      // Shared memory used to store Asub and Bsub respectively
      __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
      __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];

      // Load Asub and Bsub from device memory to shared memory
      // Each thread loads one element of each sub-matrix
      As[row][col] = GetElement(Asub, row, col);
      Bs[row][col] = GetElement(Bsub, row, col);

      // Synchronize to make sure the sub-matrices are loaded
      // before starting the computation
      __syncthreads();
      // Multiply Asub and Bsub together
      for (int e = 0; e < BLOCK_SIZE; ++e)
         Cvalue += As[row][e] * Bs[e][col];

      // Synchronize to make sure that the preceding
      // computation is done before loading two new
      // sub-matrices of A and B in the next iteration
      __syncthreads();
   }

   // Write Csub to device memory
   // Each thread writes one element
   SetElement(Csub, row, col, Cvalue);
}

Answer 4

Насколько большой блок? может случиться так, что время, необходимое для копирования небольшого объема данных в графический процессор и настройки среды, намного больше, чем время расчета.

Помните также, что CUDA выполняет jit-компиляцию при первом запуске, поэтому для точного бенчмаркинга вам нужно запускать его много раз.