OxNK x NC= N x NC * NxOxNK x NC
O = flattened output dimension
NK = Number of Kernels
NC = Number of channels = RGB
N = flattened input dimension
(This is an amended version of the previous equation for Feedforwarding a Convolutional Neural Network:
OxNxNKxNC = NxKxNC * KxOxNK)
In order to calculate this using CUDA I could not iterate NK,N, and O all in parallel without using the Tensorflow Method (GPU_1D_KERN_LOOP) and my Transkern array:
- Code: Select all
GPU_1D_KERN_LOOP(index, O_O*N_N*4){
const int nk = (index / 1 )% 4;
const int col = (index / 4 )% O_O;
const int row = (index / 4 /O_O)% N_N;
atomicAdd(&O_a[col+nk*O_O],I_a[row] * Kernel[nk*25+Transkern[row+col*N_N]]);
}
NB. size of kernel = 5x5 = 25, Number of kernels = 4
The Transkern array produces a connection matrix using the Output and Input dimensions O_O x N_N there are NK Transkern matrix's.
The TF routine is implemented using a basic config using CUDA streams and an iterator defined as follows:
(This is copied straight from TF source code!)
- Code: Select all
/*Am copying googles tensorflow here - to learn*/
namespace detail {
template <typename T>
class GpuGridRange {
struct Iterator {
__device__ Iterator(T index, T delta) : index_(index), delta_(delta) {}
__device__ T operator*() const { return index_;}
__device__ Iterator& operator++() {
index_ += delta_;
return *this;
}
__device__ bool operator!=(const Iterator& other) const {
bool greater = index_ > other.index_;
bool less = index_ < other.index_;
if(!other.delta_){
return less;
}
if(!delta_){
return greater;
}
return less || greater;
}
private:
T index_;
const T delta_;
}; //end Iterator struct
public:
__device__ GpuGridRange(T begin,T delta,T end)
: begin_(begin),delta_(delta),end_(end) {}
__device__ Iterator begin() const {return Iterator(begin_,delta_); }
__device__ Iterator end() const {return Iterator(end_,0);}
private:
T begin_;
T delta_;
T end_;
}; //end GPU class class
}; //end namespace detail
template <typename T> //Allows you to use GPU iterator with all data types
__device__ detail::GpuGridRange<T> GpuGridRangeX(T count) {
return detail::GpuGridRange<T>(
/*begin*/blockIdx.x * blockDim.x + threadIdx.x,
/*delta*/gridDim.x * blockDim.x, /*end*/count
);
}
You can see the familiar CUDA code being used by the iterator defined above. This iterates for any data type across dimension x of the CUDA grid, block's and threads.
It is then defined as a simple loop using the following simple define:
- Code: Select all
#define GPU_1D_KERN_LOOP(i, n) \
for (int i : ::GpuGridRangeX<int>(n))
Here is the above routine with config (Note the important inclusion of CUDA streams:
- Code: Select all
cudaStream_t stream1;
cudaStreamCreate(&stream1);
cudaMalloc((void **)&O_a, O_O * NK * sizeof(float));
cudaMemcpyAsync(O_a, Oarray, O_O * NK * sizeof(float), cudaMemcpyHostToDevice,stream1);
myroutine<<<1024,1024,0,stream1>>>(K_a,T_a,O_a,I_a,O_O,N_N);
I believe this method to be the best one for implementing CUDA for processing Deep Neural Nets and should make full use of the above methods used by TensorFlow.