Mathis/CUDANet
1
0
Fork 0
mirror of https://github.com/lordmathis/CUDANet.git synced 2025-12-22 06:14:22 +00:00
Code Issues Packages Projects Releases Wiki Activity

Add avgPool2d implementation

Browse Source
This commit is contained in:
Matúš Námešný
2025-11-21 19:39:30 +01:00
parent c83e1f0c45
commit 5679dc0a50
7 changed files with 102 additions and 67 deletions
Download Patch File Download Diff File Expand all files Collapse all files

10
include/backend.hpp
View File

@@ -62,6 +62,16 @@ class Backend {
CUDANet::Shape padding_shape,
CUDANet::Shape output_shape
) = 0;
virtual CUDANet::Tensor& avgPool2d(
const CUDANet::Tensor& input,
CUDANet::Tensor& output,
CUDANet::Shape input_shape,
CUDANet::Shape pool_shape,
CUDANet::Shape stride_shape,
CUDANet::Shape padding_shape,
CUDANet::Shape output_shape
) = 0;
};
} // namespace CUDANet

12
include/backend/cuda.cuh
View File

@@ -49,7 +49,7 @@ class CUDA : public Backend {
const CUDANet::Shape out_shape
) override;
CUDANet::Tensor& CUDA::maxPool2d(
CUDANet::Tensor& maxPool2d(
const CUDANet::Tensor& input,
CUDANet::Tensor& output,
CUDANet::Shape input_shape,
@@ -58,6 +58,16 @@ class CUDA : public Backend {
CUDANet::Shape padding_shape,
CUDANet::Shape output_shape
) override;
CUDANet::Tensor& avgPool2d(
const CUDANet::Tensor& input,
CUDANet::Tensor& output,
CUDANet::Shape input_shape,
CUDANet::Shape pool_shape,
CUDANet::Shape stride_shape,
CUDANet::Shape padding_shape,
CUDANet::Shape output_shape
) = 0;
};
} // namespace CUDANet::Backend

0
include/kernels/pooling.cuh → include/kernels/pool.cuh
View File

2
src/backends/cuda/kernels/pooling.cu → src/backends/cuda/kernels/pool.cu
View File

@@ -1,6 +1,6 @@
#include "cuda_helper.cuh"
#include "layer.hpp"
#include "pooling.cuh"
#include "pool.cuh"
using namespace CUDANet;

26
src/backends/cuda/layer_ops.cu
View File

@@ -135,5 +135,31 @@ CUDANet::Tensor& CUDA::maxPool2d(
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
return output;
}
CUDANet::Tensor& CUDA::avgPool2d(
const CUDANet::Tensor& input,
CUDANet::Tensor& output,
CUDANet::Shape input_shape,
CUDANet::Shape pool_shape,
CUDANet::Shape stride_shape,
CUDANet::Shape padding_shape,
CUDANet::Shape output_shape
) {
dim3 block(8, 8, 8);
dim3 grid(
(output_shape[0] + block.x - 1) / block.x,
(output_shape[1] + block.y - 1) / block.y,
(output_shape[2] + block.z - 1) / block.z
);
Kernels::avg_pool<<<grid, block>>>(
input.data<float>(), output.data<float>(), input_shape, output_shape, pool_shape,
stride_shape, padding_shape
);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
return output;
}

45
src/backends/cuda/layers/avg_pooling.cu
View File

