 
 	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];
+	for (int loadIdx = threadIdx.x; loadIdx < 15; 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)
 		tmp5 -= tmp4 * a_elements[0];

		row_to_start--;

		dB[g_k * row_to_start + idx] = tmp11;
		for (int loadIdx = threadIdx.x; loadIdx < 8; loadIdx += step)
			a_elements[loadIdx] = dA[offset - column_width*(last_row - 1 - row_to_start) - (loadIdx + 1) * factor];
		__syncthreads();
		tmp12 -= tmp11 * a_elements[0];
		tmp13 -= tmp11 * a_elements[1];

		row_to_start--;

		dB[g_k * row_to_start + idx] = tmp5;
		for (int loadIdx = threadIdx.x; loadIdx < 14; 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();
		tmp6 -= tmp5 * a_elements[0];
		tmp7 -= tmp5 * a_elements[1];
