Remove minor code duplication

.gitignore

@@ -30,6 +30,8 @@ BlackCat_Tensors.sln
 /include/bc_py.py
 /examples/mnist_test/mnist_test_exec_cu
 /examples/mnist_test/mnist_test_exec
+/examples/cifar10/cifar_test
+/examples/cifar10/cifar_test_cu
 /examples/mnist_test_recurrent/mnist_test_exec_cu
 /examples/mnist_test_recurrent/mnist_test_exec
 /examples/mnist_test*/mnist_test_exec*

blackcat/caffe/img2col.cu

@@ -17,6 +17,7 @@
 #include <algorithm>
 #include "img2col.h"
 #include <iostream>
+#include "../common.h"
 
 namespace bc {
 namespace caffe {

blackcat/caffe/img2col.h

@@ -8,6 +8,9 @@
 #ifndef BLACKCAT_TENSORS_CAFFE_IMG2COL_H_
 #define BLACKCAT_TENSORS_CAFFE_IMG2COL_H_
 
+#include "../common.h"
+#include <vector>
+
 /**
  * THIS FILE HAS BEEN MODIFIED FROM ITS ORIGINAL SOURCE
  */

blackcat/random.h

@@ -0,0 +1,2 @@
+#include "common.h"
+#include "random/random.h"

blackcat/random/device.cu

@@ -0,0 +1,89 @@
+/*
+ * Device.h
+ *
+ *  Created on: Dec 3, 2018
+ *      Author: joseph
+ */
+
+
+#ifdef __CUDACC__
+#ifndef DEVICE_H_
+#define DEVICE_H_
+
+#include <cuda_runtime_api.h>
+#include <cuda.h>
+#include <curand.h>
+#include <curand_kernel.h>
+#include "../common.h"
+
+namespace bc {
+
+struct device_tag;
+
+namespace random {
+namespace device_impl {
+
+__global__
+static void bc_curand_init(curandState_t* state) {
+
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+	if (i==0)
+		curand_init(0,0,0,state);
+}
+
+static constexpr unsigned float_decimal_length = 100000;
+template<class T> __global__
+static void randomize(curandState_t* state, T t, float lower_bound, float upper_bound) {
+
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+	for (; i < t.size(); i += blockDim.x * gridDim.x) {
+		curandState_t tmpstate = *state;
+		skipahead(i, &tmpstate);
+		t[i] = curand(&tmpstate) % float_decimal_length;
+		t[i] /= float_decimal_length;
+		t[i] *= (upper_bound - lower_bound);
+		t[i] += lower_bound;
+	}
+
+	__syncthreads();
+	if (i == 0)
+		skipahead(t.size(), state);
+}
+}
+
+template<class SystemTag>
+struct Random;
+
+template<>
+struct Random<device_tag> {
+
+	static curandState_t* bc_curand_state() {
+		struct bc_curandState_t {
+		    curandState_t* state = nullptr;
+
+			bc_curandState_t() {
+				BC_CUDA_ASSERT((cudaMalloc((void**) &state, sizeof(curandState_t))));
+				device_impl::bc_curand_init<<<1,1>>>(state);
+				cudaDeviceSynchronize();
+			}
+			~bc_curandState_t() {
+				BC_CUDA_ASSERT((cudaFree((void*)state)));
+			}
+		};
+		thread_local bc_curandState_t bc_state;
+		return bc_state.state;
+	}
+
+	template<class Stream, typename T>
+	static void randomize(Stream stream, T t, float lower_bound, float upper_bound) {
+		device_impl::randomize<<<calculate_block_dim(t.size()),calculate_threads(t.size()), 0, stream>>>(
+				bc_curand_state(), t, lower_bound, upper_bound);
+	}
+};
+
+}
+}
+
+
+#endif
+#endif /* DEVICE_H_ */

blackcat/random/host.h

@@ -11,25 +11,42 @@
 #include <random>
 
 namespace bc {
+
+struct host_tag;
+
 namespace random {
 
+template<class SystemTag>
+struct Random;
+
 template<>
 struct Random<host_tag> {
-	 template<typename T, typename value_type>
-	    static void randomize_kernel(T& tensor, value_type lower_bound, value_type upper_bound) {
-	 BC_omp_for__
-	        for (int i = 0; i < tensor.size(); ++i) {
-	            tensor[i] = ((value_type) (std::rand() / ((value_type) RAND_MAX + 1)) * (upper_bound - lower_bound)) + lower_bound;
-	        }
-	 BC_omp_bar__
-	    }
-
-	 template<class Stream, typename T, typename value_type>
-	    static void randomize(Stream stream, T& tensor, value_type lower_bound, value_type upper_bound) {
-		 stream.enqueue([&](){
-			 randomize_kernel(tensor, lower_bound, upper_bound);
-		 });
-	    }
+	template<class Tensor, class value_type>
+	static void randomize_kernel(
+			Tensor& tensor,
+			value_type lower_bound,
+			value_type upper_bound)
+	{
+		BC_omp_for__
+		for (int i = 0; i < tensor.size(); ++i) {
+			tensor[i] =
+					((value_type(std::rand()) / ((value_type) RAND_MAX + 1))
+					* (upper_bound - lower_bound)) + lower_bound;
+		}
+		BC_omp_bar__
+	}
+
+	template<class Stream, class Tensor, class value_type>
+		static void randomize(
+				Stream stream,
+				Tensor& tensor,
+				value_type lower_bound,
+				value_type upper_bound)
+	{
+		stream.enqueue([&](){
+			randomize_kernel(tensor, lower_bound, upper_bound);
+		});
+	 }
 };
 
 

blackcat/random/random.h

@@ -9,9 +9,7 @@
 #ifndef BC_RANDOM_RANDOM_H_
 #define BC_RANDOM_RANDOM_H_
 
-BC_DEFAULT_MODULE_BODY(random, Random)
-
 #include "host.h"
-#include "device.h"
+#include "device.cu"
 
 #endif /* RANDOM_H_ */

blackcat/tensors.h

@@ -15,8 +15,8 @@
 #include "utility.h"
 #include "operations.h"
 #include "algorithms.h"
+#include "random.h"
 
-#include "random/random.h"
 #include "tensors/tensors.h"
 
 namespace bc {

blackcat/tensors/expression_templates/blas_expression_parser/device_impl.cu

@@ -1,3 +1,5 @@
+#include <type_traits>
+
 namespace bc {
 namespace tensors {
 namespace exprs { 

blackcat/tensors/expression_templates/functions/reductions/device.cu

@@ -9,6 +9,7 @@
 #ifndef BLACKCATTENSORS_TENSORS_FUNCTIONS_REDUCTIONS_DEVICE_REDUCE_CU_
 #define BLACKCATTENSORS_TENSORS_FUNCTIONS_REDUCTIONS_DEVICE_REDUCE_CU_
 
+#include "../../../../common.h"
 #include <cuda_runtime_api.h>
 #include <cuda.h>
 

blackcat/tensors/expression_templates/nd_evaluator/device.cu

@@ -10,7 +10,7 @@
 #ifndef BC_MATHEMATICS_DEVICE_H_
 #define BC_MATHEMATICS_DEVICE_H_
 
-#include <cublas_v2.h>
+#include "../../../common.h"
 #include <cuda_runtime_api.h>
 #include <cuda.h>
 
@@ -19,10 +19,16 @@
 
 
 namespace bc {
+
+struct device_tag;
+
 namespace tensors {
 namespace exprs {
 namespace evaluator {
 
+template<class SystemTag>
+struct Evaluator;
+
 template<>
 struct Evaluator<device_tag> {
 

blackcat/tensors/expression_templates/nd_evaluator/evaluator.h

@@ -5,13 +5,10 @@
  *      Author: joseph
  */
 
-#ifndef EVALUATOR_H_
-#define EVALUATOR_H_
-
-BC_DEFAULT_MODULE_BODY(tensors { namespace exprs { namespace evaluator , Evaluator) } }
+#ifndef BLACKCAT_TENSORS_EXPRS_EVALUATOR_H_
+#define BLACKCAT_TENSORS_EXPRS_EVALUATOR_H_
 
 #include "device.cu"
 #include "host.h"
 
-
 #endif /* EVALUATOR_H_ */