sharpetronics/QGIS
1
0
Fork 0
mirror of https://github.com/qgis/QGIS.git synced 2025-04-14 00:07:35 -04:00
Code Issues Packages Projects Releases Wiki Activity

[opencl] Fix hillshade renderer with 16bit rasters

Browse Source 
Fixes #21121
This commit is contained in:
Alessandro Pasotti 2019-02-02 10:53:11 +01:00
parent 4495699bfe
commit 1aef9cfa49
2 changed files with 46 additions and 11 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

9
resources/opencl_programs/hillshade_renderer.cl
View File

@ -1,3 +1,4 @@
// Note: "float *scanLine" may be replaced by code with actual input data type
__kernel void processNineCellWindow( __global float *scanLine1,
__global float *scanLine2,
__global float *scanLine3,
@ -39,8 +40,8 @@ __kernel void processNineCellWindow( __global float *scanLine1,
if ( x32 == rasterParams[0] ) x32 = x22;
if ( x33 == rasterParams[0] ) x33 = x22;
float derX = ( ( x13 + x23 + x23 + x33 ) - ( x11 + x21 + x21 + x31 ) ) / ( 8 * rasterParams[3] );
float derY = ( ( x31 + x32 + x32 + x33 ) - ( x11 + x12 + x12 + x13 ) ) / ( 8 * -rasterParams[4]);
float derX = ( ( x13 + x23 + x23 + x33 ) - ( x11 + x21 + x21 + x31 ) ) / ( 8.0f * rasterParams[3] );
float derY = ( ( x31 + x32 + x32 + x33 ) - ( x11 + x12 + x12 + x13 ) ) / ( 8.0f * -rasterParams[4]);
if ( derX == rasterParams[0] ||
derX == rasterParams[0] )
@ -90,7 +91,7 @@ __kernel void processNineCellWindow( __global float *scanLine1,
weight_270 * val270_mul_127 +
weight_315 * val315_mul_127 +
weight_360 * val360_mul_127 ) / xx_plus_yy ) /
( 1 + rasterParams[8] * xx_plus_yy );
( 1.0f + rasterParams[8] * xx_plus_yy );
res = clamp( 1.0f + cang_mul_127, 0.0f, 255.0f );
}
}
@ -99,7 +100,7 @@ __kernel void processNineCellWindow( __global float *scanLine1,
res = ( rasterParams[9] -
( derY * rasterParams[6] -
derX * rasterParams[7] )) /
sqrt( 1 + rasterParams[8] *
sqrt( 1.0f + rasterParams[8] *
( derX * derX + derY * derY ) );
res = res <= 0.0f ? 1.0f : 1.0f + res;
}

48
src/core/raster/qgshillshaderenderer.cpp
View File

@ -30,7 +30,9 @@
#include <chrono>
#include "qgssettings.h"
#endif
#include "qgsexception.h"
#include "qgsopenclutils.h"
#include "qdebug.h"
#endif
QgsHillshadeRenderer::QgsHillshadeRenderer( QgsRasterInterface *input, int band, double lightAzimuth, double lightAngle ):
@ -165,7 +167,8 @@ QgsRasterBlock *QgsHillshadeRenderer::block( int bandNo, const QgsRectangle &ext
bool useOpenCL( QgsOpenClUtils::enabled()
&& QgsOpenClUtils::available()
&& ( ! mRasterTransparency || mRasterTransparency->isEmpty() )
&& mAlphaBand <= 0 );
&& mAlphaBand <= 0
&& inputBlock->dataTypeSize() <= 4 );
// Check for sources
QString source;
if ( useOpenCL )
@ -190,6 +193,37 @@ QgsRasterBlock *QgsHillshadeRenderer::block( int bandNo, const QgsRectangle &ext
std::size_t inputDataTypeSize = inputBlock->dataTypeSize();
std::size_t outputDataTypeSize = outputBlock->dataTypeSize();
// Buffer scanline, 1px height, 2px wider
QString typeName;
switch ( inputBlock->dataType() )
{
case Qgis::DataType::Byte:
typeName = QStringLiteral( "unsigned char" );
break;
case Qgis::DataType::UInt16:
typeName = QStringLiteral( "unsigned int" );
break;
case Qgis::DataType::Int16:
typeName = QStringLiteral( "short" );
break;
case Qgis::DataType::UInt32:
typeName = QStringLiteral( "unsigned int" );
break;
case Qgis::DataType::Int32:
typeName = QStringLiteral( "int" );
break;
case Qgis::DataType::Float32:
typeName = QStringLiteral( "float" );
break;
default:
throw QgsException( QStringLiteral( "Unsupported data type for OpenCL processing.") );
}
if ( inputBlock->dataType() != Qgis::DataType::Float32 )
{
source.replace(QStringLiteral( "__global float *scanLine" ), QStringLiteral( "__global %1 *scanLine" ).arg( typeName ));
}
qDebug() << source;
// Data type for input is Float32 (4 bytes)
std::size_t scanLineWidth( inputBlock->width() + 2 );
std::size_t inputSize( inputDataTypeSize * inputBlock->width() );
@ -236,7 +270,6 @@ QgsRasterBlock *QgsHillshadeRenderer::block( int bandNo, const QgsRectangle &ext
// Whether use multidirectional
rasterParams.push_back( static_cast<float>( mMultiDirectional ) ); // 17
cl::Buffer rasterParamsBuffer( queue, rasterParams.begin(), rasterParams.end(), true, false, nullptr );
cl::Buffer scanLine1Buffer( ctx, CL_MEM_READ_ONLY, bufferSize, nullptr, nullptr );
cl::Buffer scanLine2Buffer( ctx, CL_MEM_READ_ONLY, bufferSize, nullptr, nullptr );
@ -245,13 +278,14 @@ QgsRasterBlock *QgsHillshadeRenderer::block( int bandNo, const QgsRectangle &ext
// Note that result buffer is an image
cl::Buffer resultLineBuffer( ctx, CL_MEM_WRITE_ONLY, outputDataTypeSize * width, nullptr, nullptr );
static cl::Program program;
static std::once_flag programBuilt;
std::call_once( programBuilt, [ = ]()
static std::map<Qgis::DataType, cl::Program> programCache;
cl::Program program = programCache[inputBlock->dataType()];
if (! program.get() )
{
// Create a program from the kernel source
program = QgsOpenClUtils::buildProgram( source, QgsOpenClUtils::ExceptionBehavior::Throw );
} );
programCache[inputBlock->dataType()] = QgsOpenClUtils::buildProgram( source, QgsOpenClUtils::ExceptionBehavior::Throw );
program = programCache[inputBlock->dataType()];
}
// Disable program cache when developing and testing cl program
// program = QgsOpenClUtils::buildProgram( ctx, source, QgsOpenClUtils::ExceptionBehavior::Throw );