 }
 
 __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
-	
-	block_result[threadIdx.x] = 0;
-	__syncthreads();
+	extern __shared__ float block_result[]; // best value: blockDim.x + blockDim.y
 	
 	int i = blockIdx.x * blockDim.x + threadIdx.x;
 	int j = blockIdx.y * blockDim.y + threadIdx.y;
 	
+	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) * d_v[j]);
+	atomicAdd(&block_result[threadIdx.x], d_m.get(i, j) * block_result[blockDim.x + threadIdx.y]);
 	__syncthreads();
 	
 	if (threadIdx.y == 0)

	
	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_shared_kernel<<<num_blocks, num_threads, CHECKBOARD_BLOCK_MAX_HEIGHT>>>(d_m, d_v, d_r);
	
	Vector h_r(d_r, CPU);
	return h_r;

	std::cout << "Measured time for <GPU naive>: " << ((float) t_chrono)/CLOCKS_PER_SEC << "s" << std::endl;
	
	t_chrono = clock();
	for (int i(0) ; i != NUM_TESTS * 100 ; i++) {
		productMatrixVectorGPU_shared(h_m, h_v);
	}
	t_chrono = clock() - t_chrono;
