 		row_to_start--;
 
 		dB[g_k * row_to_start + idx] = tmp10;
-		for (int loadIdx = threadIdx.x; loadIdx < 9; loadIdx += step)
-			a_elements[loadIdx] = dA[offset - column_width*(last_row - 1 - row_to_start) - (loadIdx + 1) * factor];
+		for (int loadIdx = threadIdx.x; loadIdx < 9; loadIdx += step) {
+			if (row_to_start - (loadIdx + 1) >= first_row)
+				a_elements[loadIdx] = dA[offset - column_width*(last_row - 1 - row_to_start) - (loadIdx + 1) * factor];
+			else
+				a_elements[loadIdx] = (T)0;
+		}
 		__syncthreads();
 		tmp11 -= tmp10 * a_elements[0];
 		tmp12 -= tmp10 * a_elements[1];

		}

		row_to_start = last_row - 1 - 20 * i;
		for (int loadIdx = threadIdx.x; loadIdx < 19; loadIdx += step)
			a_elements[loadIdx] = dA[offset - column_width*(last_row - 1 - row_to_start) - (loadIdx + 1) * factor];
		__syncthreads();
		dB[g_k * row_to_start + idx] = tmp0;
		tmp1 -= tmp0 * a_elements[0];

		row_to_start--;

		dB[g_k * row_to_start + idx] = tmp9;
		for (int loadIdx = threadIdx.x; loadIdx < 10; loadIdx += step) {
			if (row_to_start - (loadIdx + 1) >= first_row)
				a_elements[loadIdx] = dA[offset - column_width*(last_row - 1 - row_to_start) - (loadIdx + 1) * factor];
			else
				a_elements[loadIdx] = (T)0;
		}
		__syncthreads();
		tmp10 -= tmp9 * a_elements[0];
		tmp11 -= tmp9 * a_elements[1];