@@ -1,45 +0,0 @@
#include "avg_pooling.hpp"
#include "cuda_helper.cuh"
#include "pooling.cuh"
using namespace CUDANet::Layers;
void AvgPooling2d::initCUDA() {
d_output = nullptr;
CUDA_CHECK(cudaMalloc(
(void**)&d_output,
sizeof(float) * outputSize.first * outputSize.second * nChannels
));
}
void AvgPooling2d::delCUDA() {
cudaFree(d_output);
}
float* AvgPooling2d::forwardCUDA(const float* d_input) {
dim3 block(8, 8, 8);
dim3 grid(
(outputSize.first + block.x - 1) / block.x,
(outputSize.second + block.y - 1) / block.y,
(nChannels + block.z - 1) / block.z
);
Kernels::avg_pooling<<<grid, block>>>(
d_input, d_output, inputSize, outputSize, nChannels, poolingSize,
stride, padding
);
CUDA_CHECK(cudaGetLastError());
activation->activate(d_output);
CUDA_CHECK(cudaDeviceSynchronize());
return d_output;
}
void AdaptiveAvgPooling2d::initCUDA() {
cudaFree(d_output);
cudaMalloc(
(void**)&d_output,
sizeof(float) * outputSize.first * outputSize.second * nChannels
);
}

74
src/layers/avg_pooling.cpp
View File

@@ -49,23 +49,43 @@ AvgPool2d::AvgPool2d(
AvgPool2d::~AvgPool2d() {}
CUDANet::Tensor& AvgPool2d::forward(CUDANet::Tensor& input);
CUDANet::Tensor& AvgPool2d::forward(CUDANet::Tensor& input) {
output.zero();
backend->avgPool2d(
input,
output,
in_shape,
pool_shape,
stride_shape,
padding_shape,
out_shape
);
return output;
}
CUDANet::Shape AvgPool2d::input_shape();
CUDANet::Shape AvgPool2d::input_shape() {
return in_shape;
}
CUDANet::Shape AvgPool2d::output_shape();
CUDANet::Shape AvgPool2d::output_shape() {
return out_shape;
}
size_t AvgPool2d::input_size();
size_t AvgPool2d::input_size() {
return sizeof(float) * in_shape[0] * in_shape[1] * in_shape[2];
}
size_t AvgPool2d::output_size();
size_t AvgPool2d::output_size() {
return sizeof(float) * out_shape[0] * out_shape[1] * out_shape[3];
}
void AvgPool2d::set_weights(void* input);
void AvgPool2d::set_weights(void* input) {}
CUDANet::Tensor& AvgPool2d::get_weights();
CUDANet::Tensor& AvgPool2d::get_weights() {}
void AvgPool2d::set_biases(void* input);
void AvgPool2d::set_biases(void* input) {}
CUDANet::Tensor& AvgPool2d::get_biases();
CUDANet::Tensor& AvgPool2d::get_biases() {}
AdaptiveAvgPool2d::AdaptiveAvgPool2d(
@@ -73,15 +93,29 @@ AdaptiveAvgPool2d::AdaptiveAvgPool2d(
CUDANet::Shape output_shape,
CUDANet::Backend *backend
)
: AvgPool2d(input_shape, {1, 1}, {1, 1}, {0, 0}, backend) {
stride_shape = {
input_shape[0] / output_shape[0],
input_shape[1] / output_shape[1]
};
pool_shape = {
input_shape[0] - (output_shape[0] - 1) * stride_shape[0],
input_shape[1] - (output_shape[1] - 1) * stride_shape[1]
};
padding_shape = {(pool_shape[0] - 1) / 2, (pool_shape[1] - 1) / 2};
: AvgPool2d(
input_shape,
// pool_shape
{
input_shape[0] - (output_shape[0] - 1) * (input_shape[0] / output_shape[0]),
input_shape[1] - (output_shape[1] - 1) * (input_shape[1] / output_shape[1])
},
// stride_shape
{
input_shape[0] / output_shape[0],
input_shape[1] / output_shape[1]
},
// padding_shape
{
(input_shape[0] - (output_shape[0] - 1) * (input_shape[0] / output_shape[0]) - 1) / 2,
(input_shape[1] - (output_shape[1] - 1) * (input_shape[1] / output_shape[1]) - 1) / 2
},
backend
) {
out_shape = output_shape;
}
output = CUDANet::Tensor(
Shape{out_shape[0] * out_shape[1] * out_shape[2]},
CUDANet::DType::FLOAT32, backend
);
}