View metadata, citation and similar papers at core.ac.uk

brought to you by

CORE

provided by DSpace at University of West Bohemia

ISSN 2464-4617(print) ISSN 2464-4625(CD)CSRN 2702

Computer Science Research Notes
http://www.WSCG.eu

CUDA-based SeqSLAM for Real-Time Place Recognition
Safa Ouerghi
Carthage Univ,
Sup’Com, GRESCOM,
2083 El Ghazela,
Tunisiasafa.ouerghi@supcom.tn

Remi Boutteau
Normandie Univ,
UNIROUENESIGELEC, IRSEEM
76000 Rouen, France
Remi.Boutteau@esigelec.fr

Fethi Tlili
Carthage Univ,
Sup’Com, GRESCOM,
2083 El Ghazela,
Tunisia
fethi.tlili@supcom.tn

Xavier Savatier
Normandie Univ,
UNIROUENESIGELEC, IRSEEM
76000 Rouen, France
Xavier.Savatier@esigelec.fr

ABSTRACT
Vehicle localization is a fundamental issue in autonomous navigation that has been extensively studied by the
Robotics community. An important paradigm for vehicle localization is based on visual place recognition which
relies on learning a database, then consecutively trying to find matchings between this database and the actual
visual input. An increasing interest has been directed to visual place recognition in varying conditions like day
and night cycles and seasonal changes. A major approach dealing with such challenges is Sequence SLAM (SeqSLAM) based on matching a sequence of images to the database instead of a single image. This algorithm allows
global pose recovery at the expense of a higher computational time. To solve this problem with a certain amountof speedup, we propose in this work, a CUDA-based solution for real-time place recognition with SeqSLAM.
We design a mapping of SeqSLAM to CUDA architecture and we describe, in detail, our hardware-specific implementation considerations as well as the parallelization methods. Performance analysis against existing CPU
implementation is also given, showing a speedup to six times faster than the CPU for common sized databases.
More speedup could be obtained when dealing with bigger databases.

Keywords
Place recognition, global localization, SeqSLAM, robotics, CUDA, GPU.

1

INTRODUCTION

results achieved by SeqSLAM, building a full matching
matrix introduces substantial computational complexityand enhancing it then computing scores for all straightlines presents an additional computational bottleneck.

Visual place recognition systems have become increasingly important for a variety of mobile robotic platforms and applications. A typical setup of a visual
recognition system contains a given database of images (the map) and the robot consecutively tries to findmatchings between this database and the actual visual
input (query images). The query is either a single image, which is most common, or an organized set of images i.e. a sequence. Recently, there has been an increasing focus on visual place recognition in varying
conditions which promoted the design of vision-based
systems that can work across common real-world perceptual changes such as varying weather conditions andday and night cycles [4, 16, 10, 11].

However, many problems are being solved using
programmable graphics hardware to achieve a certain
amount of speedup and the usage of GPU computation
has become a popular topic in the community. In fact,
the Compute Unified Device Architecture (CUDA)
has enabled graphics processors to be explicitly
programmed as general purpose shared-memory
multi-core processors with a high level of parallelism[6].
In this paper, we focus on an efficient parallelizationof the state-of-the-art visual place recognition SeqSLAM to achieve a certain amout of speedup. We consider a single GPU implementation and we describe thestrategies to efficiently map the problem components to
CUDA architecture.

A well-known successful approach is SeqSLAM (Sequence SLAM) for visual place recognition introducedby Milford and Weyth in [10] and also described in [9].
This work puts forward the concept of camera-based
GPS, which means performing global localization withensuring to never lose again and aims at matching image sequences under strong seasonal changes. SeqSLAM is based on the computation of image-by-image
dissimilarity scores between all query and database images, stored in a so-called difference matrix, then, computing a straight-line path through the full matching
matrix to finally select the path with the smallest sumof dissimilarity scores. Despite the quite remarkable

Short Papers Proceedings

The paper proceeds as follows: Section II presents
some related work. Section III summarises the SeqSLAM algorithm, and brievly reviews GPU architecture
and performance considerations. Implementation details are, subsequently, presented in Section IV. Finally,
Section IV, presents the results, discusses the outcomesof this paper and presents suggestions for future work.

131

ISBN 978-80-86943-50-3

