check point

Author: dwha

cifar10.cpp

@@ -193,12 +193,12 @@ int cifar10()
 	trainingDataFilenames.push_back("cifar-10/data_batch_5.bin");
 	std::vector<ff::CudaTensor> trainingImages;
 	std::vector<ff::CudaTensor> trainingLabels;
-	LoadCifar10(kBatchSize, 50000, false, trainingDataFilenames, trainingImages, trainingLabels);
+	LoadCifar10(kBatchSize, 5000, false, trainingDataFilenames, trainingImages, trainingLabels);
 	std::vector<std::string> testDataFilenames;
 	testDataFilenames.push_back("cifar-10/test_batch.bin");
 	std::vector<ff::CudaTensor> testImages;
 	std::vector<ff::CudaTensor> testLabels;
-	LoadCifar10(kBatchSize, 10000, false, testDataFilenames, testImages, testLabels);
+	LoadCifar10(kBatchSize, 1000, false, testDataFilenames, testImages, testLabels);
 
 #if 1
 	ff::CudaNn nn;

ffCudaNn.cpp

@@ -500,7 +500,7 @@ namespace ff
 			int nJobs = _wG._dataSize;
 			int numBlocks = (nJobs + K_THREAD_PER_BLOCK - 1) / K_THREAD_PER_BLOCK;
 			dim3 blocks(numBlocks), threads(K_THREAD_PER_BLOCK);
-			BackwardConv2d_Wg_Cuda <<<blocks, threads>>> (
+			BackwardConv2d_Wg_Cuda <<<blocks, threads >>> (
 				_wG._dataGpu, _pX->_dataGpu, yG->_dataGpu,
 				_wG._d3, _wG._d2,
 				_y._d3, _y._d1, _y._d0, _pX->_d1, _pX->_d0,
@@ -864,13 +864,14 @@ namespace ff
 		int ch = blockIdx.x;
 		int image = threadIdx.x;
 
-		__shared__ float meanArr[BLOCK_SIZE];
-		meanArr[image] = 0.0f;
 		int mDash = nImages * nRow * nCol;
 		int imageStride = nChannel * nRow * nCol;
 		int channelStride = nRow * nCol;
 		int currChBaseIndex = ch * channelStride;
 		int baseIndex = image * imageStride + currChBaseIndex;
+
+		__shared__ float meanArr[BLOCK_SIZE];
+		meanArr[image] = 0.0f;
 		for (int i = 0; i < channelStride; ++i)
 		{
 			meanArr[image] += x[baseIndex + i];
@@ -900,10 +901,13 @@ namespace ff
 
 		float alpha = w[ch * 2 + 0];
 		float beta = w[ch * 2 + 1];
-		meanAndVariance[ch * 2 + 0] = mean;
-		meanAndVariance[ch * 2 + 1] = variance;
-		meanAndVarianceAcc[(ch + 1) * 2 + 0] += mean;
-		meanAndVarianceAcc[(ch + 1) * 2 + 1] += variance;
+		if (threadIdx.x == 0)
+		{
+			meanAndVariance[ch * 2 + 0] = mean;
+			meanAndVariance[ch * 2 + 1] = variance;
+			meanAndVarianceAcc[(ch + 1) * 2 + 0] += mean;
+			meanAndVarianceAcc[(ch + 1) * 2 + 1] += variance;
+		}
 		float d = rsqrtf(variance + 1e-8f);
 		for (int i = 0; i < channelStride; ++i)
 		{
@@ -947,8 +951,7 @@ namespace ff
 
 		if (_nn->IsTraining())
 		{
-			++_accCount;
-			if (_accCount <= 24)
+			if (++_accCount <= 24)
 			{
 				ForwardBatchNorm2d_Train_0_Cuda <<< 1, 1 >>> (_meanAndVarianceAcc._dataGpu);
 			}

main.cpp

@@ -7,6 +7,7 @@ int cifar10();
 int simple()
 {
 #if 1
+	float learningRate = 0.01f;
 	ff::CudaNn nn;
 	nn.AddFc(1000, 4096);
 	nn.AddFc(4096, 1024);
@@ -22,6 +23,7 @@ int simple()
 	x.SetRandom();
 	y.SetRandom();
 #else
+	float learningRate = 0.001f;
 	ff::CudaNn nn;
 	nn.AddConv2d(3, 1, 8, 1, 1);		// 8 * 8 * 8
 	nn.AddRelu();
@@ -39,7 +41,6 @@ int simple()
 	y.SetRandom();
 #endif
 
-	float learningRate = 0.0001f;
 	const ff::CudaTensor* yPred = nullptr;
 	for (int i = 0; i < 10000; ++i)
 	{

mnist.cpp

@@ -125,7 +125,7 @@ int mnist()
 	LoadMnistData("mnist/train-images.idx3-ubyte", "mnist/train-labels.idx1-ubyte", kBatchSize, trainingImages, trainingLabels);
 	LoadMnistData("mnist/t10k-images.idx3-ubyte", "mnist/t10k-labels.idx1-ubyte", kBatchSize, testImages, testLabels);
 
-#if 0
+#if 1
 	float learningRate = 0.001f;
 	ff::CudaNn nn;
 	nn.AddFc(28 * 28, 2048);
@@ -142,7 +142,7 @@ int mnist()
 		testImages[i].Reshape(28, 28, 1, testImages[i]._dataSize / (28 * 28));
 	}
 
-	float learningRate = 0.0001f;
+	float learningRate = 0.001f;
 	ff::CudaNn nn;
 	nn.AddConv2d(3, 1, 4, 1, 1);
 	nn.AddBatchNorm2d(4);