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

Combine padding and conv kernel

Browse Source
This commit is contained in:
Matúš Námešný
2024-03-18 19:53:40 +01:00
parent e6d3757312
commit 6cf604423a
6 changed files with 14 additions and 209 deletions
Download Patch File Download Diff File Expand all files Collapse all files

90
src/kernels/convolution.cu
View File

@@ -2,83 +2,6 @@
#include "convolution.cuh"
/*
Pads matrix width x height x n_channels to width + 2 * padding x height + 2 *
padding x n_channels Matrix is represented as a pointer to a vector
For example:
w = 2
h = 3
n = 2
p = 1
Channel 0:
0 1
2 3
4 5
Channel 1:
6 7
8 9
10 11
Is represented as:
0 1 2 3 4 5 6 7 8 9 10 11
Padded result (as a continuous vector):
0.0f, 0.0f, 0.0f, 0.0f,
0.0f, 0.0f, 1.0f, 0.0f,
0.0f, 2.0f, 3.0f, 0.0f,
0.0f, 4.0f, 5.0f, 0.0f,
0.0f, 0.0f, 0.0f, 0.0f,
0.0f, 0.0f, 0.0f, 0.0f,
0.0f, 6.0f, 7.0f, 0.0f,
0.0f, 8.0f, 9.0f, 0.0f,
9.0f, 10.0f, 11.0f, 0.0f,
0.0f, 0.0f, 0.0f, 0.0f
Args:
d_input: Pointer to input vector representing matrix
d_padded: Pointer to output vector representing padded matrix (needs to be
pre-allocated)
w: Width of input matrix
h: Height of input matrix
n: Number of channels in input matrix
p: Padding
*/
__global__ void CUDANet::Kernels::padding(
const float* __restrict__ d_input,
float* __restrict__ d_padded,
const unsigned int w,
const unsigned int h,
const unsigned int n,
const unsigned int p
) {
int tid = blockDim.x * blockIdx.x + threadIdx.x;
if (tid >= (w + 2 * p) * (h + 2 * p) * n) {
return;
}
int idx = tid;
// unravel index into padded matrix
int i_n = idx / ((w + 2 * p) * (h + 2 * p));
int i_h = idx % ((w + 2 * p) * (h + 2 * p)) / (w + 2 * p);
int i_w = idx % (w + 2 * p);
// if i is in the padding region
if (i_w < p || i_w >= (w + p) || i_h < p || i_h >= (h + p)) {
d_padded[tid] = 0.0f;
} else {
// Get index into input vector
int i_orig = i_n * w * h + (i_h - p) * w + (i_w - p);
d_padded[tid] = d_input[i_orig];
}
}
__global__ void CUDANet::Kernels::convolution(
const float* __restrict__ d_input,
const float* __restrict__ d_kernel,
@@ -108,12 +31,21 @@ __global__ void CUDANet::Kernels::convolution(
for (int c = 0; c < nChannels; c++) {
for (int k = 0; k < kernelSize; k++) {
for (int l = 0; l < kernelSize; l++) {
// if i, j is in the padding region
if (i * stride + k < paddingSize ||
i * stride + k >= (inputSize + paddingSize) ||
j * stride + l < paddingSize ||
j * stride + l >= (inputSize + paddingSize)) {
continue;
}
int kernelIndex = f * kernelSize * kernelSize * nChannels +
c * kernelSize * kernelSize + k * kernelSize +
l;
int inputIndex = c * inputSize * inputSize +
(i * stride + k) * inputSize +
(j * stride + l);
(i * stride + k - paddingSize) * inputSize +
(j * stride + l - paddingSize);
sum += d_kernel[kernelIndex] * d_input[inputIndex];
}

20
src/layers/conv2d.cu
View File

@@ -23,6 +23,7 @@ Layers::Conv2d::Conv2d(
kernelSize(kernelSize),
stride(stride),
numFilters(numFilters) {
switch (padding) {
case SAME:
outputSize = inputSize;
@@ -64,12 +65,6 @@ Layers::Conv2d::Conv2d(
(void**)&d_biases, sizeof(float) * outputSize * outputSize * numFilters
));
d_padded = nullptr;
CUDA_CHECK(cudaMalloc(
(void**)&d_padded, sizeof(float) * (inputSize + 2 * paddingSize) *
(inputSize + 2 * paddingSize) * inputChannels
));
toCuda();
}
@@ -77,7 +72,6 @@ Layers::Conv2d::~Conv2d() {
cudaFree(d_output);
cudaFree(d_weights);
cudaFree(d_biases);
cudaFree(d_padded);
}
void Layers::Conv2d::initializeWeights() {
@@ -113,18 +107,10 @@ void Layers::Conv2d::toCuda() {
}
float* Layers::Conv2d::forward(const float* d_input) {
// Pad input
int THREADS_PER_BLOCK = (inputSize + 2 * paddingSize) *
(inputSize + 2 * paddingSize) * inputChannels;
Kernels::padding<<<1, THREADS_PER_BLOCK>>>(
d_input, d_padded, inputSize, inputSize, inputChannels, paddingSize
);
// Convolve
THREADS_PER_BLOCK = outputSize * outputSize * numFilters;
int THREADS_PER_BLOCK = outputSize * outputSize * numFilters;
Kernels::convolution<<<1, THREADS_PER_BLOCK>>>(
d_padded, d_weights, d_output, inputSize + 2 * paddingSize, inputChannels, paddingSize,
d_input, d_weights, d_output, inputSize, inputChannels, paddingSize,
kernelSize, stride, numFilters, outputSize
);