ISSN 2464-4617(print) ISSN 2464-4625(CD)

2

CSRN 2702

RELATED WORKand patch-normalised to enhance contrast. Then, SeqSLAM builds a difference matrix holding SAD (Sumof Absolute Differences) scores between all query anddatabase image sequences. A key processing step is
to normalize the image difference values within their(spatially) local image neighborhoods. Subsequently,
a search for diagonals of low difference values isperformed over the defined sequence length. However,
SeqSLAM assumes similar speeds in repeated route
traversals and negligible accelerations, limiting its
performance in certain application contexts wherethese criteria are not met.

Visual place recognition is a well-defined but verybroad problem. State of the art techniques for visual
place recognition based on features, fall generally into
two categories: those that selectively extract parts of
the image and those that use the whole image without
a selection phase. Examples of the first category
are local feature descriptors such as scale-invariant
feature transforms (SIFT) [7] and speeded-up robustfeatures (SURF) [1]. In the second category, we findglobal descriptors such as HOG [3] and GIST [13]that use to process the whole image regardless of its
content. In fact, a considerable emphasis was placedon accelerating image descriptors for visual place
recognition using CUDA GPUs that many problems
are recently being solved with. In [20], Scale-invariant
feature transform (SIFT) has been accelerated using
CUDA GPU. Speeded-up robust features (SURF) has
also been ported to CUDA to achieve a certain amountof speedup [19]. In addition, GIST and HOG global
descriptors have been GPU accelerated in [18] and [15]respectively. We find, as well, another broad categoryof methods for visual place recognition based on a pure
image retrieval. This process can be accelerated using
inverted indices where image ID numbers are stored
against words that appear in the image. FabMAP,
a well known technique for visual place recognition
along a 1000-km path [2] has relied on an inverted
index with a bag-of-words model. Inverted index
technique has been recently implemented on GPU
demonstrating its efficiency [21]. The bag-of-words
technique has also been GPU-implemented [17].
Furthermore, image retrieval systems can be enhancedby adding topological information as both FabMap
and SeqSLAM do. Topological maps have in turnbeen improved by incorporating metric odometry
information such as SMART [14] and CAT-SLAM [8].
The metric information within the topological placenode can be stored as a sparse landmark map or as adense occupancy grid map. However, dense spatialmodeling has only become feasible in the past few
years with the advent of GPU technology [12].

3BACKGOUND

3.1

SeqSLAM

3.2

GPU architecture summary and performance considerations

NVIDIA GPUs comprise a set of SMs (Streaming Multiprocessors). Each SM consists of many SPs (Streaming Processors) and SFUs (Special Function Units),
which are simple but energy-efficient processing units
that execute instructions in a SIMD fashion (Single Instruction Multiple Data). The main memory of GPUs,
called global memory, can be accessed from all the SPs
in every SM. Furthermore, each SM has a set of registers and a small memory (shared memory). This sharedmemory is much faster than global memory. Each kernel is invoked by a set of threads that are grouped
into thread blocks. The blocks are distributed by the
hardware among the available SMs. Depending on theamount of required resources, each SM may be able tosimultaneously execute several blocks. Each block has
assigned an amount of shared memory that allows theexchanging of data among threads of the same block.
A more detailed description of NVIDIA’s GPU architecture can be found in [5]. The main important features influencing GPU performance are synchronization barriers, the trade-off between available threading and shared resources, coalescence issues in global
memory access and shared memory bank conflicts.

4

CUDA-BASED SEQSLAM IMPLEMENTATIONIn this section, we present the mapping strategies of
SeqSLAM to the CUDA architecture. In total, we
deal with three different kernels, each having a specificlaunch configuration that optimizes the use of the hardware. The first kernel is dedicated to the difference matrix computation, the second to contrast enhancementof the obtained difference matrix and the last one performs route searching in order to find the best match.

The SeqSLAM algorithm demonstrated impressive
place recognition performance across significant
condition variance such as seasonal changes and incase of low quality imagery (low resolution, low depth,
and image blur). In order to increase the discriminative
nature of the observation and to avoid the problem offalse-positives, location is represented in SeqSLAM asa sequence of images, rather than a single image fromone pose. Images are, beforehand, resolution-reduced

Short Papers Proceedings

Computer Science Research Notes
http://www.WSCG.eu

4.1

