Skip to content

Commit 1aef9cf

Browse files
committed
[opencl] Fix hillshade renderer with 16bit rasters
Fixes #21121
1 parent 4495699 commit 1aef9cf

File tree

2 files changed

+46
-11
lines changed

2 files changed

+46
-11
lines changed

resources/opencl_programs/hillshade_renderer.cl

Lines changed: 5 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,3 +1,4 @@
1+
// Note: "float *scanLine" may be replaced by code with actual input data type
12
__kernel void processNineCellWindow( __global float *scanLine1,
23
__global float *scanLine2,
34
__global float *scanLine3,
@@ -39,8 +40,8 @@ __kernel void processNineCellWindow( __global float *scanLine1,
3940
if ( x32 == rasterParams[0] ) x32 = x22;
4041
if ( x33 == rasterParams[0] ) x33 = x22;
4142

42-
float derX = ( ( x13 + x23 + x23 + x33 ) - ( x11 + x21 + x21 + x31 ) ) / ( 8 * rasterParams[3] );
43-
float derY = ( ( x31 + x32 + x32 + x33 ) - ( x11 + x12 + x12 + x13 ) ) / ( 8 * -rasterParams[4]);
43+
float derX = ( ( x13 + x23 + x23 + x33 ) - ( x11 + x21 + x21 + x31 ) ) / ( 8.0f * rasterParams[3] );
44+
float derY = ( ( x31 + x32 + x32 + x33 ) - ( x11 + x12 + x12 + x13 ) ) / ( 8.0f * -rasterParams[4]);
4445

4546
if ( derX == rasterParams[0] ||
4647
derX == rasterParams[0] )
@@ -90,7 +91,7 @@ __kernel void processNineCellWindow( __global float *scanLine1,
9091
weight_270 * val270_mul_127 +
9192
weight_315 * val315_mul_127 +
9293
weight_360 * val360_mul_127 ) / xx_plus_yy ) /
93-
( 1 + rasterParams[8] * xx_plus_yy );
94+
( 1.0f + rasterParams[8] * xx_plus_yy );
9495
res = clamp( 1.0f + cang_mul_127, 0.0f, 255.0f );
9596
}
9697
}
@@ -99,7 +100,7 @@ __kernel void processNineCellWindow( __global float *scanLine1,
99100
res = ( rasterParams[9] -
100101
( derY * rasterParams[6] -
101102
derX * rasterParams[7] )) /
102-
sqrt( 1 + rasterParams[8] *
103+
sqrt( 1.0f + rasterParams[8] *
103104
( derX * derX + derY * derY ) );
104105
res = res <= 0.0f ? 1.0f : 1.0f + res;
105106
}

src/core/raster/qgshillshaderenderer.cpp

Lines changed: 41 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -30,7 +30,9 @@
3030
#include <chrono>
3131
#include "qgssettings.h"
3232
#endif
33+
#include "qgsexception.h"
3334
#include "qgsopenclutils.h"
35+
#include "qdebug.h"
3436
#endif
3537

3638
QgsHillshadeRenderer::QgsHillshadeRenderer( QgsRasterInterface *input, int band, double lightAzimuth, double lightAngle ):
@@ -165,7 +167,8 @@ QgsRasterBlock *QgsHillshadeRenderer::block( int bandNo, const QgsRectangle &ext
165167
bool useOpenCL( QgsOpenClUtils::enabled()
166168
&& QgsOpenClUtils::available()
167169
&& ( ! mRasterTransparency || mRasterTransparency->isEmpty() )
168-
&& mAlphaBand <= 0 );
170+
&& mAlphaBand <= 0
171+
&& inputBlock->dataTypeSize() <= 4 );
169172
// Check for sources
170173
QString source;
171174
if ( useOpenCL )
@@ -190,6 +193,37 @@ QgsRasterBlock *QgsHillshadeRenderer::block( int bandNo, const QgsRectangle &ext
190193
std::size_t inputDataTypeSize = inputBlock->dataTypeSize();
191194
std::size_t outputDataTypeSize = outputBlock->dataTypeSize();
192195
// Buffer scanline, 1px height, 2px wider
196+
QString typeName;
197+
switch ( inputBlock->dataType() )
198+
{
199+
case Qgis::DataType::Byte:
200+
typeName = QStringLiteral( "unsigned char" );
201+
break;
202+
case Qgis::DataType::UInt16:
203+
typeName = QStringLiteral( "unsigned int" );
204+
break;
205+
case Qgis::DataType::Int16:
206+
typeName = QStringLiteral( "short" );
207+
break;
208+
case Qgis::DataType::UInt32:
209+
typeName = QStringLiteral( "unsigned int" );
210+
break;
211+
case Qgis::DataType::Int32:
212+
typeName = QStringLiteral( "int" );
213+
break;
214+
case Qgis::DataType::Float32:
215+
typeName = QStringLiteral( "float" );
216+
break;
217+
default:
218+
throw QgsException( QStringLiteral( "Unsupported data type for OpenCL processing.") );
219+
}
220+
221+
if ( inputBlock->dataType() != Qgis::DataType::Float32 )
222+
{
223+
source.replace(QStringLiteral( "__global float *scanLine" ), QStringLiteral( "__global %1 *scanLine" ).arg( typeName ));
224+
}
225+
qDebug() << source;
226+
193227
// Data type for input is Float32 (4 bytes)
194228
std::size_t scanLineWidth( inputBlock->width() + 2 );
195229
std::size_t inputSize( inputDataTypeSize * inputBlock->width() );
@@ -236,7 +270,6 @@ QgsRasterBlock *QgsHillshadeRenderer::block( int bandNo, const QgsRectangle &ext
236270
// Whether use multidirectional
237271
rasterParams.push_back( static_cast<float>( mMultiDirectional ) ); // 17
238272

239-
240273
cl::Buffer rasterParamsBuffer( queue, rasterParams.begin(), rasterParams.end(), true, false, nullptr );
241274
cl::Buffer scanLine1Buffer( ctx, CL_MEM_READ_ONLY, bufferSize, nullptr, nullptr );
242275
cl::Buffer scanLine2Buffer( ctx, CL_MEM_READ_ONLY, bufferSize, nullptr, nullptr );
@@ -245,13 +278,14 @@ QgsRasterBlock *QgsHillshadeRenderer::block( int bandNo, const QgsRectangle &ext
245278
// Note that result buffer is an image
246279
cl::Buffer resultLineBuffer( ctx, CL_MEM_WRITE_ONLY, outputDataTypeSize * width, nullptr, nullptr );
247280

248-
static cl::Program program;
249-
static std::once_flag programBuilt;
250-
std::call_once( programBuilt, [ = ]()
281+
static std::map<Qgis::DataType, cl::Program> programCache;
282+
cl::Program program = programCache[inputBlock->dataType()];
283+
if (! program.get() )
251284
{
252285
// Create a program from the kernel source
253-
program = QgsOpenClUtils::buildProgram( source, QgsOpenClUtils::ExceptionBehavior::Throw );
254-
} );
286+
programCache[inputBlock->dataType()] = QgsOpenClUtils::buildProgram( source, QgsOpenClUtils::ExceptionBehavior::Throw );
287+
program = programCache[inputBlock->dataType()];
288+
}
255289

256290
// Disable program cache when developing and testing cl program
257291
// program = QgsOpenClUtils::buildProgram( ctx, source, QgsOpenClUtils::ExceptionBehavior::Throw );

0 commit comments

Comments
 (0)