Memoria en CUDA

Problema 1

Considera el cálculo suma de matrices en el que cada elemento de la matriz resultado es la suma de los elementos correspondientes de las matrices de entrada. ¿Es posible utilizar memoria compartida para reducir el acceso a memoria global?. Pista: Analiza los elementos accedidos por cada hilo. ¿Existe algún patrón de acceso entre los hilos?

Problema 2

Considera el cálculo multiplicación de matrices. ¿Es posible utilizar memoria compartida para este caso, reduciendo el acceso a memoria global?

Problema 3

¿Qué tipo de comportamiento incorrecto puede ocurrir si olvidamos utilizar la instrucción __syncthreads() en el kernel que se muestra a continuación?. Existen dos llamadas a la función __syncthreads() justifica cada una de ellas.

__global__
void matrix_mul_kernel(float* Md, float* Nd, float* Pd, const int cWidth)
{
  __shared__ float Mds[TILE_WIDTH][TILE_WIDTH];
  __shared__ float Nds[TILE_WIDTH][TILE_WIDTH];

  int bx_ = blockIdx.x;
  int by_ = blockIdx.y;
  int tx_ = threadIdx.x;
  int ty_ = threadIdx.y;

  int row_ = by_ * TILE_WIDTH + ty;
  int col_ = bx_ * TILE_WIDTH + tx;

  float p_value_ = 0.0f;

  for (int m = 0; m < cWidth / TILE_WIDTH; ++m)
  {
    Mds[ty_][tx_] = Md[row_ * cWidth + (m * TILE_WIDTH + tx_)];
    Nds[ty_][tx_] = Nd[(m * TILE_WIDTH + ty_) * cWidth + col_];
    __syncthreads();

    for (int k = 0; k < TILE_WIDTH; ++k)
      p_value_ += Mds[ty_][k] * Nds[k][tx_];

    __syncthreads();
  }

  Pd[row_ * cWidth + col_] = p_value_;
}