Difference Matrix Computation

The SeqSLAM algorithm uses Sum of Absolute Differences (SAD) to compare each query image Iq fromthe query sequence Q = (I1 , ..., IQ ) with Q = |Q| toeach database image ID from the database Sequence

132

ISBN 978-80-86943-50-3

ISSN 2464-4617(print) ISSN 2464-4625(CD)CSRN 2702

D = (I1 , ..., ID ) with D = |D| and calculates the difference score:d=1|IQ − ID |Rx Ry

input : D, Q
output: Difference Matrix M__shared__ uchar1 Bu fD [TileD × Rx × Ry ];__shared__ uchar1 Bu fQ [TileQ × Rx × Ry ];__shared__ int temp[Rx × Ry /2], Mt [TileD × TileQ ];blockIdx.x ← Bx , blockIdx.y ← By ;sd ← min(TileD , D − Bx × TileD );sq ← min(TileQ , D − By × TileQ );threadIdx.x ← Idx, Dim ← Rx × Ry ;for i ← Idx to sd × Dim do
Bu fD [i] ← D[i + Bx × TileD × Dim];i ← i + 8 × warpSize;endsynchronise the threads ;for i ← Idx to sq × Dim do
Bu fQ [i] ← Q[i + By × TileD × Dim];i ← i + 8 × warpSize;endsynchronise the threads ;for i ← 0 to sd dofor j ← 0 to sq do
Initialize temp ;for k ← Idx to Dim do
id ← mod(k, 8 ×WarpiSize);temp[id] ← temp[id] + |Bu fD [i × Dim +k] − Bu fQ [ j × Dim + k]|;i ← i + 8 × warpSize;endsynchronise the threads ;Mt [i × sq + j] ← Reduction(temp);
end
end
synchronise the threads ;Copy Mt to global memory;(1)

where d is the difference score and Rx and Ry are thehorizontal and vertical image dimensions, respectively.
The difference scores are assembled into the so-called
difference matrix M of size D × Q.
We assume that the database and query sequences reside in the GPU’s global memory as unsigned chartypes. In fact, database images are cropped, converted
to grayscale and resolution-reduced during the preprocessing step done offline and could therefore easily fit
into GPU’s global memory. On another note, it has
been reported in [16, 10] that, the longer the query sequences are, the more difficult is to localize the precise
match of a specific frame. This stems from the factthat although longer sequences are more distinguishable, the weight of each frame decreases with the sequence length. These considerations show that both the
database and the query sequences could be stored in
global memory, the first at the initialization of the GPU
after an offline preprocessing, and the second, online
acquired, is transferred from host to GPU after preprocessing as well.
Our parallelization strategy is based on a common strategy called tiling. This strategy consists in the partitionof the data into subsets called tiles, such that each tile
fits into the shared memory of a block. We apply thistechnique to the difference matrix computation where
each tile of the difference matrix is handled by a 2dimensional block. The tiling process requires for eachblock to load TileQ query images into shared memory
and TileD database images as described in Algorithm
1. Due to shared memory size limitations, we use theunsigned char type for the allocated buffers in sharedmemory. Although the use of 1-byte unsigned chartype is not recommended in shared memory as bankconflicts occur when threads in the same half warp access the same 4 or 8 bytes wide memory banks, thelimited size of shared memory imposes its use. In our
implementation, we use the buit-in type uchar1 denoting an unsigned char. We empirically adjust the pair
(TileD , TileQ ) through a time vs (TileD , TileQ ) evaluation. The best performance, overall, was achieved by(TileD , TileQ )= (4, 4).

4.2

Algorithm 1: Parallel difference matrix computation

is normalised within a fixed range l by subtracting the
local mean and dividing by the local standard deviation
to enhance the local matching contrast. This process
gives the new difference matrix M̂:M̂i =Mi − M̄lσl(2)where M̄l is the local mean and σl is the local standard
deviation, in a range l templates around template i.
We apply, as well, the tiling technique to enhance
the previously calculated difference matrix, stored in
global memory in a row major order. The tile size
is fixed to Tile×Q. The threads within each block
handle the computation of a tile and are designed to be
two-dimensional. In our design, each block of index
0<blockIdx.x<gridDim.x-1 will load Tile+l of M.
The blocks handling the border tiles, i.e. block0 andblockgridDim.x-1 only load Tile+ 12 l.

