Skip to content

Commit e4e114a

Browse files
griwodzCarsten Griwodzsimogasp
authored
[debug] Add const qualifier to ops that copy between host and device
* make planes const in logging * write functions use const planes * add const to memcpyFromDevice and memcpyToHost * missing const method in implementation --------- Co-authored-by: Carsten Griwodz <[email protected]> Co-authored-by: Simone Gasparini <[email protected]>
1 parent cafb65e commit e4e114a

File tree

5 files changed

+45
-44
lines changed

5 files changed

+45
-44
lines changed

‎src/popsift/common/plane_2d.cu‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -164,7 +164,7 @@ __host__
164164
void PlaneBase::memcpyToHost( void* dst, int dst_pitch,
165165
void* src, int src_pitch,
166166
short cols, short rows,
167-
int elemSize )
167+
int elemSize ) const
168168
{
169169
assert( dst != 0 );
170170
assert( dst_pitch != 0 );
@@ -185,7 +185,7 @@ void PlaneBase::memcpyToHost( void* dst, int dst_pitch,
185185
void* src, int src_pitch,
186186
short cols, short rows,
187187
int elemSize,
188-
cudaStream_t stream )
188+
cudaStream_t stream ) const
189189
{
190190
assert( dst != 0 );
191191
assert( dst_pitch != 0 );

‎src/popsift/common/plane_2d.h‎

Lines changed: 28 additions & 27 deletions
Original file line numberDiff line numberDiff line change
@@ -65,10 +65,10 @@ struct PlaneBase
6565
void memcpyToDevice( void* dst, int dst_pitch, void* src, int src_pitch, short cols, short rows, int elemSize, cudaStream_t stream );
6666

6767
__host__
68-
void memcpyToHost( void* dst, int dst_pitch, void* src, int src_pitch, short cols, short rows, int elemSize );
68+
void memcpyToHost( void* dst, int dst_pitch, void* src, int src_pitch, short cols, short rows, int elemSize ) const;
6969

7070
__host__
71-
void memcpyToHost( void* dst, int dst_pitch, void* src, int src_pitch, short cols, short rows, int elemSize, cudaStream_t stream );
71+
void memcpyToHost( void* dst, int dst_pitch, void* src, int src_pitch, short cols, short rows, int elemSize, cudaStream_t stream ) const;
7272

7373
#ifdef PLANE2D_CUDA_OP_DEBUG
7474
__host__
@@ -130,17 +130,18 @@ template <typename T> struct PitchPlane2D : public PlaneT<T>
130130

131131
/** cuda memcpy from parameter (plane allocated on device) to
132132
* this (plane allocated on host) */
133-
__host__ inline void memcpyFromDevice( PitchPlane2D<T>& devPlane,
133+
__host__ inline void memcpyFromDevice( const PitchPlane2D<T>& devPlane,
134134
short cols, short rows );
135-
__host__ inline void memcpyFromDevice( PitchPlane2D<T>& devPlane,
136-
short cols, short rows, cudaStream_t stream );
135+
__host__ inline void memcpyFromDevice( const PitchPlane2D<T>& devPlane,
136+
short cols, short rows,
137+
cudaStream_t stream );
137138

138139
/** cuda memcpy from this (plane allocated on device) to
139140
* parameter (plane allocated on host) */
140141
__host__ inline void memcpyToHost( PitchPlane2D<T>& hostPlane,
141-
short cols, short rows );
142+
short cols, short rows ) const;
142143
__host__ inline void memcpyToHost( PitchPlane2D<T>& hostPlane,
143-
short cols, short rows, cudaStream_t stream );
144+
short cols, short rows, cudaStream_t stream ) const;
144145

