1 // Copyright (c) 2018, ETH Zurich and UNC Chapel Hill.
2 // All rights reserved.
3 //
4 // Redistribution and use in source and binary forms, with or without
5 // modification, are permitted provided that the following conditions are met:
6 //
7 // * Redistributions of source code must retain the above copyright
8 // notice, this list of conditions and the following disclaimer.
9 //
10 // * Redistributions in binary form must reproduce the above copyright
11 // notice, this list of conditions and the following disclaimer in the
12 // documentation and/or other materials provided with the distribution.
13 //
14 // * Neither the name of ETH Zurich and UNC Chapel Hill nor the names of
15 // its contributors may be used to endorse or promote products derived
16 // from this software without specific prior written permission.
17 //
18 // THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
19 // AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
20 // IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
21 // ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDERS OR CONTRIBUTORS BE
22 // LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
23 // CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
24 // SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
25 // INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
26 // CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
27 // ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
28 // POSSIBILITY OF SUCH DAMAGE.
29 //
30 // Author: Johannes L. Schoenberger (jsch-at-demuc-dot-de)
31
32 #ifndef COLMAP_SRC_MVS_CUDA_ROTATE_H_
33 #define COLMAP_SRC_MVS_CUDA_ROTATE_H_
34
35 #include <cuda_runtime.h>
36
37 namespace colmap {
38 namespace mvs {
39
40 // Rotate the input matrix by 90 degrees in counter-clockwise direction.
41 template <typename T>
42 void CudaRotate(const T* input, T* output, const int width, const int height,
43 const int pitch_input, const int pitch_output);
44
45 ////////////////////////////////////////////////////////////////////////////////
46 // Implementation
47 ////////////////////////////////////////////////////////////////////////////////
48
49 #ifdef __CUDACC__
50
51 #define TILE_DIM_ROTATE 32
52
53 namespace internal {
54
55 template <typename T>
CudaRotateKernel(T * output_data,const T * input_data,const int width,const int height,const int input_pitch,const int output_pitch)56 __global__ void CudaRotateKernel(T* output_data, const T* input_data,
57 const int width, const int height,
58 const int input_pitch,
59 const int output_pitch) {
60 int input_x = blockDim.x * blockIdx.x + threadIdx.x;
61 int input_y = blockDim.y * blockIdx.y + threadIdx.y;
62
63 if (input_x >= width || input_y >= height) {
64 return;
65 }
66
67 int output_x = input_y;
68 int output_y = width - 1 - input_x;
69
70 *((T*)((char*)output_data + output_y * output_pitch) + output_x) =
71 *((T*)((char*)input_data + input_y * input_pitch) + input_x);
72 }
73
74 } // namespace internal
75
76 template <typename T>
CudaRotate(const T * input,T * output,const int width,const int height,const int pitch_input,const int pitch_output)77 void CudaRotate(const T* input, T* output, const int width, const int height,
78 const int pitch_input, const int pitch_output) {
79 dim3 block_dim(TILE_DIM_ROTATE, 1, 1);
80 dim3 grid_dim;
81 grid_dim.x = (width - 1) / TILE_DIM_ROTATE + 1;
82 grid_dim.y = height;
83
84 internal::CudaRotateKernel<<<grid_dim, block_dim>>>(
85 output, input, width, height, pitch_input, pitch_output);
86 }
87
88 #undef TILE_DIM_ROTATE
89
90 #endif // __CUDACC__
91
92 } // namespace mvs
93 } // namespace colmap
94
95 #endif // COLMAP_SRC_MVS_CUDA_ROTATE_H_
96