Difference Matrix Contrast Enhancement

After creating a difference matrix comparing all previously seen locations to each other using SAD scores,
SeqSLAM employs local neighbourhood normalisation
to remove biases from lighting variations between route
traversals. Each element Mi in the difference matrix M

Short Papers Proceedings

Computer Science Research Notes
http://www.WSCG.eu

The cuda kernel begins with loading a tile of M intoa preallocated shared memory of size (Tile+l) × Q.

133

ISBN 978-80-86943-50-3

ISSN 2464-4617(print) ISSN 2464-4625(CD)CSRN 2702

Computer Science Research Notes
http://www.WSCG.eu

input : Difference Matrix M, loutput: Enhanced Difference Matrix M̂

Next, the computation of Q means for Tile rows of
the enhanced M̂ is handled by Q×Tile threads where
threadIdx.x<Tile and threadIdx.y<Q. Each thread sums
(1 + l) values of M in a column where the y-coordinate
is threadIdx.y and the x-coordinate starts from threadIdx.x to finally divide the sum by the number of values
summed up i.e., (1 + l). After the kernel has finishedcomputing the means, threads are synchronised to subsequently perform the computation of standard deviations the same way they did for the means computation
thus described using the formula of standard deviationand the means already calculated. However, border tiles
handled by blockIdx.x=0 and blockIdx.x=gridDim.x-1
are special cases. In this way, for the first block, if
threadIdx.x≤ 21 l the local neighbourhood would be below (1 + l) and have thus to adapt its range with eachrow of M̂, i.e. each threadIdx.x. The same principle applies to the last block. If threaIdx.x≥ 12 l the range has
to be adapted as well. In our implementation, we compute the range for each thread using thread registers. Finally, each block writes in column major the tile of M̂ toglobal memory. Implementation details thus described
are presented in Algorithm 2. However, we only describe the means computation given that the same principle applies to standard deviations computation. The
final straightforward step is also not presented and it
consists in substracting the mean and dividing by thestandard deviation each element Mi to obtain M̂i .

__shared__ M̂sh [Tile][Q], Msh [l + Tile][Q];__shared__ Mean[Tile][Q];Initialize Msh to 0;
Idx ← threadIdx.x, Idy ← threadIdx.y,
Bid ← blockIdx.x, grid ← gridDim.x;a ← max(0, Tile × Blid − 12 l);b ← min(D, Tile × Blid + (Tile − 1) + 12 l);if Idx ≤ b − a and Idy < Q thenMsh [Idx][Idy] ← M[a + Idx][Idy];
endsynchronise the threads ;if (Idx ≤ Q) and (Idy < Tile) and (Bid = 0) thena ← max(0, Idx − 21 l);b ← min(1 + l, 1 + Idx + 12 l);Mean[Idx][Idy] ← 1b . ∑bi=0 Msh [a + i][Idy];
endsynchronise the threads ;if (Idx ≤ Q) and (Idy < Tile) and (0 < Bid < grid)then1Mean[Idx][Idy] ← 1+l. ∑1+li=0 Msh [Idx + i][Idy];
endsynchronise the threads ;if (Idx ≤ Q) and (Idy < Tile) and (Bid = grid − 1) thenb ← (1 + 12 l) + min(D − Bid × Tile − 1 − Idx, 12 l);Mean[Idx][Idy] ← 1b . ∑1+l
i=0 Msh [Idx + i][Idy];
endsynchronise the threads ;4.3

Algorithm 2: Parallel mean computation for Difference Matrix contrast enhancement

Route Searching

The inputs to the localization algorithm are the query
sequence locally acquired and the database Sequencethat serves as the visual memory of the robot. The algorithm attempts to find a subsequence of equal length
in the database that is the most similar to all regarded
subsequences. This is performed by building the image
difference matrix as stated above, contrast enhancing it
and then searching for the minimizing sub-route referring to connected regions of low difference in the difference matrix resembling to lines shapes. This process ispresented in Figure 1 where, for each starting search position in left column (marked by a green dot), a range ofslopes is traversed (highlighted by the semi-transparent
green area) modelizing possible other sub-routes thatwill be traversed for other slopes. For each matrix entry that is traversed during a sub-route traversal, its difference value is added to an accumulative value thatis initialized with zero in the beginning. This value iscalled sub-route score. Since there is a range of possible
slopes for each starting point there is an equal numberof scores for each sub-route starting point. When allscores are calculated for one starting point the minimal
score of them is selected. Finally, the sub-route with thesmallest score and the second smallest score are used tocompute the best database matching to the input query
sequence.