145146
__host__ __device__ inline const T* ptr( int y ) const {
146147
return (const T*)( (const char*)this->data + y * _pitchInBytes );
@@ -219,7 +220,7 @@ inline void PitchPlane2D<T>::memcpyFromHost( PitchPlane2D<T>& hostPlane, short c
219220

220221
template <typename T>
221222
__host__
222-
inline void PitchPlane2D<T>::memcpyFromDevice( PitchPlane2D<T>& devPlane, short cols, short rows )
223+
inline void PitchPlane2D<T>::memcpyFromDevice( const PitchPlane2D<T>& devPlane, short cols, short rows )
223224
{
224225
PlaneBase::memcpyToHost( this->data, this->_pitchInBytes,
225226
devPlane.data, devPlane._pitchInBytes,
@@ -229,7 +230,7 @@ inline void PitchPlane2D<T>::memcpyFromDevice( PitchPlane2D<T>& devPlane, short
229230

230231
template <typename T>
231232
__host__
232-
inline void PitchPlane2D<T>::memcpyFromDevice( PitchPlane2D<T>& devPlane, short cols, short rows, cudaStream_t stream )
233+
inline void PitchPlane2D<T>::memcpyFromDevice( const PitchPlane2D<T>& devPlane, short cols, short rows, cudaStream_t stream )
233234
{
234235
PlaneBase::memcpyToHost( this->data, this->_pitchInBytes,
235236
devPlane.data, devPlane._pitchInBytes,
@@ -240,14 +241,14 @@ inline void PitchPlane2D<T>::memcpyFromDevice( PitchPlane2D<T>& devPlane, short
240241

241242
template <typename T>
242243
__host__
243-
inline void PitchPlane2D<T>::memcpyToHost( PitchPlane2D<T>& hostPlane, short cols, short rows )
244+
inline void PitchPlane2D<T>::memcpyToHost( PitchPlane2D<T>& hostPlane, short cols, short rows ) const
244245
{
245246
hostPlane.memcpyFromDevice( *this, cols, rows );
246247
}
247248

248249
template <typename T>
249250
__host__
250-
inline void PitchPlane2D<T>::memcpyToHost( PitchPlane2D<T>& hostPlane, short cols, short rows, cudaStream_t stream )
251+
inline void PitchPlane2D<T>::memcpyToHost( PitchPlane2D<T>& hostPlane, short cols, short rows, cudaStream_t stream ) const
251252
{
252253
hostPlane.memcpyFromDevice( *this, cols, rows, stream );
253254
}
@@ -319,17 +320,17 @@ template <typename T> class Plane2D : public PitchPlane2D<T>
319320

320321
/** cuda memcpy from parameter (plane allocated on device) to
321322
* this (plane allocated on host) */
322-
__host__ inline void memcpyFromDevice( Plane2D<T>& devPlane );
323-
__host__ inline void memcpyFromDevice( PitchPlane2D<T>& devPlane );
324-
__host__ inline void memcpyFromDevice( Plane2D<T>& devPlane, cudaStream_t stream );
325-
__host__ inline void memcpyFromDevice( PitchPlane2D<T>& devPlane, cudaStream_t stream );
323+
__host__ inline void memcpyFromDevice( const Plane2D<T>& devPlane );
324+
__host__ inline void memcpyFromDevice( const PitchPlane2D<T>& devPlane );
325+
__host__ inline void memcpyFromDevice( const Plane2D<T>& devPlane, cudaStream_t stream );
326+
__host__ inline void memcpyFromDevice( const PitchPlane2D<T>& devPlane, cudaStream_t stream );
326327

327328
/** cuda memcpy from this (plane allocated on device) to
328329
* parameter (plane allocated on host) */
329-
__host__ inline void memcpyToHost( Plane2D<T>& hostPlane );
330-
__host__ inline void memcpyToHost( Plane2D<T>& hostPlane, cudaStream_t stream );
331-
__host__ inline void memcpyToHost( PitchPlane2D<T>& hostPlane );
332-
__host__ inline void memcpyToHost( PitchPlane2D<T>& hostPlane, cudaStream_t stream );
330+
__host__ inline void memcpyToHost( Plane2D<T>& hostPlane ) const;
331+
__host__ inline void memcpyToHost( Plane2D<T>& hostPlane, cudaStream_t stream ) const;
332+
__host__ inline void memcpyToHost( PitchPlane2D<T>& hostPlane ) const;
333+
__host__ inline void memcpyToHost( PitchPlane2D<T>& hostPlane, cudaStream_t stream ) const;
333334

334335
__host__ __device__
335336
inline short getCols( ) const { return _cols; }
@@ -460,7 +461,7 @@ inline void Plane2D<T>::memcpyFromHost( PitchPlane2D<T>& hostPlane, cudaStream_t
460461

461462
template <typename T>
462463
__host__
463-
inline void Plane2D<T>::memcpyFromDevice( Plane2D<T>& devPlane )
464+
inline void Plane2D<T>::memcpyFromDevice( const Plane2D<T>& devPlane )
464465
{
465466
assert( devPlane._cols == this->_cols );
466467
assert( devPlane._rows == this->_rows );
@@ -469,14 +470,14 @@ inline void Plane2D<T>::memcpyFromDevice( Plane2D<T>& devPlane )
469470

470471
template <typename T>
471472
__host__
472-
inline void Plane2D<T>::memcpyFromDevice( PitchPlane2D<T>& devPlane )
473+
inline void Plane2D<T>::memcpyFromDevice( const PitchPlane2D<T>& devPlane )
473474
{
474475
PitchPlane2D<T>::memcpyFromDevice( devPlane, this->_cols, this->_rows );
475476
}
476477

477478
template <typename T>
478479
__host__
479-
inline void Plane2D<T>::memcpyFromDevice( Plane2D<T>& devPlane, cudaStream_t stream )
480+
inline void Plane2D<T>::memcpyFromDevice( const Plane2D<T>& devPlane, cudaStream_t stream )
480481
{
481482
assert( devPlane._cols == this->_cols );
482483
assert( devPlane._rows == this->_rows );
@@ -485,35 +486,35 @@ inline void Plane2D<T>::memcpyFromDevice( Plane2D<T>& devPlane, cudaStream_t str
485486

486487
template <typename T>
487488
__host__
488-
inline void Plane2D<T>::memcpyFromDevice( PitchPlane2D<T>& devPlane, cudaStream_t stream )
489+
inline void Plane2D<T>::memcpyFromDevice( const PitchPlane2D<T>& devPlane, cudaStream_t stream )
489490
{
490491
PitchPlane2D<T>::memcpyFromDevice( devPlane, this->_cols, this->_rows, stream );
491492
}
492493

493494
template <typename T>
494495
__host__
495-
inline void Plane2D<T>::memcpyToHost( Plane2D<T>& hostPlane )
496+
inline void Plane2D<T>::memcpyToHost( Plane2D<T>& hostPlane ) const
496497
{
497498
hostPlane.memcpyFromDevice( *this );
498499
}
499500

500501
template <typename T>
501502
__host__
502-
inline void Plane2D<T>::memcpyToHost( PitchPlane2D<T>& hostPlane )
503+
inline void Plane2D<T>::memcpyToHost( PitchPlane2D<T>& hostPlane ) const
503504
{
504505
hostPlane.memcpyFromDevice( *this, this->_cols, this->_rows );
505506
}
506507

507508
template <typename T>
508509
__host__
509-
inline void Plane2D<T>::memcpyToHost( Plane2D<T>& hostPlane, cudaStream_t stream )
510+
inline void Plane2D<T>::memcpyToHost( Plane2D<T>& hostPlane, cudaStream_t stream ) const
510511
{
511512
hostPlane.memcpyFromDevice( *this, stream );
512513
}
513514

514515
template <typename T>
515516
__host__
516-
inline void Plane2D<T>::memcpyToHost( PitchPlane2D<T>& hostPlane, cudaStream_t stream )
517+
inline void Plane2D<T>::memcpyToHost( PitchPlane2D<T>& hostPlane, cudaStream_t stream ) const
517518
{
518519
hostPlane.memcpyFromDevice( *this, this->_cols, this->_rows, stream );
519520
}

‎src/popsift/common/write_plane_2d.cu‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -17,7 +17,7 @@ using namespace std;
1717
namespace popsift {
1818

1919
__host__
20-
void write_plane2D( const char* filename, bool onDevice, Plane2D_float& f )
20+
void write_plane2D( const char* filename, bool onDevice, const Plane2D_float& f )
2121
{
2222
if( onDevice ) {
2323
// cerr << __FILE__ << ":" << __LINE__ << ": copying from device" << endl;
@@ -32,7 +32,7 @@ void write_plane2D( const char* filename, bool onDevice, Plane2D_float& f )
3232
}
3333

3434
__host__
35-
void write_plane2Dunscaled( const char* filename, bool onDevice, Plane2D_float& f, int offset )
35+
void write_plane2Dunscaled( const char* filename, bool onDevice, const Plane2D_float& f, int offset )
3636
{
3737
if( onDevice ) {
3838
// cerr << __FILE__ << ":" << __LINE__ << ": copying from device" << endl;
@@ -47,7 +47,7 @@ void write_plane2Dunscaled( const char* filename, bool onDevice, Plane2D_float&
4747
}
4848

4949
__host__
50-
void write_plane2D( const char* filename, Plane2D_float& f )
50+
void write_plane2D( const char* filename, const Plane2D_float& f )
5151
{
5252
// cerr << "Enter " << __FUNCTION__ << endl;
5353

@@ -107,7 +107,7 @@ void write_plane2D( const char* filename, Plane2D_float& f )
107107
}
108108

109109
__host__
110-
void write_plane2Dunscaled( const char* filename, Plane2D_float& f, int offset )
110+
void write_plane2Dunscaled( const char* filename, const Plane2D_float& f, int offset )
111111
{
112112
int rows = f.getRows();
113113
int cols = f.getCols();
@@ -139,7 +139,7 @@ void write_plane2Dunscaled( const char* filename, Plane2D_float& f, int offset )
139139
}
140140

141141
__host__
142-
void dump_plane2Dfloat( const char* filename, bool onDevice, Plane2D_float& f )
142+
void dump_plane2Dfloat( const char* filename, bool onDevice, const Plane2D_float& f )
143143
{
144144
if( onDevice ) {
145145
// cerr << __FILE__ << ":" << __LINE__ << ": copying from device" << endl;
@@ -154,7 +154,7 @@ void dump_plane2Dfloat( const char* filename, bool onDevice, Plane2D_float& f )
154154
}
155155

156156
__host__
157-
void dump_plane2Dfloat( const char* filename, Plane2D_float& f )
157+
void dump_plane2Dfloat( const char* filename, const Plane2D_float& f )
158158
{
159159
int rows = f.getRows();
160160
int cols = f.getCols();

‎src/popsift/common/write_plane_2d.h‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -11,13 +11,13 @@
1111

1212
namespace popsift {
1313

14-
void write_plane2D( const char* filename, bool onDevice, Plane2D_float& f );
15-
void write_plane2D( const char* filename, Plane2D_float& f );
14+
void write_plane2D( const char* filename, bool onDevice, const Plane2D_float& f );
15+
void write_plane2D( const char* filename, const Plane2D_float& f );
1616

17-
void write_plane2Dunscaled( const char* filename, bool onDevice, Plane2D_float& f, int offset=0 );
18-
void write_plane2Dunscaled( const char* filename, Plane2D_float& f, int offset=0 );
17+
void write_plane2Dunscaled( const char* filename, bool onDevice, const Plane2D_float& f, int offset=0 );
18+
void write_plane2Dunscaled( const char* filename, const Plane2D_float& f, int offset=0 );
1919

20-
void dump_plane2Dfloat( const char* filename, bool onDevice, Plane2D_float& f );
21-
void dump_plane2Dfloat( const char* filename, Plane2D_float& f );
20+
void dump_plane2Dfloat( const char* filename, bool onDevice, const Plane2D_float& f );
21+
void dump_plane2Dfloat( const char* filename, const Plane2D_float& f );
2222
} // namespace popsift
2323

‎src/popsift/sift_octave.cu‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -153,7 +153,7 @@ void Octave::download_and_save_array( const char* basename, int octave )
153153
POP_CUDA_FATAL_TEST(err, "cudaMemcpy3D failed: ");
154154

155155
for( int l = 0; l<_levels; l++ ) {
156-
Plane2D_float p(width, height, &array[l*width*height], width * sizeof(float));
156+
const Plane2D_float p(width, height, &array[l*width*height], width * sizeof(float));
157157

158158
ostringstream ostr;
159159
ostr << "dir-octave/" << basename << "-o-" << octave << "-l-" << l << ".pgm";
@@ -173,7 +173,7 @@ void Octave::download_and_save_array( const char* basename, int octave )
173173
POP_CUDA_FATAL_TEST(err, "cudaMemcpy3D failed: ");
174174

175175
for( int l = 0; l<_levels; l++ ) {
176-
Plane2D_float p(width, height, &array[l*width*height], width * sizeof(float));
176+
const Plane2D_float p(width, height, &array[l*width*height], width * sizeof(float));
177177

178178
ostringstream ostr;
179179
ostr << "dir-interm/" << basename << "-o-" << octave << "-l-" << l << ".pgm";
@@ -189,7 +189,7 @@ void Octave::download_and_save_array( const char* basename, int octave )
189189
POP_CUDA_FATAL_TEST(err, "cudaMemcpy3D failed: ");
190190

191191
for (int l = 0; l<_levels - 1; l++) {
192-
Plane2D_float p(width, height, &array[l*width*height], width * sizeof(float));
192+
const Plane2D_float p(width, height, &array[l*width*height], width * sizeof(float));
193193

194194
ostringstream ostr;
195195
ostr << "dir-dog/d-" << basename << "-o-" << octave << "-l-" << l << ".pgm";

0 commit comments

Comments
 (0)