Разбор программы

// Perform "naive" square matrix multiplication
//
#include <stdio.h>
 
#define BLOCK_SIZE  16          // submatrix size
#define N           1024        // matrix size is N*N
 
 
// KERNEL //
__global__ void matMult ( float * a, float * b, int n, float * c ){
  int   bx  = blockIdx.x;       // block index
  int   by  = blockIdx.y;
  int   tx  = threadIdx.x;      // thread index
  int   ty  = threadIdx.y;
  float sum = 0.0f;         // подэлемент
  int     ia  = n * BLOCK_SIZE * by + n * ty;   // a [i][0]
  int   ib  = BLOCK_SIZE * bx + tx;
    
  // умножение матриц
  for ( int k = 0; k < n; k++ )
    sum += a [ia + k] * b [ib + k*n];
            
  // блок глобальной памяти для матрицы
  // каждый поток в один элемент
  int ic = n * BLOCK_SIZE * by + BLOCK_SIZE * bx;
    
  c [ic + n * ty + tx] = sum;
}
 
 
// HOST CODE //
int main ( int argc, char *  argv [] ){
  int   numBytes = N * N * sizeof ( float );
  
  // выделение памяти хоста
  float * a = new float [N*N];
  float * b = new float [N*N];
  float * c = new float [N*N];
 
  for ( int i = 0; i < N; i++ )
    for ( int j = 0; j < N; j++ ){
      a [i] = 0.0f;
      b [i] = 1.0f;
    }
        
  // выделение памяти устройства
  float * adev = NULL;
  float * bdev = NULL;
  float * cdev = NULL;
    
  cudaMalloc ( (void**)&adev, numBytes );
  cudaMalloc ( (void**)&bdev, numBytes );
  cudaMalloc ( (void**)&cdev, numBytes );
 
  // конфигурация запуска ядра
  dim3 threads ( BLOCK_SIZE, BLOCK_SIZE );
  dim3 blocks  ( N / threads.x, N / threads.y);
 
  // создание куда таймера
  cudaEvent_t start, stop;
  float gpuTime = 0.0f;
 
  cudaEventCreate ( &start );
  cudaEventCreate ( &stop );
    
  // asynchronously issue work to the GPU (all to stream 0)
  cudaEventRecord ( start, 0 );
  cudaMemcpy      ( adev, a, numBytes, cudaMemcpyHostToDevice );
  cudaMemcpy      ( bdev, b, numBytes, cudaMemcpyHostToDevice );
  
  matMult<<<blocks, threads>>> ( adev, bdev, N, cdev );
    
  cudaMemcpy      ( c, cdev, numBytes, cudaMemcpyDeviceToHost );
  cudaEventRecord ( stop, 0 );
 
  cudaEventSynchronize ( stop );
  cudaEventElapsedTime ( &gpuTime, start, stop );
 
  
  
 
  // время работы
  printf("time spent executing by the GPU: %.2f millseconds\n", gpuTime );
 
  // освобождение ресурсов
  cudaEventDestroy ( start );
  cudaEventDestroy ( stop  );
  cudaFree         ( adev  );
  cudaFree         ( bdev  );
  cudaFree         ( cdev  );
 
  delete a;
  delete b;
  delete c;
  system("pause");
 
  return 0;
}