Short Papers Proceedings

Figure 1: Minimal slope search in the difference matrix.
One of the main important features influencing GPU
performance is coalescence issues in global memory
access. However, subroute scores computation require
multiple non coalesced global memory accesses (tosum up values in diagonals). We hence aim to design
a novel subroute score computation that enables coalesced accesses.
One common solution to coalescence issues is the useof shared memory to load data from global memory in acoalesced fashion, then, have contiguous threads stride
through it. Unlike global memory, there is no penalty
for strided access in shared memory especially when
there is no bank conflict. We assume that the difference

134

ISBN 978-80-86943-50-3

ISSN 2464-4617(print) ISSN 2464-4625(CD)CSRN 2702

matrix sizes are bigger than the shared memory of a single block, which would be the case for the majority ofreal applications. It is also worth noting that dependingon the database size is not a limiting factor in our implementation as the whole process of SeqSLAM is basedon a learning phase where several learnable parameters
have to be tuned.
We present our design using the mapping vector technique that assigns CUDA ressources including blocks,
threads and registers to SeqSLAM data. We will first
begin with presenting the global memory mapping vector followed by the SM mapping vector. The main
structure of SeqSLAM; the difference Matrix M of size
D × Q, is stored in global memory in a column major
order according to the following mapping vector:of slope possibilities. The process is, hence, repeated(1 + f loor(D/128)) times to account for the D possible subroutes. The 128 × k subscores are then written
to global memory via an atomicAdd between the active
blocks. We note that the atomicAdd is done inside theloop due to the limited size of shared memory and only
128 × k floats are allocated for Scores vector. The final
Score for each sub-route would be equal to the sum of
the calculated subscores for the slope possibility k.

5[M0 M1 ... MD−1 MD MD+1 ... M2D−1 ... MD(Q−1)−1 ... MDQ−1 ]
{z
}|
{z
} |
|
{z
}
col 0

col 1

Computer Science Research Notes
http://www.WSCG.eu

col Q-1(3)

5.1

This means that columns of size D are contiguously
stored in global memory. Hence M(0, 0) starts at location 0, M(0, 1) at location D and M(i, j) at location
j × D + i.

RESULTSIn this section, we describe the conducted experiments
to evaluate the performance of porting SeqSLAM to
CUDA GPU. We will first begin with exposing theexperimental setup, then, discuss the performance expressed in terms of timing of our implementation versus
exsisting CPU implementation.

Experimental Setup

