diff --git a/include/kernels/activations.cuh b/include/kernels/activations.cuh index 5f7935b..cbd91a2 100644 --- a/include/kernels/activations.cuh +++ b/include/kernels/activations.cuh @@ -10,4 +10,10 @@ relu_kernel(const float* __restrict__ src, float* __restrict__ dst, int len); __global__ void linear_kernel(const float* __restrict__ src, float* __restrict__ dst, int len); +enum Activation { + SIGMOID, + RELU, + LINEAR +}; + #endif // ACTIVATIONS_H \ No newline at end of file diff --git a/include/layers/dense.cuh b/include/layers/dense.cuh index b1d7fbf..12be07d 100644 --- a/include/layers/dense.cuh +++ b/include/layers/dense.cuh @@ -16,7 +16,7 @@ class Dense : public ILayer { Dense( int inputSize, int outputSize, - std::string activation, + Activation activation, cublasHandle_t cublasHandle ); ~Dense(); @@ -37,7 +37,7 @@ class Dense : public ILayer { std::vector weights; std::vector biases; - std::string activation; + Activation activation; void initializeWeights(); void initializeBiases(); diff --git a/src/layers/dense.cu b/src/layers/dense.cu index 227b53b..2bacd59 100644 --- a/src/layers/dense.cu +++ b/src/layers/dense.cu @@ -13,7 +13,7 @@ Layers::Dense::Dense( int inputSize, int outputSize, - std::string activation, + Activation activation, cublasHandle_t cublasHandle ) : inputSize(inputSize), @@ -68,18 +68,24 @@ void Layers::Dense::forward(const float* d_input, float* d_output) { int threadsPerBlock = 256; int blocksPerGrid = (outputSize + threadsPerBlock - 1) / threadsPerBlock; - if (activation == "sigmoid") { - sigmoid_kernel<<>>( - d_output, d_output, outputSize - ); - } else if (activation == "relu") { - relu_kernel<<>>( - d_output, d_output, outputSize - ); - } else { - linear_kernel<<>>( - d_output, d_output, outputSize - ); + switch (activation) { + case SIGMOID: + sigmoid_kernel<<>>( + d_output, d_output, outputSize + ); + break; + + case RELU: + relu_kernel<<>>( + d_output, d_output, outputSize + ); + break; + + default: + linear_kernel<<>>( + d_output, d_output, outputSize + ); + break; } CUDA_CHECK(cudaDeviceSynchronize()); diff --git a/test/layers/test_dense.cu b/test/layers/test_dense.cu index 8221ca1..323125b 100644 --- a/test/layers/test_dense.cu +++ b/test/layers/test_dense.cu @@ -17,7 +17,7 @@ class DenseLayerTest : public CublasTestFixture { std::vector& biases, float*& d_input, float*& d_output, - std::string activation + Activation activation ) { // Create Dense layer Layers::Dense denseLayer( @@ -63,7 +63,7 @@ TEST_F(DenseLayerTest, Init) { // std::cout << "Dense layer: input size = " << inputSize << ", // output size = " << outputSize << std::endl; Layers::Dense denseLayer( - inputSize, outputSize, "sigmoid", cublasHandle + inputSize, outputSize, SIGMOID, cublasHandle ); } } @@ -81,7 +81,7 @@ TEST_F(DenseLayerTest, setWeights) { {1.3f, 0.5f, 0.0f, 1.7f} }; - Layers::Dense denseLayer(inputSize, outputSize, "sigmoid", cublasHandle); + Layers::Dense denseLayer(inputSize, outputSize, SIGMOID, cublasHandle); denseLayer.setWeights(weights); } @@ -108,8 +108,7 @@ TEST_F(DenseLayerTest, ForwardUnitWeightMatrixLinear) { float* d_output; Layers::Dense denseLayer = commonTestSetup( - inputSize, outputSize, input, weights, biases, d_input, d_output, - "linear" + inputSize, outputSize, input, weights, biases, d_input, d_output, LINEAR ); denseLayer.forward(d_input, d_output); @@ -145,7 +144,7 @@ TEST_F(DenseLayerTest, ForwardRandomWeightMatrixRelu) { float* d_output; Layers::Dense denseLayer = commonTestSetup( - inputSize, outputSize, input, weights, biases, d_input, d_output, "relu" + inputSize, outputSize, input, weights, biases, d_input, d_output, RELU ); denseLayer.forward(d_input, d_output); @@ -188,7 +187,7 @@ TEST_F(DenseLayerTest, ForwardRandomWeightMatrixSigmoid) { Layers::Dense denseLayer = commonTestSetup( inputSize, outputSize, input, weights, biases, d_input, d_output, - "sigmoid" + SIGMOID ); denseLayer.forward(d_input, d_output);