Skip to content

Commit b09df53

Browse files
elpasoelpaso
committedAug 8, 2018
Use OpenCL command queue
·
ltr-3_40final-3_4_0
1 parent 16a49cd commit b09df53

File tree

8 files changed

+112
-266
lines changed

8 files changed

+112
-266
lines changed
 

‎resources/opencl_programs/aspect.cl

Lines changed: 39 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4,13 +4,19 @@ __kernel void processNineCellWindow( __global float *scanLine1,
44
__global float *scanLine2,
55
__global float *scanLine3,
66
__global float *resultLine,
7+
<<<<<<< 16a49cddaa18cb6d0b12335fe24c68cda183e1c0
78
__global float *rasterParams // mInputNodataValue, mOutputNodataValue, mZFactor, mCellSizeX, mCellSizeY
89
)
910
{
11+
=======
12+
__global float *rasterParams
13+
) {
14+
>>>>>>> Use OpenCL command queue
1015

1116
// Get the index of the current element
1217
const int i = get_global_id(0);
1318

19+
<<<<<<< 16a49cddaa18cb6d0b12335fe24c68cda183e1c0
1420
if ( scanLine2[i+1] == rasterParams[0] )
1521
{
1622
resultLine[i] = rasterParams[1];
@@ -36,5 +42,38 @@ __kernel void processNineCellWindow( __global float *scanLine1,
3642
{
3743
resultLine[i] = 180.0f + atan2pi( derX, derY ) * 180.0f;
3844
}
45+
=======
46+
// Do the operation
47+
//return (( (x31 - x11) + 2 * (x32 - x12) + (x33 - x13) ) / (8 * mCellSizeX))
48+
float derX = calcFirstDer( scanLine1[i], scanLine2[i], scanLine3[i],
49+
scanLine1[i+1], scanLine2[i+1], scanLine3[i+1],
50+
scanLine1[i+2], scanLine2[i+2], scanLine3[i+2],
51+
rasterParams[0], rasterParams[1], rasterParams[2], rasterParams[3]
52+
);
53+
//return (((x11 - x13) + 2 * (x21 - x23) + (x31 - x33)) / ( 8 * mCellSizeY));
54+
float derY = calcFirstDer( scanLine1[i+2], scanLine1[i+1], scanLine1[i],
55+
scanLine2[i+2], scanLine2[i+1], scanLine2[i],
56+
scanLine3[i+2], scanLine3[i+1], scanLine3[i],
57+
rasterParams[0], rasterParams[1], rasterParams[2], rasterParams[4]
58+
);
59+
60+
61+
if ( derX == rasterParams[1] || derY == rasterParams[1] ||
62+
( derX == 0.0f && derY == 0.0f) )
63+
{
64+
resultLine[i] = rasterParams[1];
65+
}
66+
else
67+
{
68+
// 180.0 / M_PI = 57.29577951308232
69+
float aspect = atan2( derX, derY ) * 57.29577951308232;
70+
if ( aspect < 0 )
71+
resultLine[i] = 90.0f - aspect;
72+
else if (aspect > 90.0f)
73+
// 360 + 90 = 450
74+
resultLine[i] = 450.0f - aspect;
75+
else
76+
resultLine[i] = 90.0 - aspect;
77+
>>>>>>> Use OpenCL command queue
3978
}
4079
}

‎resources/opencl_programs/slope.cl

Lines changed: 32 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4,13 +4,19 @@ __kernel void processNineCellWindow( __global float *scanLine1,
44
__global float *scanLine2,
55
__global float *scanLine3,
66
__global float *resultLine,
7+
<<<<<<< 16a49cddaa18cb6d0b12335fe24c68cda183e1c0
78
__global float *rasterParams // mInputNodataValue, mOutputNodataValue, mZFactor, mCellSizeX, mCellSizeY
89
)
910
{
11+
=======
12+
__global float *rasterParams
13+
) {
14+
>>>>>>> Use OpenCL command queue
1015

1116
// Get the index of the current element
1217
const int i = get_global_id(0);
1318

19+
<<<<<<< 16a49cddaa18cb6d0b12335fe24c68cda183e1c0
1420
if ( scanLine2[i+1] == rasterParams[0] )
1521
{
1622
resultLine[i] = rasterParams[1];
@@ -38,5 +44,31 @@ __kernel void processNineCellWindow( __global float *scanLine1,
3844
res = atanpi( res );
3945
resultLine[i] = res * 180.0f;
4046
}
47+
=======
48+
// Do the operation
49+
//return (( (x31 - x11) + 2 * (x32 - x12) + (x33 - x13) ) / (8 * mCellSizeX))
50+
float derX = calcFirstDer( scanLine1[i], scanLine2[i], scanLine3[i],
51+
scanLine1[i+1], scanLine2[i+1], scanLine3[i+1],
52+
scanLine1[i+2], scanLine2[i+2], scanLine3[i+2],
53+
rasterParams[0], rasterParams[1], rasterParams[2], rasterParams[3]
54+
);
55+
//return (((x11 - x13) + 2 * (x21 - x23) + (x31 - x33)) / ( 8 * mCellSizeY));
56+
float derY = calcFirstDer( scanLine1[i+2], scanLine1[i+1], scanLine1[i],
57+
scanLine2[i+2], scanLine2[i+1], scanLine2[i],
58+
scanLine3[i+2], scanLine3[i+1], scanLine3[i],
59+
rasterParams[0], rasterParams[1], rasterParams[2], rasterParams[4]
60+
);
61+
62+
63+
if ( derX == rasterParams[1] || derY == rasterParams[1] )
64+
{
65+
resultLine[i] = rasterParams[1];
66+
}
67+
else
68+
{
69+
float res = sqrt( derX * derX + derY * derY );
70+
res = atanpi( res );
71+
resultLine[i] = res * 180.0;
72+
>>>>>>> Use OpenCL command queue
4173
}
4274
}

‎src/analysis/raster/aspect.cl

Lines changed: 0 additions & 114 deletions
This file was deleted.

‎src/analysis/raster/qgsninecellfilter.cpp

Lines changed: 38 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -217,10 +217,20 @@ int QgsNineCellFilter::processRasterGPU( const QString &source, QgsFeedback *fee
217217

218218
// Prepare context and queue
219219
cl::Context ctx = QgsOpenClUtils::context();
220+
<<<<<<< 16a49cddaa18cb6d0b12335fe24c68cda183e1c0
220221
cl::CommandQueue queue( ctx );
221222

222223
//keep only three scanlines in memory at a time, make room for initial and final nodata
223224
QgsOpenClUtils::CPLAllocator<float> scanLine( xSize + 2 );
225+
=======
226+
cl::Context::setDefault( ctx );
227+
cl::CommandQueue queue( ctx );
228+
229+
//keep only three scanlines in memory at a time, make room for initial and final nodata
230+
QgsOpenClUtils::CPLAllocator<float> scanLine1( xSize + 2 );
231+
QgsOpenClUtils::CPLAllocator<float> scanLine2( xSize + 2 );
232+
QgsOpenClUtils::CPLAllocator<float> scanLine3( xSize + 2 );
233+
>>>>>>> Use OpenCL command queue
224234
QgsOpenClUtils::CPLAllocator<float> resultLine( xSize );
225235

226236
// Cast to float (because double just crashes on some GPUs)
@@ -289,7 +299,18 @@ int QgsNineCellFilter::processRasterGPU( const QString &source, QgsFeedback *fee
289299
{
290300
QgsDebugMsg( "Raster IO Error" );
291301
}
302+
<<<<<<< 16a49cddaa18cb6d0b12335fe24c68cda183e1c0
292303
queue.enqueueWriteBuffer( scanLine2Buffer, CL_TRUE, 0, bufferSize, scanLine.get() );
304+
=======
305+
}
306+
else
307+
{
308+
//normally fetch only scanLine3 and release scanline 1 if we move forward one row
309+
scanLine1.reset( scanLine2.release() );
310+
scanLine2.reset( scanLine3.release() );
311+
scanLine3.reset( xSize + 2 );
312+
}
313+
>>>>>>> Use OpenCL command queue
293314

294315
// Read scanline3: second real raster row
295316
if ( GDALRasterIO( rasterBand, GF_Read, 0, i + 1, xSize, 1, &scanLine[1], xSize, 1, GDT_Float32, 0, 0 ) != CE_None )
@@ -320,6 +341,23 @@ int QgsNineCellFilter::processRasterGPU( const QString &source, QgsFeedback *fee
320341
queue.enqueueWriteBuffer( *scanLineBuffer[rowIndex[2]], CL_TRUE, 0, bufferSize, scanLine.get() ); // row 0
321342
}
322343
}
344+
<<<<<<< 16a49cddaa18cb6d0b12335fe24c68cda183e1c0
345+
=======
346+
// Set first and last extra colums to nodata
347+
scanLine1[0] = scanLine1[xSize + 1] = mInputNodataValue;
348+
scanLine2[0] = scanLine2[xSize + 1] = mInputNodataValue;
349+
scanLine3[0] = scanLine3[xSize + 1] = mInputNodataValue;
350+
351+
// TODO: There is room for further optimization here: instead of replacing the buffers
352+
// we could just replace just hthe new one (the top row) and switch the order
353+
// of buffer arguments in the kernell call.
354+
errorCode = cl::enqueueWriteBuffer( scanLine1Buffer, CL_TRUE, 0,
355+
sizeof( float ) * ( xSize + 2 ), scanLine1.get() );
356+
errorCode = cl::enqueueWriteBuffer( scanLine2Buffer, CL_TRUE, 0,
357+
sizeof( float ) * ( xSize + 2 ), scanLine2.get() );
358+
errorCode = cl::enqueueWriteBuffer( scanLine3Buffer, CL_TRUE, 0,
359+
sizeof( float ) * ( xSize + 2 ), scanLine3.get() );
360+
>>>>>>> Use OpenCL command queue
323361

324362
kernel( cl::EnqueueArgs(
325363
queue,

‎src/analysis/raster/slope.cl

Lines changed: 0 additions & 106 deletions
This file was deleted.

0 commit comments

Comments
 (0)
Please sign in to comment.