 	}
 }
 
-template<int I1, int I2, int I3, int I4>
-__global__ void preact_c(float input[I4][I3+I1-1][I3+I1-1], float preact[I2*I4][I3][I3], float weight[I2][I1][I1])
-{
-	FOR6(I1, I1, I2, I3, I3, I4, atomicAdd(&preact[i6 * I2 + i3][i4][i5], weight[i3][i1][i2] * input[i6][i4 + i1][i5 + i2]));
-}
-
-template<int I1, int I2>
-__global__ void bias_c(float preact[I1][I2][I2], float bias[I1])
-{
-	FOR3(I1, I2, I2, preact[i1][i2][i3] += bias[i1]);
-}
-
-template<int I1, int I2, int I3>
-__global__ void preact_s(float input[I2][I1*I3][I1*I3], float preact[I2][I3][I3], float weight[I1][I1])
-{
-	FOR5(I1, I1, I2, I3, I3, atomicAdd(&preact[i3][i4][i5], weight[i1][i2] * input[i3][i4 * I1 + i1][i5 * I1 + i2]));
-}
-
-template<int I1, int I2>
-__global__ void bias_s(float preact[I1][I2][I2], float bias[1])
-{
-	FOR3(I1, I2, I2, preact[i1][i2][i3] += bias[0]);
-}
-
-__global__ void preact_f(float input[36][4][4], float preact[10], float weight[10][36][4][4])
-{
-	FOR4(10, 36, 4, 4, atomicAdd(&preact[i1], weight[i1][i2][i3][i4] * input[i2][i3][i4]));
+__global__ void preact_c1(float input[28][28], float preact[6][24][24], float weight[6][5][5])
+{
+	const int pos = blockIdx.x * blockDim.x + threadIdx.x;
+	const int size = blockDim.x * gridDim.x;
+
+	const int N = 5*5*6*24*24;
+
+	for (int n = N * pos / size; n < N * (pos+1) / size; ++n) {
+		int idx = n;
+		const int i1 = ((idx /= 1	) % 5);
+		const int i2 = ((idx /= 5	) % 5);
+		const int i3 = ((idx /= 5	) % 6);
+		const int i4 = ((idx /= 6	) % 24);
+		const int i5 = ((idx /= 24	) % 24);
+
+		atomicAdd(&preact[i3][i4][i5], weight[i3][i1][i2] * input[i4 + i1][i5 + i2]);
+	}
+}
+
+__global__ void bias_c1(float preact[6][24][24], float bias[6])
+{
+	const int pos = blockIdx.x * blockDim.x + threadIdx.x;
+	const int size = blockDim.x * gridDim.x;
+
+	const int N = 6*24*24;
+
+	for (int n = N * pos / size; n < N * (pos+1) / size; ++n) {
+		int idx = n;
+		const int i1 = ((idx /= 1	) % 6);
+		const int i2 = ((idx /= 6	) % 24);
+		const int i3 = ((idx /= 24	) % 24);
+
+		preact[i1][i2][i3] += bias[i1];
+	}
+}
+
+__global__ void preact_s1(float input[6][24][24], float preact[6][6][6], float weight[1][4][4])
+{
+	const int pos = blockIdx.x * blockDim.x + threadIdx.x;
+	const int size = blockDim.x * gridDim.x;
+
+	const int N = 4*4*6*6*6;
+
+	for (int n = N * pos / size; n < N * (pos+1) / size; ++n) {
+		int idx = n;
+		const int i1 = ((idx /= 1	) % 4);
+		const int i2 = ((idx /= 4	) % 4);
+		const int i3 = ((idx /= 4	) % 6);
+		const int i4 = ((idx /= 6	) % 6);
+		const int i5 = ((idx /= 6	) % 6);
+
+		atomicAdd(&preact[i3][i4][i5], weight[0][i1][i2] * input[i3][i4 * 4 + i1][i5 * 4 + i2]);
+	}
+}
+
+__global__ void bias_s1(float preact[6][6][6], float bias[1])
+{
+	const int pos = blockIdx.x * blockDim.x + threadIdx.x;
+	const int size = blockDim.x * gridDim.x;
+
+	const int N = 6*6*6;
+
+	for (int n = N * pos / size; n < N * (pos+1) / size; ++n) {
+		int idx = n;
+		const int i1 = ((idx /= 1	) % 6);
+		const int i2 = ((idx /= 6	) % 6);
+		const int i3 = ((idx /= 6	) % 6);
+
+		preact[i1][i2][i3] += bias[0];
+	}
+}
+
+__global__ void preact_f(float input[6][6][6], float preact[10], float weight[10][6][6][6])
+{
+	const int pos = blockIdx.x * blockDim.x + threadIdx.x;
+	const int size = blockDim.x * gridDim.x;
+
+	const int N = 10*6*6*6;
+
+	for (int n = N * pos / size; n < N * (pos+1) / size; ++n) {
+		int idx = n;
+		const int i1 = ((idx /= 1	) % 10);
+		const int i2 = ((idx /= 10	) % 6);
+		const int i3 = ((idx /= 6	) % 6);
+		const int i4 = ((idx /= 6	) % 6);
+
+		atomicAdd(&preact[i1], weight[i1][i2][i3][i4] * input[i2][i3][i4]);
+	}
 }
 
 __global__ void bias_f(float preact[10], float bias[10])

#include <cublas_v2.h>

const static float dt = 1.0E-01f;
const static float threshold = 1.3E-01f;

struct Layer {
	float *output;

			l_f.bp_clear();
			l_s1.bp_clear();
			l_c1.bp_clear();

			makeError<<<10, 1>>>(l_f.d_preact, l_f.output, train_set[i].label, 10);

			bp_weight_f<<<64, 64>>>((float (*)[6][6][6])l_f.d_weight, l_f.d_preact, (float (*)[6][6])l_s1.output);
			bp_bias_f<<<64, 64>>>(l_f.bias, l_f.d_preact);

			bp_output_s1<<<64, 64>>>((float (*)[6][6])l_s1.d_output, (float (*)[6][6][6])l_f.weight, l_f.d_preact);
			bp_preact_s1<<<64, 64>>>((float (*)[6][6])l_s1.d_preact, (float (*)[6][6])l_s1.d_output, (float (*)[6][6])l_s1.preact);
			bp_weight_s1<<<64, 64>>>((float (*)[4][4])l_s1.d_weight, (float (*)[6][6])l_s1.d_preact, (float (*)[24][24])l_c1.output);
			bp_bias_s1<<<64, 64>>>(l_s1.bias, (float (*)[6][6])l_s1.d_preact);

			bp_output_c1<<<64, 64>>>((float (*)[24][24])l_c1.d_output, (float (*)[4][4])l_s1.weight, (float (*)[6][6])l_s1.d_preact);
			bp_preact_c1<<<64, 64>>>((float (*)[24][24])l_c1.d_preact, (float (*)[24][24])l_c1.d_output, (float (*)[24][24])l_c1.preact);
			bp_weight_c1<<<64, 64>>>((float (*)[5][5])l_c1.d_weight, (float (*)[24][24])l_c1.d_preact, (float (*)[28])l_input.output);
			bp_bias_c1<<<64, 64>>>(l_c1.bias, (float (*)[24][24])l_c1.d_preact);


			apply_grad<<<64, 64>>>(l_f.weight, l_f.d_weight, l_f.M * l_f.N);
			apply_grad<<<64, 64>>>(l_s1.weight, l_s1.d_weight, l_s1.M * l_s1.N);
			apply_grad<<<64, 64>>>(l_c1.weight, l_c1.d_weight, l_c1.M * l_c1.N);
		}
