Skip to content
GitLab
Projects
Groups
Snippets
/
Help
Help
Support
Community forum
Keyboard shortcuts
?
Submit feedback
Contribute to GitLab
Sign in
Toggle navigation
Menu
Open sidebar
Houman Mirzaalian Dastjerdi
waLBerla
Commits
8d37d83c
Commit
8d37d83c
authored
Jun 20, 2018
by
Houman Mirzaalian Dastjerdi
Browse files
Clean code of Blur image for merge requests
parent
12fea53b
Pipeline
#10015
failed with stage
in 185 minutes and 9 seconds
Changes
3
Pipelines
1
Hide whitespace changes
Inline
Side-by-side
apps/tutorials/cuda/02_BlurImage_cuda.cpp
View file @
8d37d83c
...
...
@@ -13,8 +13,8 @@
// You should have received a copy of the GNU General Public License along
// with waLBerla (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
//
//! \file 03_
GameOfLife
.cpp
//! \author
Martin Bauer <martin.bauer
@fau.de>
//! \file 03_
BlurImage_cuda
.cpp
//! \author
Houman Mirzaalian Dastjerdi <houman.mirzaalian
@fau.de>
//
//======================================================================================================================
...
...
@@ -59,11 +59,6 @@ const int blurKernelWidth = 3;
const
float
blurKernelSigma
=
1000.
;
//------------------------
ScalarField
*
createField
(
IBlock
*
const
block
,
StructuredBlockStorage
*
const
storage
)
{
return
new
ScalarField
(
...
...
@@ -76,16 +71,13 @@ ScalarField * createField( IBlock* const block, StructuredBlockStorage* const st
make_shared
<
cuda
::
HostFieldAllocator
<
double
>
>
()
// allocator for host pinned memory
);
}
//------------------------
class
CUDA_BLUR
{
public:
CUDA_BLUR
(
BlockDataID
gpuFieldSrcID
,
BlockDataID
gpuFieldDstID
,
int
numRows
,
int
numCols
,
const
int
filterWidth
)
:
gpuFieldSrcID_
(
gpuFieldSrcID
),
gpuFieldDstID_
(
gpuFieldDstID
),
numRows_
(
numRows
),
numCols_
(
numCols
),
filterWidth_
(
filterWidth
)
:
gpuFieldSrcID_
(
gpuFieldSrcID
),
gpuFieldDstID_
(
gpuFieldDstID
),
filterWidth_
(
filterWidth
)
{
//create and fill the filter we will convolve with
auto
h_filter
=
new
float
[
blurKernelWidth
*
blurKernelWidth
];
...
...
@@ -108,8 +100,8 @@ public:
}
}
cudaMalloc
(
&
d_filt
t
er_
,
blurKernelWidth
*
blurKernelWidth
*
sizeof
(
float
));
cudaMemcpy
(
d_filt
t
er_
,
h_filter
,
blurKernelWidth
*
blurKernelWidth
*
sizeof
(
float
),
cudaMemcpyHostToDevice
);
cudaMalloc
(
&
d_filter_
,
blurKernelWidth
*
blurKernelWidth
*
sizeof
(
float
));
cudaMemcpy
(
d_filter_
,
h_filter
,
blurKernelWidth
*
blurKernelWidth
*
sizeof
(
float
),
cudaMemcpyHostToDevice
);
}
...
...
@@ -122,16 +114,7 @@ public:
myKernel
.
addFieldIndexingParam
(
cuda
::
FieldIndexing
<
double
>::
xyz
(
*
srcCudaField
)
);
myKernel
.
addFieldIndexingParam
(
cuda
::
FieldIndexing
<
double
>::
xyz
(
*
dstCudaField
)
);
// typedef typename boost::function_traits<decltype(gaussian_blur)>::arg5_type arg5_type;
// WALBERLA_LOG_DEVEL("Type is" << typeid(arg5_type).name());
// auto equal_res = boost::is_same<const float *const, typename boost::function_traits<decltype(gaussian_blur)>::arg5_type>::value;
// WALBERLA_LOG_DEVEL("Types are equal " << equal_res);
myKernel
.
addParam
<
int
>
(
numRows_
);
myKernel
.
addParam
<
int
>
(
numCols_
);
myKernel
.
addParam
<
float
*>
(
d_filtter_
);
myKernel
.
addParam
<
float
*>
(
d_filter_
);
myKernel
.
addParam
<
int
>
(
filterWidth_
);
myKernel
();
...
...
@@ -142,25 +125,17 @@ private:
BlockDataID
gpuFieldSrcID_
;
BlockDataID
gpuFieldDstID_
;
int
filterWidth_
;
int
numRows_
;
int
numCols_
;
float
*
d_filtter_
;
float
*
d_filter_
;
};
//------------------------
int
main
(
int
argc
,
char
**
argv
)
{
walberla
::
Environment
env
(
argc
,
argv
);
geometry
::
GrayScaleImage
image
(
"image.png"
);
int
numRows
=
image
.
height
();
int
numCols
=
image
.
width
();
// Create blocks
shared_ptr
<
StructuredBlockForest
>
blocks
=
blockforest
::
createUniformBlockGrid
(
uint_t
(
2
)
,
uint_t
(
2
),
uint_t
(
1
),
// number of blocks in x,y,z direction
...
...
@@ -192,12 +167,12 @@ int main( int argc, char ** argv )
commScheme
.
addDataToCommunicate
(
make_shared
<
Packing
>
(
gpuFieldSrcID
)
);
// Create Timeloop
const
uint_t
numberOfTimesteps
=
uint_t
(
5
00
);
// number of timesteps for non-gui runs
const
uint_t
numberOfTimesteps
=
uint_t
(
1
00
);
// number of timesteps for non-gui runs
SweepTimeloop
timeloop
(
blocks
,
numberOfTimesteps
);
// Registering the sweep
timeloop
.
add
()
<<
BeforeFunction
(
commScheme
,
"Communication"
)
<<
Sweep
(
CUDA_BLUR
(
gpuFieldSrcID
,
gpuFieldDstID
,
numRows
,
numCols
,
blurKernelWidth
),
"GameOfLifeSweep"
);
<<
Sweep
(
CUDA_BLUR
(
gpuFieldSrcID
,
gpuFieldDstID
,
blurKernelWidth
),
"GameOfLifeSweep"
);
timeloop
.
add
()
<<
Sweep
(
cuda
::
fieldCpyFunctor
<
ScalarField
,
GPUField
>
(
cpuFieldID
,
gpuFieldDstID
)
);
...
...
apps/tutorials/cuda/02_BlurImage_cuda.cu
View file @
8d37d83c
...
...
@@ -14,28 +14,20 @@ namespace walberla {
__global__
void
gaussian_blur
(
cuda
::
FieldAccessor
<
double
>
inputChannel
,
cuda
::
FieldAccessor
<
double
>
outputChannel
,
int
numRows
,
int
numCols
,
float
*
filter
,
int
filterWidth
)
{
inputChannel
.
set
(
blockIdx
,
threadIdx
);
outputChannel
.
set
(
blockIdx
,
threadIdx
);
/*
int px = blockIdx.x * blockDim.x + threadIdx.x;
int py = blockIdx.y * blockDim.y + threadIdx.y;
if (px >= numCols || py >= numRows) {
return;
}*/
double
c
=
0.0
f
;
for
(
int
fx
=
0
;
fx
<
filterWidth
;
fx
++
)
{
for
(
int
fy
=
0
;
fy
<
filterWidth
;
fy
++
)
{
int
imagex
=
fx
-
filterWidth
/
2
;
int
imagey
=
fy
-
filterWidth
/
2
;
//imagex = min(max(imagex, 0), numCols - 1);
//imagey = min(max(imagey, 0), numRows - 1);
c
+=
(
filter
[
fy
*
filterWidth
+
fx
]
*
inputChannel
.
getNeighbor
(
imagex
,
imagey
,
0
)
);
}
...
...
apps/tutorials/cuda/02_BlurImage_cuda.h
View file @
8d37d83c
//
// Created by po60nani on 6/18/18.
//
#pragma once
//
// Created by po60nani on 6/13/18.
//
#include
<stdio.h>
#include
<cuda.h>
...
...
@@ -19,12 +10,5 @@ namespace walberla {
__global__
void
gaussian_blur
(
cuda
::
FieldAccessor
<
double
>
inputChannel
,
cuda
::
FieldAccessor
<
double
>
outputChannel
,
int
numRows
,
int
numCols
,
float
*
filter
,
int
filterWidth
);
}
Write
Preview
Supports
Markdown
0%
Try again
or
attach a new file
.
Cancel
You are about to add
0
people
to the discussion. Proceed with caution.
Finish editing this message first!
Cancel
Please
register
or
sign in
to comment