For the experiments described in the following, we extracted still frames from the original video of the Nordland dataset1 , downsampled them to 64 × 32 and converted them into grayscale. It is worth noting that in
The proposed mapping vector of data on the SM related litterature, the downscaling was 64 × 32 for not
ressources is the following:
only the Nordland dataset, but also the Alderley2 . The
downscaling was even greater for other datasets such as
the Nurburgring dataset, equal to 32 × 24 and the ap[M0 , ..., M2D−1 , M2D , ..., M4D−1 , ..., M2×i×D , ..., M2×(i+1)×D−1 , proach of SeqSLAM is fundamentally based on impor{z
} |
{z
}
|
{z
}
|
block 0
block 1tant downscaling ratios.
block i..., MD(Q−2+MOD(Q,2))−1 , ..., MDQ−1 ]Furthermore, all the data reside in the GPU device|{z}memory at the beginning of each test, so there are noblock ceil(Q/2)(4) data transfers to CPU during the benchmarks to prevent
interactions with other factors in the study. The perThis means that each block i performs a coalesced formance of the experiments for Parallel SeqSLAM ismemory access to global memory in order to load data measured in time t in milliseconds(ms). The system on
from location 2 × i × D to location 2 × (i + 1) × D − 1 which our implementation was evaluated is equipped
into its shared memory which is equivalent to 2 with an i7 CPU running at up to 3.5 GHz, the intel i7columns of the difference Matrix M. The last block CORE. The CUDA device is an NVIDIA GeForce GTXof index blockIdx.x = ceil(Q/2) loads the remaining 850M running at 876 MHz with 4096 MB of GDDR5columns which would be either two if MOD(Q, 2) = 0 device memory. The evaluation has been performed
or only one column if not. Thus, a total number of with CUDA version 7.5 integrated with VisualStudioceil(Q/2) + 1 one-dimensional blocks is used in the 2012. At the first execution of SeqSLAM, memory allaunch configuration of the kernel and 4 warps of locations have to be performed. This is required only
once and takes about 10ms. All the experiments were
threads per block.
run for 5 database sequence lengths as presented in TaAs shown in Algorithm 3, the function of each block ble 1. Morever, 3 query sequences were used with 3
is, therefore, to compute a subscore for each sub-route different lengths as presented in Table 2. As stated bewhere the used values are only strided by D+offset fore, database sequences were obtained using the openwith offset∈ {0, 1, 2}. Within a block, the subscore of Source Code OpenseqSLAM3 by varying the parama given sub-route for all slope possibilities is handled eter imageSkip. However, the CPU-based code usedby a thread. We note that simultaneous accesses of
threads within a block to shared memory are done to 1 http://nrkbeta.no/2013/01/15/nordlandsbanen-minute-byconsecutive values, i.e. a column of the matrix M to minute-season-by-season/
minimize bank conflicts. As only 4 warps are used 2 https://wiki.qut.edu.au/display/cyphy/
per block, 128 × k subscores are computed in paral- Michael+Milford+Datasets+and+Downloads
lel by the active threads, where k designs the number 3 https://openslam.org/openseqslam.html

Short Papers Proceedings

135

ISBN 978-80-86943-50-3

ISSN 2464-4617(print) ISSN 2464-4625(CD)CSRN 2702

Computer Science Research Notes
http://www.WSCG.eu

in benchmarking is a C++/OpenCV port of OpenSeqSLAM 4 .
Database
D1
D2
D3
D4
D5

input : M̂, movmin , movmax
output: Matchindex , Matchscoreslopes ← movmax − movmin + 1;
Idx ← threadIdx.x;
Bid ← blockIdx.x;__shared__ Icremindices [slopes][Q], Score[slopes][Q];__shared__ Bu f f er[2 × D], Score[D][128];
if Idx = 0 thenSlopes ← FindSlopePossibilities(movmin , movmax );Icremindices ← SlopeIndices(Slopes);end
synchronise the threads ;for i ← Idx + 2 × D × Bid to 2 × D × (1 + Bid ) do
Bu f f er[i − 2 × D × Bid ] ← M̂[i];i ← i + 4 × warpSize;endsynchronise the threads ;for s ← Idx to D do
intindices[slopes][Q];for i ← 0 to slopes dofor j ← 0 to Q do
indices[i][ j] ← j × D + s +min(Icremindices [i × Q + j], D − s);endendsynchronise the threads ;for i ← 0 to slopes dosum ← 0; for j ← 0 to 2 doindex ← indices[i][ j + 2 × Bid ];sum ← sum + Bu f f er[index − 2 × D × Bid ];end
Subscore[i][mod(s, 128)] ← sum;end
synchronise the threads ;atomicAdd of Subscore in Temp of size D × slopes
in global memory;s ← s + 4 × warpSize;
end
Idx ← threadIdx.x × blockDim.x + blockIdx.x;
for k ← Idx to D do
Score[k] ← MinOverSlopes(Temp);k ← k + blockDim.x×gridDim.x;
end
synchronise the threads ;min1 ← FindMinWithShuffle(Score);
index1 ← FindIndexOfMin(min1 );Set values in Score around min1 of radius Rwindow to
max machine value;min2 ← FindMinWithShuffle(Score);
Matchindex ← index1 ;Matchscore ← min1 /min2 ;Database Seq Length
714
1428
2747
3570
5100

imageSkip
50
25
13
10
7

Table 1: Database Sequences
Query
query 1
query 2
query 3

Query Seq Length
11
20
32

Table 2: Query Sequences

5.2

CUDA based implementation vs CPU
based implementation

5.2.1

Difference Matrix Computation timing

