 	const dim3 num_threads(CHECKBOARD_BLOCK_MAX_HEIGHT, CHECKBOARD_BLOCK_MAX_WIDTH, 1);
 	const dim3 num_blocks((d_m.getHeight() + CHECKBOARD_BLOCK_MAX_HEIGHT -1)/CHECKBOARD_BLOCK_MAX_HEIGHT, (d_m.getWidth() + CHECKBOARD_BLOCK_MAX_WIDTH -1)/CHECKBOARD_BLOCK_MAX_WIDTH, 1);
 	productMatrixVectorGPU_naive_kernel<<<num_blocks, num_threads>>>(d_m, d_v, d_r);
+	cudaThreadSynchronize(); // block until the device is finished
+	
+	// check for error
+	cudaError_t error = cudaGetLastError();
+  	if(error != cudaSuccess) {
+    		std::cout << "CUDA error: " << cudaGetErrorString(error) << std::endl;
+		throw error;
+	}
 	
 	Vector h_r(d_r, CPU);
 	return h_r;
 }
 
 __global__ void productMatrixVectorGPU_shared_kernel(const Matrix d_m, const Vector d_v, Vector d_r) {
-	extern __shared__ float block_result[]; // best value: blockDim.x + blockDim.y
+	extern __shared__ float block_result[]; // need: blockDim.x + blockDim.y
 	
 	int i = blockIdx.x * blockDim.x + threadIdx.x;
 	int j = blockIdx.y * blockDim.y + threadIdx.y;
 	
+	if (i >= d_m.getHeight() || j >= d_m.getWidth())
+		return;
+	
+	__syncthreads();
+		
 	block_result[threadIdx.x] = 0;
 	block_result[blockDim.x + threadIdx.y] = d_v[j];
 	__syncthreads();
 	
-	if (i >= d_m.getHeight() || j >= d_m.getWidth())
-		return;
-	
 	atomicAdd(&block_result[threadIdx.x], d_m.get(i, j) * block_result[blockDim.x + threadIdx.y]);
 	__syncthreads();
 	

#include <string.h>
#include <iostream>


#include "Matrix.cuh"
#include "Vector.cuh"

#define CPU 1
#define GPU 0

#define MINI_NUM_TESTS 16

#define CHECKBOARD_BLOCK_MAX_WIDTH  1
#define CHECKBOARD_BLOCK_MAX_HEIGHT 256
