Purpose: For education purposes only. The code demonstrates supervised learning task using a very simple neural network.
Reference: inspired by Andrew Trask‘s post.
Here is a follow-up post featuring a little bit more complicated code:
Neural Network in C++ (Part 2: MNIST Handwritten Digits Dataset)
The core component of the code, the learning algorithm, is only 10 lines:
| __global__ void kFit( const float* X, const int X_w, const int X_h, const float* y, const int y_w, float* l1, const int l1_w, float* l_1_d, float* pred, float* pred_d, float* W0, float* W1, float* buffer) { | |
| for (unsigned i = 0; i < 50; ++i) { | |
| dSigmoid(dDot(X, W0, l1, X_h, X_w, l1_w), l1, X_h, l1_w); | |
| dSigmoid(dDot(l1, W1, pred, X_h, l1_w, y_w), pred, X_h, y_w); | |
| dMartixByMatrixElementwise(dMartixSubstractMatrix(y, pred, pred_d, X_h, y_w), dSigmoid_d(pred, buffer, X_h, y_w), pred_d, X_h, y_w ); | |
| dMartixByMatrixElementwise(dDot_m1_m2T(pred_d, W1, l_1_d, X_h, y_w, l1_w), dSigmoid_d(l1, buffer, X_h, l1_w), l_1_d, X_h, l1_w); | |
| dDot_m1T_m2( l1, pred_d, W1, X_h, l1_w, y_w ); | |
| dDot_m1T_m2( X, l_1_d, W0, X_h, X_w, l1_w ); | |
| } | |
| } |
The loop above runs for 50 iterations (epochs) and fits the vector of attributes X to the vector of classes y. I am going to use 4 records from Iris flower dataset. The attributes (X) are sepal length, sepal width, petal length, and petal width. In my example, I have 2 (Iris Setosa (0) and Iris Virginica (1)) of 3 classes you can find in the original dataset. Predictions are stored in vector pred.
Neural network architecture. Values of vectors W0, W1, layer_1 and pred change over the course of training the network, while vectors X and y must not be changed:
| X W0 layer_1 W1 pred y | |
| 5.1 3.5 1.4 0.2 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.05 0 | |
| 4.9 3.0 1.4 0.2 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.06 0 | |
| 6.2 3.4 5.4 2.3 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.93 1 | |
| 5.9 3.0 5.1 1.8 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.92 1 | |
| 0.5 | |
| 0.5 | |
| 0.5 | |
| 0.5 |
The size of matrix X is the size of the batch by the number of attributes.
Line 3. Finding the values of the hidden layer:
| dSigmoid(dDot(X, W0, l1, X_h, X_w, l1_w), l1, X_h, l1_w); |
In order to calculate the hidden layer, first of all, we will need to multiply a 4 x 4 matrix X by a 4 x 4 matrix W0. Then, we will need to apply an activation function; in this case, we will use a sigmoid function.
A subroutine for matrix multiplication:
| __global__ void kDot(const float *m1, const float *m2, float *output, const int m1_rows , const int m1_columns, const int m2_columns ){ | |
| /* Computes the product of two matrices: m1 x m2. | |
| Inputs: | |
| m1: array, left matrix of size m1_rows x m1_columns | |
| m2: array, right matrix of size m1_columns x m2_columns (the number of rows in the right matrix | |
| must be equal to the number of the columns in the left one) | |
| output: array, the results of the computation are to be stored here: | |
| m1 * m2, product of two arrays m1 and m2, a matrix of size m1_rows x m2_columns | |
| m1_rows: int, number of rows in the left matrix m1 | |
| m1_columns: int, number of columns in the left matrix m1 | |
| m2_columns: int, number of columns in the right matrix m2 | |
| */ | |
| const int id = blockIdx.x * blockDim.x + threadIdx.x; | |
| const int r = (int)id / m2_columns; | |
| const int c = id % m2_columns; | |
| float t_output = 0.f; | |
| for( int k = 0; k < m1_columns; ++k ) { | |
| t_output += m1[ r * m1_columns + k ] * m2[ k * m2_columns + c ]; | |
| } | |
| output[ id ] = t_output; | |
| } | |
| __device__ float* dDot(const float *m1, const float *m2, float *output, const int m1_rows , const int m1_columns, const int m2_columns ) | |
| { | |
| kDot <<< m1_rows, m2_columns >>> (m1, m2, output, m1_rows , m1_columns, m2_columns ); | |
| cudaDeviceSynchronize(); | |
| return output; | |
| } |
A subroutine for the sigmoid function:
| __global__ void kSigmoid(float const *input, float *output) { | |
| /* Computes the value of the sigmoid function f(x) = 1/(1 + e^-x). | |
| Inputs: | |
| input: array | |
| output: array, the results of the computation are to be stored here | |
| */ | |
| const int id = blockIdx.x * blockDim.x + threadIdx.x; | |
| output[id] = 1.0 / (1.0 + std::exp(-input[id])); | |
| } | |
| __device__ void dSigmoid(float const *input, float *output, const int height, const int width){ | |
| kSigmoid <<< height, width >>> (input, output); | |
| cudaDeviceSynchronize(); | |
| } |
Sigmoid function (red) and its first derivative (blue graph):

