From 48ba09b28db1395f52cd96a426bc74f3e9a33c7d Mon Sep 17 00:00:00 2001 From: LordMathis Date: Tue, 27 Feb 2024 18:51:22 +0100 Subject: [PATCH] Format source code using clang-format --- .clang-format | 20 +++++++ include/functions/activations.cuh | 11 ++-- include/layers/conv.cuh | 35 +++++++----- include/layers/dense.cuh | 55 ++++++++++-------- include/layers/ilayer.cuh | 19 ++++--- src/functions/activations.cu | 56 +++++++++--------- src/layers/dense.cu | 72 ++++++++++++++++-------- src/utils/cuda_helper.cu | 6 +- test/layers/test_dense.cu | 75 ++++++++++++++++--------- test/test_utils/test_cublas_fixture.cu | 2 +- test/test_utils/test_cublas_fixture.cuh | 4 +- test/utils/test_functions.cu | 12 ++-- 12 files changed, 229 insertions(+), 138 deletions(-) create mode 100644 .clang-format diff --git a/.clang-format b/.clang-format new file mode 100644 index 0000000..e6e9c2e --- /dev/null +++ b/.clang-format @@ -0,0 +1,20 @@ +Language: Cpp +Standard: c++20 +BasedOnStyle: Google +AccessModifierOffset: -2 +AlignAfterOpenBracket: BlockIndent +AlignArrayOfStructures: Left +AlignConsecutiveAssignments: + Enabled: true +AlignConsecutiveBitFields: + Enabled: true +AlignConsecutiveDeclarations: + Enabled: true +AlignConsecutiveMacros: + Enabled: true +AllowAllArgumentsOnNextLine: false +AllowAllParametersOfDeclarationOnNextLine: false +AllowShortFunctionsOnASingleLine: Empty +AlwaysBreakBeforeMultilineStrings: false +BinPackParameters: false +IndentWidth: 4 \ No newline at end of file diff --git a/include/functions/activations.cuh b/include/functions/activations.cuh index 638756c..b557681 100644 --- a/include/functions/activations.cuh +++ b/include/functions/activations.cuh @@ -7,8 +7,11 @@ __device__ float sigmoid(float a); __device__ float relu(float a); __device__ float linear(float a); -__global__ void sigmoid_kernel(const float* __restrict__ src, float* __restrict__ dst, int len); -__global__ void relu_kernel(const float* __restrict__ src, float* __restrict__ dst, int len); -__global__ void linear_kernel(const float* __restrict__ src, float* __restrict__ dst, int len); +__global__ void +sigmoid_kernel(const float* __restrict__ src, float* __restrict__ dst, int len); +__global__ void +relu_kernel(const float* __restrict__ src, float* __restrict__ dst, int len); +__global__ void +linear_kernel(const float* __restrict__ src, float* __restrict__ dst, int len); -#endif // ACTIVATIONS_H \ No newline at end of file +#endif // ACTIVATIONS_H \ No newline at end of file diff --git a/include/layers/conv.cuh b/include/layers/conv.cuh index eed7d76..81bd1ad 100644 --- a/include/layers/conv.cuh +++ b/include/layers/conv.cuh @@ -5,22 +5,27 @@ namespace Layers { - class Conv { - public: - Conv(int inputSize, int outputSize, int kernelSize, cublasHandle_t cublasHandle); - ~Conv(); +class Conv { + public: + Conv( + int inputSize, + int outputSize, + int kernelSize, + cublasHandle_t cublasHandle + ); + ~Conv(); - void forward(const float* input, float* output); + void forward(const float* input, float* output); - private: - int inputSize; - int outputSize; - int kernelSize; - cublasHandle_t cublasHandle; - float* d_weights; - float* d_biases; - }; + private: + int inputSize; + int outputSize; + int kernelSize; + cublasHandle_t cublasHandle; + float* d_weights; + float* d_biases; +}; -} // namespace Layers +} // namespace Layers -#endif // CONV_LAYER_H +#endif // CONV_LAYER_H diff --git a/include/layers/dense.cuh b/include/layers/dense.cuh index 9393154..b1d7fbf 100644 --- a/include/layers/dense.cuh +++ b/include/layers/dense.cuh @@ -1,42 +1,49 @@ #ifndef DENSE_LAYER_H #define DENSE_LAYER_H -#include -#include #include + +#include #include +#include + #include "ilayer.cuh" namespace Layers { - class Dense : public ILayer { - public: - Dense(int inputSize, int outputSize, std::string activation, cublasHandle_t cublasHandle); - ~Dense(); +class Dense : public ILayer { + public: + Dense( + int inputSize, + int outputSize, + std::string activation, + cublasHandle_t cublasHandle + ); + ~Dense(); - void forward(const float* input, float* output); - void setWeights(const std::vector>& weights); - void setBiases(const std::vector& biases); + void forward(const float* input, float* output); + void setWeights(const std::vector>& weights); + void setBiases(const std::vector& biases); - private: - int inputSize; - int outputSize; + private: + int inputSize; + int outputSize; - cublasHandle_t cublasHandle; + cublasHandle_t cublasHandle; - float* d_weights; - float* d_biases; + float* d_weights; + float* d_biases; - std::vector weights; - std::vector biases; + std::vector weights; + std::vector biases; - std::string activation; + std::string activation; - void initializeWeights(); - void initializeBiases(); - void toCuda(); - }; + void initializeWeights(); + void initializeBiases(); + void toCuda(); +}; -} // namespace Layers +} // namespace Layers -#endif // DENSE_LAYER_H +#endif // DENSE_LAYER_H diff --git a/include/layers/ilayer.cuh b/include/layers/ilayer.cuh index a0007f0..40db713 100644 --- a/include/layers/ilayer.cuh +++ b/include/layers/ilayer.cuh @@ -3,19 +3,20 @@ #define I_LAYER_H #include + #include namespace Layers { - class ILayer { - public: - virtual ~ILayer() {} +class ILayer { + public: + virtual ~ILayer() {} - virtual void forward(const float* input, float* output) = 0; - virtual void setWeights(const std::vector>& weights) = 0; - virtual void setBiases(const std::vector& biases) = 0; - }; + virtual void forward(const float* input, float* output) = 0; + virtual void setWeights(const std::vector>& weights) = 0; + virtual void setBiases(const std::vector& biases) = 0; +}; -} // namespace Layers +} // namespace Layers -#endif // I_LAYERH \ No newline at end of file +#endif // I_LAYERH \ No newline at end of file diff --git a/src/functions/activations.cu b/src/functions/activations.cu index 57399b4..f48dc1c 100644 --- a/src/functions/activations.cu +++ b/src/functions/activations.cu @@ -1,44 +1,48 @@ -#include "activations.cuh" #include -__device__ float sigmoid(float a) -{ - return 1.0 / (1.0 + exp (-a)); +#include "activations.cuh" + +__device__ float sigmoid(float a) { + return 1.0 / (1.0 + exp(-a)); } -__device__ float relu(float a) -{ +__device__ float relu(float a) { return a < 0.0 ? 0.0 : a; } -__device__ float linear(float a) -{ +__device__ float linear(float a) { return a; } -__global__ void sigmoid_kernel(const float* __restrict__ src, float* __restrict__ dst, int len) { - int stride = gridDim.x * blockDim.x; - int tid = blockDim.x * blockIdx.x + threadIdx.x; +__global__ void sigmoid_kernel( + const float* __restrict__ src, + float* __restrict__ dst, + int len +) { + int stride = gridDim.x * blockDim.x; + int tid = blockDim.x * blockIdx.x + threadIdx.x; - for (int i = tid; i < len; i += stride) { - dst[i] = sigmoid(src[i]); - } + for (int i = tid; i < len; i += stride) { + dst[i] = sigmoid(src[i]); + } } -__global__ void relu_kernel(const float* __restrict__ src, float* __restrict__ dst, int len) { - int stride = gridDim.x * blockDim.x; - int tid = blockDim.x * blockIdx.x + threadIdx.x; +__global__ void +relu_kernel(const float* __restrict__ src, float* __restrict__ dst, int len) { + int stride = gridDim.x * blockDim.x; + int tid = blockDim.x * blockIdx.x + threadIdx.x; - for (int i = tid; i < len; i += stride) { - dst[i] = relu(src[i]); - } + for (int i = tid; i < len; i += stride) { + dst[i] = relu(src[i]); + } } -__global__ void linear_kernel(const float* __restrict__ src, float* __restrict__ dst, int len) { - int stride = gridDim.x * blockDim.x; - int tid = blockDim.x * blockIdx.x + threadIdx.x; +__global__ void +linear_kernel(const float* __restrict__ src, float* __restrict__ dst, int len) { + int stride = gridDim.x * blockDim.x; + int tid = blockDim.x * blockIdx.x + threadIdx.x; - for (int i = tid; i < len; i += stride) { - dst[i] = linear(src[i]); - } + for (int i = tid; i < len; i += stride) { + dst[i] = linear(src[i]); + } } diff --git a/src/layers/dense.cu b/src/layers/dense.cu index 37b6dc3..51a2aa0 100644 --- a/src/layers/dense.cu +++ b/src/layers/dense.cu @@ -1,16 +1,25 @@ -#include "dense.cuh" -#include "cuda_helper.cuh" -#include "activations.cuh" -#include -#include #include +#include + #include -#include +#include #include +#include -Layers::Dense::Dense(int inputSize, int outputSize, std::string activation, cublasHandle_t cublasHandle) - : inputSize(inputSize), outputSize(outputSize), cublasHandle(cublasHandle), activation(activation) { +#include "activations.cuh" +#include "cuda_helper.cuh" +#include "dense.cuh" +Layers::Dense::Dense( + int inputSize, + int outputSize, + std::string activation, + cublasHandle_t cublasHandle +) + : inputSize(inputSize), + outputSize(outputSize), + cublasHandle(cublasHandle), + activation(activation) { // Allocate memory for weights and biases weights.resize(outputSize * inputSize); biases.resize(outputSize); @@ -19,10 +28,12 @@ Layers::Dense::Dense(int inputSize, int outputSize, std::string activation, cubl initializeBiases(); d_weights = nullptr; - d_biases = nullptr; + d_biases = nullptr; // Allocate GPU memory for weights and biases - CUDA_CHECK(cudaMalloc((void**)&d_weights, sizeof(float) * inputSize * outputSize)); + CUDA_CHECK( + cudaMalloc((void**)&d_weights, sizeof(float) * inputSize * outputSize) + ); CUDA_CHECK(cudaMalloc((void**)&d_biases, sizeof(float) * outputSize)); toCuda(); @@ -44,30 +55,47 @@ void Layers::Dense::initializeBiases() { void Layers::Dense::forward(const float* d_input, float* d_output) { const float alpha = 1.0f; - const float beta = 1.0f; + const float beta = 1.0f; - CUBLAS_CHECK(cublasSgemv(cublasHandle, CUBLAS_OP_N, inputSize, outputSize, &alpha, d_weights, inputSize, d_input, 1, &beta, d_output, 1)); - CUBLAS_CHECK(cublasSaxpy(cublasHandle, outputSize, &alpha, d_biases, 1, d_output, 1)); + CUBLAS_CHECK(cublasSgemv( + cublasHandle, CUBLAS_OP_N, inputSize, outputSize, &alpha, d_weights, + inputSize, d_input, 1, &beta, d_output, 1 + )); + CUBLAS_CHECK( + cublasSaxpy(cublasHandle, outputSize, &alpha, d_biases, 1, d_output, 1) + ); int threadsPerBlock = 256; - int blocksPerGrid = (outputSize + threadsPerBlock - 1) / threadsPerBlock; + int blocksPerGrid = (outputSize + threadsPerBlock - 1) / threadsPerBlock; if (activation == "sigmoid") { - sigmoid_kernel<<>>(d_output, d_output, outputSize); + sigmoid_kernel<<>>( + d_output, d_output, outputSize + ); } else if (activation == "relu") { - relu_kernel<<>>(d_output, d_output, outputSize); + relu_kernel<<>>( + d_output, d_output, outputSize + ); } else { - linear_kernel<<>>(d_output, d_output, outputSize); + linear_kernel<<>>( + d_output, d_output, outputSize + ); } - } void Layers::Dense::toCuda() { - CUBLAS_CHECK(cublasSetMatrix(outputSize, inputSize, sizeof(float), weights.data(), outputSize, d_weights, outputSize)); - CUBLAS_CHECK(cublasSetVector(biases.size(), sizeof(float), biases.data(), 1, d_biases, 1)); + CUBLAS_CHECK(cublasSetMatrix( + outputSize, inputSize, sizeof(float), weights.data(), outputSize, + d_weights, outputSize + )); + CUBLAS_CHECK(cublasSetVector( + biases.size(), sizeof(float), biases.data(), 1, d_biases, 1 + )); } -void Layers::Dense::setWeights(const std::vector>& weights_input) { +void Layers::Dense::setWeights( + const std::vector>& weights_input +) { int numWeights = inputSize * outputSize; if (weights.size() != numWeights) { @@ -77,7 +105,7 @@ void Layers::Dense::setWeights(const std::vector>& weights_in for (int j = 0; j < inputSize; ++j) { for (int i = 0; i < outputSize; ++i) { - int idx = IDX2C(i, j, outputSize); + int idx = IDX2C(i, j, outputSize); weights[idx] = weights_input[i][j]; } } diff --git a/src/utils/cuda_helper.cu b/src/utils/cuda_helper.cu index 6ba2a9a..72afaae 100644 --- a/src/utils/cuda_helper.cu +++ b/src/utils/cuda_helper.cu @@ -1,8 +1,10 @@ +#include +#include + #include #include + #include "cuda_helper.cuh" -#include -#include cudaDeviceProp initializeCUDA(cublasHandle_t& cublasHandle) { int deviceCount; diff --git a/test/layers/test_dense.cu b/test/layers/test_dense.cu index cecd1ad..7110ad2 100644 --- a/test/layers/test_dense.cu +++ b/test/layers/test_dense.cu @@ -1,14 +1,24 @@ -#include "gtest/gtest.h" #include #include + #include + #include "activations.cuh" #include "dense.cuh" +#include "gtest/gtest.h" #include "test_cublas_fixture.cuh" class DenseLayerTest : public CublasTestFixture { -protected: - Layers::Dense commonTestSetup(int inputSize, int outputSize, std::vector& input, std::vector>& weights, std::vector& biases, float*& d_input, float*& d_output) { + protected: + Layers::Dense commonTestSetup( + int inputSize, + int outputSize, + std::vector& input, + std::vector>& weights, + std::vector& biases, + float*& d_input, + float*& d_output + ) { // Create Dense layer Layers::Dense denseLayer(inputSize, outputSize, "linear", cublasHandle); @@ -24,7 +34,9 @@ protected: EXPECT_EQ(cudaStatus, cudaSuccess); // Copy input to device - cublasStatus = cublasSetVector(input.size(), sizeof(float), input.data(), 1, d_input, 1); + cublasStatus = cublasSetVector( + input.size(), sizeof(float), input.data(), 1, d_input, 1 + ); EXPECT_EQ(cublasStatus, CUBLAS_STATUS_SUCCESS); return denseLayer; @@ -36,28 +48,27 @@ protected: cudaFree(d_output); } - cudaError_t cudaStatus; + cudaError_t cudaStatus; cublasStatus_t cublasStatus; }; TEST_F(DenseLayerTest, Init) { - for (int i = 1; i < 100; ++i) { for (int j = 1; j < 100; ++j) { - - int inputSize = i; + int inputSize = i; int outputSize = j; - // std::cout << "Dense layer: input size = " << inputSize << ", output size = " << outputSize << std::endl; - Layers::Dense denseLayer(inputSize, outputSize, "linear", cublasHandle); - } + // std::cout << "Dense layer: input size = " << inputSize << ", + // output size = " << outputSize << std::endl; + Layers::Dense denseLayer( + inputSize, outputSize, "linear", cublasHandle + ); + } } } TEST_F(DenseLayerTest, setWeights) { - - - int inputSize = 4; + int inputSize = 4; int outputSize = 5; std::vector> weights = { @@ -71,17 +82,17 @@ TEST_F(DenseLayerTest, setWeights) { Layers::Dense denseLayer(inputSize, outputSize, "linear", cublasHandle); denseLayer.setWeights(weights); - } TEST_F(DenseLayerTest, ForwardUnitWeightMatrix) { - - int inputSize = 3; + int inputSize = 3; int outputSize = 3; std::vector input = {1.0f, 2.0f, 3.0f}; - std::vector> weights(inputSize, std::vector(outputSize, 0.0f)); + std::vector> weights( + inputSize, std::vector(outputSize, 0.0f) + ); for (int i = 0; i < inputSize; ++i) { for (int j = 0; j < outputSize; ++j) { if (i == j) { @@ -94,11 +105,15 @@ TEST_F(DenseLayerTest, ForwardUnitWeightMatrix) { float* d_input; float* d_output; - Layers::Dense denseLayer = commonTestSetup(inputSize, outputSize, input, weights, biases, d_input, d_output); + Layers::Dense denseLayer = commonTestSetup( + inputSize, outputSize, input, weights, biases, d_input, d_output + ); denseLayer.forward(d_input, d_output); std::vector output(outputSize); - cublasStatus = cublasGetVector(outputSize, sizeof(float), d_output, 1, output.data(), 1); + cublasStatus = cublasGetVector( + outputSize, sizeof(float), d_output, 1, output.data(), 1 + ); EXPECT_EQ(cublasStatus, CUBLAS_STATUS_SUCCESS); // Check if the output is a zero vector @@ -110,7 +125,7 @@ TEST_F(DenseLayerTest, ForwardUnitWeightMatrix) { } TEST_F(DenseLayerTest, ForwardRandomWeightMatrix) { - int inputSize = 5; + int inputSize = 5; int outputSize = 4; std::vector input = {1.0f, 2.0f, 3.0f, 4.0f, 5.0f}; @@ -120,23 +135,29 @@ TEST_F(DenseLayerTest, ForwardRandomWeightMatrix) { {1.0f, 0.3f, 1.8f, 2.0f, 0.5f}, {0.2f, 1.5f, 0.9f, 0.6f, 0.0f}, {0.8f, 0.4f, 0.1f, 1.1f, 1.7f} - }; + }; std::vector biases = {0.2f, 0.5f, 0.7f, 1.1f}; float* d_input; - float* d_output; + float* d_output; + + Layers::Dense denseLayer = commonTestSetup( + inputSize, outputSize, input, weights, biases, d_input, d_output + ); - Layers::Dense denseLayer = commonTestSetup(inputSize, outputSize, input, weights, biases, d_input, d_output); - denseLayer.forward(d_input, d_output); std::vector output(outputSize); - cublasStatus = cublasGetVector(outputSize, sizeof(float), d_output, 1, output.data(), 1); + cublasStatus = cublasGetVector( + outputSize, sizeof(float), d_output, 1, output.data(), 1 + ); EXPECT_EQ(cublasStatus, CUBLAS_STATUS_SUCCESS); std::vector expectedOutput = {10.4f, 13.0f, 8.9f, 9.3f}; for (int i = 0; i < outputSize; ++i) { - EXPECT_NEAR(output[i], expectedOutput[i], 1e-4); // Allow small tolerance for floating-point comparison + EXPECT_NEAR( + output[i], expectedOutput[i], 1e-4 + ); // Allow small tolerance for floating-point comparison } commonTestTeardown(d_input, d_output); diff --git a/test/test_utils/test_cublas_fixture.cu b/test/test_utils/test_cublas_fixture.cu index 7403cf5..a1b8708 100644 --- a/test/test_utils/test_cublas_fixture.cu +++ b/test/test_utils/test_cublas_fixture.cu @@ -1,5 +1,5 @@ -#include "gtest/gtest.h" #include "cublas_v2.h" +#include "gtest/gtest.h" #include "test_cublas_fixture.cuh" cublasHandle_t CublasTestFixture::cublasHandle; diff --git a/test/test_utils/test_cublas_fixture.cuh b/test/test_utils/test_cublas_fixture.cuh index 9e927df..51ed7a3 100644 --- a/test/test_utils/test_cublas_fixture.cuh +++ b/test/test_utils/test_cublas_fixture.cuh @@ -1,8 +1,8 @@ -#include "gtest/gtest.h" #include "cublas_v2.h" +#include "gtest/gtest.h" class CublasTestFixture : public ::testing::Test { -protected: + protected: static cublasHandle_t cublasHandle; static void SetUpTestSuite(); diff --git a/test/utils/test_functions.cu b/test/utils/test_functions.cu index 8431916..b711e17 100644 --- a/test/utils/test_functions.cu +++ b/test/utils/test_functions.cu @@ -1,16 +1,16 @@ -#include "gtest/gtest.h" #include #include + #include + #include "functions.cuh" +#include "gtest/gtest.h" #include "test_cublas_fixture.cuh" class FunctionsTest : public CublasTestFixture { -protected: - cudaError_t cudaStatus; + protected: + cudaError_t cudaStatus; cublasStatus_t cublasStatus; }; -TEST_F(FunctionsTest, sigmoid) { - -} \ No newline at end of file +TEST_F(FunctionsTest, sigmoid) {} \ No newline at end of file