In the first experiment, we measured the execution
time of difference matrix computation kernel using
the query and the database sequences presented inTable 1 and Table 2 respectively. We selected the pair
(TileD , TileQ ) = (4, 4). The evaluation is depicted
in Figure 2 for query = 32 showing an increasingly
important speedup of CUDA-based parallel SeqSLAM
with the database sequence length. The speedup isaveraged at 4×.

Figure 2: Difference matrix computaion.

5.2.2

Difference Matrix Enhancement timing

We subsequently evaluated the performance of parallel SeqSLAM contrast enhancement over Sequential
contrast enhancement. The optimal tile selected is
Tile = 20 and the performance for both sequential andparallel implementations is presented in Figure 3 forquery = 32. The obtained speedup exceeds 16× for D1
and almost 13× for D5. Overall, a good speedup wasshown with parallel seqSLAM averaged at 14×.

Algorithm 3: Parallel Route Searching

4 https://github.com/subokita/OpenSeqSLAM

Short Papers Proceedings

136

ISBN 978-80-86943-50-3

ISSN 2464-4617(print) ISSN 2464-4625(CD)CSRN 2702

Computer Science Research Notes
http://www.WSCG.eu(a) GPU vs CPU time

Figure 3: Difference matrix enhancement.

5.2.3

Route searching timing

The third experiment was dedicated to the comparisonof the Route searching execution time for both CPUbased and CUDA GPU- based implementations. The
speedup of GPU over CPU is clearly visible averaged
at almost 6× for the datasets used in the evaluation andis depicted in Figure 4 for query = 32. An interesting
point is that the route searching is not a botteleneck for
CPU when attempting to localize a single image usinga short query sequence as recommended. However, we
aim to design a GPU only solution and a good speedup
was though shown.

(b) Speedup

Figure 5: Performance of CUDA based SeqSLAM.
freedom along its track and corresponding images fromtwo traverses in different conditions overlap almost perfectly. This condition is only met for certain applications where the camera has only one degree of freedom
and SeqSLAM remains a major approach from which
many recent approaches are being built upon. Hence,since our aim is to enhance timing and not localization
precision, we have relied on the Nordland dataset images as input for different database and query sequences
sizes. Furthermore, we aimed to design a GPU only
solution as in a real application for a robotic platform,
several modules are required and we can put some tasks
that do not depend on the results of the kernel on the
CPU while the GPU is executing a different task inorder to achieve a better performance and reduce thehardware complexity. In fact, this is possible thanks tothe asynchronous execution feature between the CPU
and the GPU of CUDA systems meaning that the control returns to the CPU immediately after the kernel islaunched.

Figure 4: Route searching.

5.2.4

Performance anlalysis of CUDA based SeqSLAM

In Figure 5, we show the performance results of parallel SeqSLAM using a CUDA GPU. Firstly, in Figure
5a, we compare the mean computation time of CPU and
GPU implementations, for the different database andquery sequences used in this evaluation. We show acomputation time even more important for CPU reaching 298ms for D5 and query3 against 51ms for GPU. In
figure 5b we demonstrate that the speedup is about 6×compared to the CPU implementation.

5.3

6

In this paper, we have proposed a parallel CUDA GPU
based implementation of Sequence SLAM (SeqSLAM)
for visual place recognition. SeqSLAM is a well-known
successful approach for mobile localization and place
recognition in varying conditions like seasonal changes
and day and night cycles when certain conditions are
met. Our parallelization method was based on the allocation of the three major steps of the approach tothree GPU kernels, each of which with a specific launch

Discussion

A special feature of the Nordland dataset is that the
viewpoint for the database and query sequences is exactly the same. The camera has only one degree of

Short Papers Proceedings

CONCLUSION

137

ISBN 978-80-86943-50-3

ISSN 2464-4617(print) ISSN 2464-4625(CD)CSRN 2702

configuration promoting the best possible performance.
Experimentations showed promising results thanks tothe parallel exploitation of CUDA ressources that offered a good overall speedup averaged at 6 times better than its CPU counterpart. As future work, we plan
to apply other algorithmic techniques instead of simple
image difference to deal better with a reasonable viewpoint change. We plan also to use a multi-GPU system to scale with bigger databases and achieve a greater
speedup.

