ENH: Forward substitution works in 3D Trans Thomas kernel

preconditioner.h

@@ -865,13 +865,13 @@ __global__ void thomas_kernel3D_XT(int const m, FT const alpha, FT const alpha_2
     }
 
 
-//    FT x_reg                 = sh_x[TILE_SIZE-1][tid_l];
-//    sh_x[TILE_SIZE-1][tid_l] = x_reg;
-//
-//    for (int i = TILE_SIZE-2; i >= 0; --i) {
-//            x_reg = sh_x[i][tid_l] - dev_c_prime[m - TILE_SIZE + i] * x_reg;
-//            sh_x[i][tid_l] = x_reg;
-//    }
+    FT x_reg                 = sh_x[TILE_SIZE-1][tid_l];
+    sh_x[TILE_SIZE-1][tid_l] = x_reg;
+
+    for (int i = TILE_SIZE-2; i >= 0; --i) {
+            x_reg = sh_x[i][tid_l] - dev_c_prime[m - TILE_SIZE + i] * x_reg;
+            sh_x[i][tid_l] = x_reg;
+    }
 
 
 //    for (int tile = TILES - 1; tile > 0; tile--) {
@@ -916,7 +916,7 @@ __global__ void thomas_kernel3D_XT(int const m, FT const alpha, FT const alpha_2
 
 	return;
 
-	FT beta = alpha_23;
+//	FT beta = alpha_23;
 
 	// dev_c_prime - precalculated factors
 	// b - input