Line 4. Finding the matrix with predictions pred. In order to do so, we will need to multiply a 4 x 8 matrix l1 by a 8 x 1 matrix W1. Then, we will need to apply an activation function:
| dSigmoid(dDot(l1, W1, pred, X_h, l1_w, y_w), pred, X_h, y_w); |
Line 5. Determine the vector of prediction errors pred_d. First, subtract pred from y. Then, calculate sigmoid( pred ) and, finally, multiply (elementwise) the result of these two operations.
| dMartixByMatrixElementwise(dMartixSubstractMatrix(y, pred, pred_d, X_h, y_w), dSigmoid_d(pred, buffer, X_h, y_w), pred_d, X_h, y_w ); |
CUDA kernel which finds the difference between two matrices:
| __global__ void kMartixSubstractMatrix(const float *m1, const float *m2, float *output) { | |
| /* Computes the (elementwise) difference between two arrays | |
| Inputs: | |
| m1: array | |
| m2: array | |
| output: array,the results of the computation are to be stored here | |
| */ | |
| const int id = blockIdx.x * blockDim.x + threadIdx.x; | |
| output[id] = m1[id] - m2[id]; | |
| } | |
| __device__ float* dMartixSubstractMatrix(const float *m1, const float *m2, float *output, const int width, const int height){ | |
| kMartixSubstractMatrix <<< width, height >>> ( m1, m2, output ); | |
| cudaDeviceSynchronize(); | |
| return output; | |
| } |
Elemetwise multiplicaton of two vectors:
| __global__ void kMartixByMatrixElementwise(const float *m1, const float *m2, float *output) { | |
| /* Computes the product of two arrays (elementwise multiplication). | |
| Inputs: | |
| m1: array | |
| m2: array | |
| output: array,the results of the multiplication are to be stored here | |
| */ | |
| const int id = blockIdx.x * blockDim.x + threadIdx.x; | |
| output[id] = m1[id] * m2[id]; | |
| } | |
| __device__ float* dMartixByMatrixElementwise(const float *m1, const float *m2, float *output, const int width, const int height){ | |
| kMartixByMatrixElementwise <<< width, height >>> ( m1, m2, output ); | |
| cudaDeviceSynchronize(); | |
| return output; | |
| } |
Line 6. Back propagate the prediction errors to l_1_d. First, multiply pred_d by transposed W1. Then, calculate sigmoid( l1 ) and, finally, multiply (elementwise) the result of these two operations.
| dMartixByMatrixElementwise(dDot_m1_m2T(pred_d, W1, l_1_d, X_h, y_w, l1_w), dSigmoid_d(l1, buffer, X_h, l1_w), l_1_d, X_h, l1_w); |
A subroutine that multiplies matrix by transposed matrix:
| __global__ void kDot_m1_m2T(const float *m1, const float *m2, float *output, const int m1_columns, const int m2_rows ){ | |
| /* Updates the output matrix with the product of two matrices: m1 and m2 transposed. | |
| Inputs: | |
| m1: array, left matrix of size m1_rows x m1_columns | |
| m2: array, right matrix of size m2_rows x m1_columns (m2 transposed will be of size m1_columns x m2_rows) | |
| output: array, the results of the computation are to be stored here: | |
| m1 * m2, product of two arrays m1 and m2, a matrix of size m1_rows x m2_rows | |
| m1_columns: int, number of columns in the left matrix m1 | |
| m2_rows: int, number of rows in the left matrix m2 | |
| */ | |
| const int id = blockIdx.x * blockDim.x + threadIdx.x; | |
| const int r = (int)id / m2_rows; | |
| const int c = id % m2_rows; | |
| float t_output = 0.0; | |
| int id_T; | |
| for( int k = 0; k < m1_columns; ++k ) { | |
| id_T = c * m1_columns + k; | |
| t_output += m1[ r * m1_columns + k ] * m2[ id_T ]; | |
| } | |
| output[ id ] = t_output; | |
| } | |
| __device__ float* dDot_m1_m2T(const float *m1, const float *m2, float *output, const int m1_rows , const int m1_columns, const int m2_rows ) | |
| { | |
| kDot_m1_m2T <<< m1_rows, m2_rows >>> ( m1, m2, output, m1_columns, m2_rows ); | |
| cudaDeviceSynchronize(); | |
| return output; |
Line 7. Update weights W1 with the result of matrix multiplication of transposed l1 and pred_d:
This line computes weight updates. In order to do that, we need to perform matrix multiplication of transposed matrix X by matrix pred_delta.
| vector W_delta = dot(transpose( &X[0], 4, 4 ), pred_delta, 4, 4, 1); |
Line 8. Update weights W0 with the result of matrix multiplication of transposed X and l_1_d:
| dDot_m1T_m2( X, l_1_d, W0, X_h, X_w, l1_w ); |
Complete code:
Output:
| Prediction[0] : 0.060997 True Value[0] : 0.000000 Error[0] : 0.060997 | |
| Prediction[1] : 0.076193 True Value[1] : 0.000000 Error[1] : 0.076193 | |
| Prediction[2] : 0.927551 True Value[2] : 1.000000 Error[2] : -0.072449 | |
| Prediction[3] : 0.918263 True Value[3] : 1.000000 Error[3] : -0.081737 |
Compile…
| nvcc -arch=sm_50 -rdc=true -lcudadevrt onehiddenlayerperceptron.cu -o perceptron |
… and run
| ./perceptron |






