 
 	if (row_to_start < first_row) return;
 	dB[g_k * row_to_start + idx] = tmp14;
-	for (int loadIdx = threadIdx.x; loadIdx < 5; loadIdx += step)
-		a_elements[loadIdx] = dA[offset - column_width*(last_row - 1 - row_to_start) - (loadIdx + 1) * factor];
+	for (int loadIdx = threadIdx.x; loadIdx < 5; 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();
 	if (row_to_start - 1 >= first_row)
 		tmp15 -= tmp14 * a_elements[0];


	if (row_to_start < first_row) return;
	dB[g_k * row_to_start + idx] = tmp4;
	for (int loadIdx = threadIdx.x; loadIdx < 15; loadIdx += step)
		a_elements[loadIdx] = dA[offset - column_width*(last_row - 1 - row_to_start) - (loadIdx + 1) * factor];
	__syncthreads();
	if (row_to_start - 1 >= first_row)
		tmp5 -= tmp4 * a_elements[0];


	if (row_to_start < first_row) return;
	dB[g_k * row_to_start + idx] = tmp3;
	for (int loadIdx = threadIdx.x; loadIdx < 16; 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();
	if (row_to_start- 1 >= first_row)
		tmp4 -= tmp3 * a_elements[0];