7[12] R. A. Newcombe, S. J. Lovegrove, and A. J.
Davison. Dtam: Dense tracking and mapping
in real-time. Proc. of International Conference
on Computer Vision (ICCV), Barcelona, Spain,
pages 2320–2327, 2011.
[13] A. Oliva and A. Torralba. Building the gist of ascene: the role of global image features in recognition. In. Visual Perception, Progress in Brain
Research, 155, Part B:23 – 36, 2006.
[14] E. Pepperell, P. Corke, and M. Milford. Allenvironment visual place recognition with
SMART. Proc. of IEEE International Conference on Robotics and Automation (ICRA), Hong
Kong, China, pages 1612–1618, 2014.
[15] V. Prisacariu and I. Reid. fasthog - a real-time
gpu implementation of hog. Technical Report2310/09, Department of Engineering Science,
Oxford University, 2009.
[16] N. Sunderhauf, P. Neubert, and P. Protzel. Predicting the change, a step towards life-long operationin everyday environments. 2013.
[17] K. E. A. van de Sande, T. Gevers, and C. G. M.
Snoek. Empowering visual categorization withthe gpu. In. IEEE Transactions on Multimedia(IEEE Trans. Multimedia), 13:60–70, 2011.
[18] Y. Wang, Z. Feng, H. Guo, C. He, and Y. Yang.
Scene recognition acceleration using cuda andopenmp. Proc. of First International Conferenceon Information Science and Engineering (ICISE),
Nanjing, China, pages 1422–1425, 2009.
[19] W. Yan, X. Shi, X. Yan, and L. Wang. Computing opensurf on opencl and general purpose gpu.
In. International Journal of Advanced Robotic
Systems (IJARS), 10:375, 2013.
[20] Z. Yonglong, M. Kuizhi, J. Xiang, and D. Peixiang. Parallelization and optimization of sift on
gpu using cuda. Proc. of IEEE International
Conference on High Performance Computing andCommunications (HPCC), Zhangjiajie, China,
pages 1351–1358, 2013.
[21] J. Zhou, Q. Guo, H. V. Jagadish, W. Luan,
A. K. H. Tung, Y. Yang, and Y. Zheng. Generic inverted index on the GPU. In. Computing Research
Repository (CoRR), abs/1603.08390, 2016.

REFERENCES[1]

H. Bay, A. Ess, T. Tuytelaars, and L. Van Gool.
Speeded-up robust features (surf). In. Elsevier
Computer Vision Image Understanding, 110:346–
359, 2008.
[2] M. Cummins and P. Newman. Appearance-onlyslam at large scale with fab-map 2.0. In. The International Journal of Robotics Research (IJRR),
30:1100–1123, 2011.
[3] N. Dalal and B. Triggs. Histograms of oriented
gradients for human detection. Proc. of IEEE
Computer Society Conference on Computer Vision and Pattern Recognition (CVPR), Washington, DC, USA, 110:886–893, 2005.
[4] E. Johns and G. Z. Yang. Feature co-occurrence
maps: Appearance-based localisation throughout
the day. pages 3212–3218, 2013.
[5] D. B. Kirk and W. W. Hwu. Programming Massively Parallel Processors: a Hands-on Approach.
2010.
[6] E. Lindholm, J. Nickolls, S. Oberman, and
J. Montrym. Nvidia tesla: A unified graphics
and computing architecture. In. IEEE Microwave
Magazine (IEEE Micro), 28:39–55, 2008.
[7] D. G. Lowe. Object recognition from local scaleinvariant features. Proc. of the International Conference on Computer Vision (ICCV), Washington,
DC, USA, 2:1150–, 1999.
[8] W. Maddern, M. Milford, and G. Wyeth. Catslam: probabilistic localisation and mapping using a continuous appearance-based trajectory. In.
The International Journal of Robotics Research(IJRR), 31:429–451, 2012.
[9] M. Milford. Vision-based place recognition: how
low can you go? In. The International Journal of
Robotics Research (IJRR), 32:766–789, 2013.
[10] M. Milford and G. Wyeth. Seqslam: Visual routebased navigation for sunny summer days andstormy winter nights. pages 1643–1649, 2012.
[11] T. Naseer, L. Spinello, W. Burgard, and C. Stachnis. Robust visual robot localization across seasons using network flows. pages 2564–2570,
2014.

Short Papers Proceedings

Computer Science Research Notes
http://www.WSCG.eu

138

ISBN 978-80-86943-50-3