Skip to content

Commit c55d0d8

Browse files
committed
Fix cuda resize kernel if image_dim > neuron_dim
Related to #33
1 parent 1e9e8ca commit c55d0d8

5 files changed

Lines changed: 47 additions & 1 deletion

File tree

src/CudaLib/resize_kernel.h

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -22,7 +22,7 @@ void resize_kernel(T *dst, T const *src, uint32_t dst_dim, uint32_t src_dim, uin
2222

2323
uint32_t src_margin = 0, dst_margin = 0;
2424
if (src_dim < dst_dim) dst_margin = (dst_dim - src_dim) * 0.5;
25-
else if (src_dim > dst_dim) src_margin = (src_dim - dst_dim) / 0.5;
25+
else if (src_dim > dst_dim) src_margin = (src_dim - dst_dim) * 0.5;
2626

2727
dst[(i + dst_margin) * dst_dim + (j + dst_margin)] =
2828
src[(i + src_margin) * src_dim + (j + src_margin)];

src/SelfOrganizingMapLib/Trainer.h

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -133,6 +133,7 @@ class Trainer<SOMLayout, DataLayout, T, false> : public TrainerBase<SOMLayout, D
133133
this->use_flip, this->interpolation, neuron_dim);
134134

135135
#ifdef PRINT_DEBUG
136+
std::cout << "spatial_transformed_images" << std::endl;
136137
for (auto&& e : spatial_transformed_images) std::cout << e << " ";
137138
std::cout << std::endl;
138139
#endif
@@ -237,6 +238,7 @@ class Trainer<SOMLayout, DataLayout, T, true> : public TrainerBase<SOMLayout, Da
237238
data.get_dimension()[0], neuron_dim, this->use_flip, this->interpolation, d_cos_alpha, d_sin_alpha);
238239

239240
#ifdef PRINT_DEBUG
241+
std::cout << "spatial_transformed_images" << std::endl;
240242
thrust::host_vector<T> spatial_transformed_images = d_spatial_transformed_images;
241243
for (auto&& e : spatial_transformed_images) std::cout << e << " ";
242244
std::cout << std::endl;

test/CudaTest/CMakeLists.txt

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -11,6 +11,7 @@ cuda_add_executable(
1111
compare_trainer_mixed.cu
1212
mapper.cu
1313
mixed_precision.cu
14+
resize.cu
1415
rotate_90_degrees_list.cu
1516
update_neurons.cu
1617
)

test/CudaTest/compare_trainer_cpu.cu

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -88,4 +88,8 @@ INSTANTIATE_TEST_CASE_P(TrainerCompareTest_all, compare_trainer_cpu,
8888
,TrainerCompareTestData(2, 64, 100, 100, 360, true)
8989
,TrainerCompareTestData(2, 64, 100, 45, 360, true)
9090
,TrainerCompareTestData(2, 124, 91, 64, 360, true)
91+
,TrainerCompareTestData(2, 4, 2, 2, 1, false)
92+
,TrainerCompareTestData(2, 4, 2, 2, 8, false)
93+
,TrainerCompareTestData(2, 4, 2, 2, 1, true)
94+
,TrainerCompareTestData(2, 4, 2, 2, 8, true)
9195
));

test/CudaTest/resize.cu

Lines changed: 39 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,39 @@
1+
/**
2+
* @file CudaTest/resize.cu
3+
* @date Nov 14, 2018
4+
* @author Bernd Doser, HITS gGmbH
5+
*/
6+
7+
#include <gtest/gtest.h>
8+
#include <thrust/device_vector.h>
9+
#include <thrust/host_vector.h>
10+
#include <vector>
11+
12+
#include "CudaLib/resize_kernel.h"
13+
14+
using namespace pink;
15+
16+
TEST(ResizeTest, small)
17+
{
18+
uint32_t neuron_dim = 2;
19+
uint32_t image_dim = 4;
20+
uint32_t min_dim = 2;
21+
22+
std::vector<uint32_t> image{ 0, 1, 2, 3,
23+
4, 5, 6, 7,
24+
8, 9, 10, 11,
25+
12, 13, 14, 15};
26+
27+
thrust::device_vector<uint32_t> d_image = image;
28+
thrust::device_vector<uint32_t> d_resized_image(4);
29+
30+
dim3 dim_block(32, 32);
31+
dim3 dim_grid(1, 1);
32+
33+
resize_kernel<<<dim_grid, dim_block>>>(thrust::raw_pointer_cast(&d_resized_image[0]),
34+
thrust::raw_pointer_cast(&d_image[0]), neuron_dim, image_dim, min_dim);
35+
36+
thrust::host_vector<uint32_t> result = d_resized_image;
37+
thrust::host_vector<uint32_t> expected = std::vector<uint32_t>{5, 6, 9, 10};
38+
EXPECT_EQ(expected, result);
39+
}

0 commit comments

Comments
 (0)