[opencl] Test with image2d

This commit is contained in:
Alessandro Pasotti 2018-07-04 14:27:41 +02:00
parent 12fa896554
commit a1a65bb4e8
6 changed files with 42 additions and 79 deletions

View File

@ -4,19 +4,13 @@ __kernel void processNineCellWindow( __global float *scanLine1,
__global float *scanLine2,
__global float *scanLine3,
__global float *resultLine,
<<<<<<< 16a49cddaa18cb6d0b12335fe24c68cda183e1c0
__global float *rasterParams // mInputNodataValue, mOutputNodataValue, mZFactor, mCellSizeX, mCellSizeY
)
{
=======
__global float *rasterParams
) {
>>>>>>> Use OpenCL command queue
// Get the index of the current element
const int i = get_global_id(0);
<<<<<<< 16a49cddaa18cb6d0b12335fe24c68cda183e1c0
if ( scanLine2[i+1] == rasterParams[0] )
{
resultLine[i] = rasterParams[1];
@ -42,38 +36,5 @@ __kernel void processNineCellWindow( __global float *scanLine1,
{
resultLine[i] = 180.0f + atan2pi( derX, derY ) * 180.0f;
}
=======
// Do the operation
//return (( (x31 - x11) + 2 * (x32 - x12) + (x33 - x13) ) / (8 * mCellSizeX))
float derX = calcFirstDer( scanLine1[i], scanLine2[i], scanLine3[i],
scanLine1[i+1], scanLine2[i+1], scanLine3[i+1],
scanLine1[i+2], scanLine2[i+2], scanLine3[i+2],
rasterParams[0], rasterParams[1], rasterParams[2], rasterParams[3]
);
//return (((x11 - x13) + 2 * (x21 - x23) + (x31 - x33)) / ( 8 * mCellSizeY));
float derY = calcFirstDer( scanLine1[i+2], scanLine1[i+1], scanLine1[i],
scanLine2[i+2], scanLine2[i+1], scanLine2[i],
scanLine3[i+2], scanLine3[i+1], scanLine3[i],
rasterParams[0], rasterParams[1], rasterParams[2], rasterParams[4]
);
if ( derX == rasterParams[1] || derY == rasterParams[1] ||
( derX == 0.0f && derY == 0.0f) )
{
resultLine[i] = rasterParams[1];
}
else
{
// 180.0 / M_PI = 57.29577951308232
float aspect = atan2( derX, derY ) * 57.29577951308232;
if ( aspect < 0 )
resultLine[i] = 90.0f - aspect;
else if (aspect > 90.0f)
// 360 + 90 = 450
resultLine[i] = 450.0f - aspect;
else
resultLine[i] = 90.0 - aspect;
>>>>>>> Use OpenCL command queue
}
}

View File

@ -4,19 +4,13 @@ __kernel void processNineCellWindow( __global float *scanLine1,
__global float *scanLine2,
__global float *scanLine3,
__global float *resultLine,
<<<<<<< 16a49cddaa18cb6d0b12335fe24c68cda183e1c0
__global float *rasterParams // mInputNodataValue, mOutputNodataValue, mZFactor, mCellSizeX, mCellSizeY
)
{
=======
__global float *rasterParams
) {
>>>>>>> Use OpenCL command queue
// Get the index of the current element
const int i = get_global_id(0);
<<<<<<< 16a49cddaa18cb6d0b12335fe24c68cda183e1c0
if ( scanLine2[i+1] == rasterParams[0] )
{
resultLine[i] = rasterParams[1];
@ -44,31 +38,5 @@ __kernel void processNineCellWindow( __global float *scanLine1,
res = atanpi( res );
resultLine[i] = res * 180.0f;
}
=======
// Do the operation
//return (( (x31 - x11) + 2 * (x32 - x12) + (x33 - x13) ) / (8 * mCellSizeX))
float derX = calcFirstDer( scanLine1[i], scanLine2[i], scanLine3[i],
scanLine1[i+1], scanLine2[i+1], scanLine3[i+1],
scanLine1[i+2], scanLine2[i+2], scanLine3[i+2],
rasterParams[0], rasterParams[1], rasterParams[2], rasterParams[3]
);
//return (((x11 - x13) + 2 * (x21 - x23) + (x31 - x33)) / ( 8 * mCellSizeY));
float derY = calcFirstDer( scanLine1[i+2], scanLine1[i+1], scanLine1[i],
scanLine2[i+2], scanLine2[i+1], scanLine2[i],
scanLine3[i+2], scanLine3[i+1], scanLine3[i],
rasterParams[0], rasterParams[1], rasterParams[2], rasterParams[4]
);
if ( derX == rasterParams[1] || derY == rasterParams[1] )
{
resultLine[i] = rasterParams[1];
}
else
{
float res = sqrt( derX * derX + derY * derY );
res = atanpi( res );
resultLine[i] = res * 180.0;
>>>>>>> Use OpenCL command queue
}
}

View File

@ -233,6 +233,11 @@ int QgsNineCellFilter::processRasterGPU( const QString &source, QgsFeedback *fee
>>>>>>> Use OpenCL command queue
QgsOpenClUtils::CPLAllocator<float> resultLine( xSize );
<<<<<<< 12fa896554321892e88ca05407cf125a7ccf92c1
=======
cl_int errorCode = 0;
>>>>>>> [opencl] Test with image2d
// Cast to float (because double just crashes on some GPUs)
std::vector<float> rasterParams;
@ -349,7 +354,7 @@ int QgsNineCellFilter::processRasterGPU( const QString &source, QgsFeedback *fee
scanLine3[0] = scanLine3[xSize + 1] = mInputNodataValue;
// TODO: There is room for further optimization here: instead of replacing the buffers
// we could just replace just hthe new one (the top row) and switch the order
// we could just replace just the new one (the top row) and switch the order
// of buffer arguments in the kernell call.
errorCode = cl::enqueueWriteBuffer( scanLine1Buffer, CL_TRUE, 0,
sizeof( float ) * ( xSize + 2 ), scanLine1.get() );

View File

@ -1087,6 +1087,7 @@ QgsOptions::QgsOptions( QWidget *parent, Qt::WindowFlags fl, const QList<QgsOpti
if ( QgsOpenClUtils::available( ) )
{
mGPUEnableCheckBox->setEnabled( true );
<<<<<<< 12fa896554321892e88ca05407cf125a7ccf92c1
for ( const auto &dev : QgsOpenClUtils::devices( ) )
{
@ -1100,13 +1101,38 @@ QgsOptions::QgsOptions( QWidget *parent, Qt::WindowFlags fl, const QList<QgsOpti
connect( mOpenClDevicesCombo, qgis::overload< int >::of( &QComboBox::currentIndexChanged ), infoUpdater );
mOpenClDevicesCombo->setCurrentIndex( mOpenClDevicesCombo->findData( QgsOpenClUtils::deviceId( QgsOpenClUtils::activeDevice() ) ) );
infoUpdater( -1 );
=======
mGPUInfoLabel->setText( QStringLiteral( "OpenCL compatible GPU found on your system:<br>"
"Name: <b>%1</b><br>"
"Vendor: <b>%2</b><br>"
"Profile: <b>%3</b><br>"
"Version: <b>%4</b><br>"
"Image support: <b>%5</b><br>"
"Max image2d width: <b>%6</b><br>"
"Max image2d height: <b>%7</b><br>"
).arg( QgsOpenClUtils::deviceInfo( QgsOpenClUtils::Info::Name ),
QgsOpenClUtils::deviceInfo( QgsOpenClUtils::Info::Vendor ),
QgsOpenClUtils::deviceInfo( QgsOpenClUtils::Info::Profile ),
QgsOpenClUtils::deviceInfo( QgsOpenClUtils::Info::Version ),
QgsOpenClUtils::deviceInfo( QgsOpenClUtils::Info::ImageSupport ),
QgsOpenClUtils::deviceInfo( QgsOpenClUtils::Info::Image2dMaxWidth ),
QgsOpenClUtils::deviceInfo( QgsOpenClUtils::Info::Image2dMaxHeight )
)
);
>>>>>>> [opencl] Test with image2d
}
else
{
mGPUEnableCheckBox->setEnabled( false );
<<<<<<< 12fa896554321892e88ca05407cf125a7ccf92c1
mGPUInfoTextBrowser->setText( tr( "An OpenCL compatible device was not found on your system.<br>"
"You may need to install additional libraries in order to enable OpenCL.<br>"
"Please check your logs for further details." ) );
=======
mGPUInfoLabel->setText( QStringLiteral( "OpenCL compatible GPU was not found on your system.<br>"
"You may need to install additional libraries in order to enable OpenCL.<br>"
"Please check your logs for further details." ) );
>>>>>>> [opencl] Test with image2d
}
@ -1662,8 +1688,12 @@ void QgsOptions::saveOptions()
#ifdef HAVE_OPENCL
// OpenCL settings
QgsOpenClUtils::setEnabled( mGPUEnableCheckBox->isChecked() );
<<<<<<< 12fa896554321892e88ca05407cf125a7ccf92c1
QString preferredDevice( mOpenClDevicesCombo->currentData().toString() );
QgsOpenClUtils::storePreferredDevice( preferredDevice );
=======
>>>>>>> [opencl] Test with image2d
#endif
// Gdal skip driver list

View File

@ -359,6 +359,9 @@ bool QgsColorRampShader::shade( double value, int *returnRedValue, int *returnGr
mLUTInitialized = true;
}
if ( mLUT.empty() )
return false;
// overflow indicates that value > maximum value + DOUBLE_DIFF_THRESHOLD
// that way idx can point to the last valid item
bool overflow = false;
@ -377,10 +380,6 @@ bool QgsColorRampShader::shade( double value, int *returnRedValue, int *returnGr
overflow = true;
}
}
else if ( lutIndex < 0 )
{
return false;
}
else
{
// get initial value from LUT

View File

@ -238,7 +238,7 @@ class CORE_EXPORT QgsRasterBlock
* \param row row index
* \param column column index
* \returns true if value is no data */
bool isNoData( int row, int column )
bool isNoData( int row, int column ) const
{
return isNoData( static_cast< qgssize >( row ) * mWidth + column );
}
@ -248,7 +248,7 @@ class CORE_EXPORT QgsRasterBlock
* \param row row index
* \param column column index
* \returns true if value is no data */
bool isNoData( qgssize row, qgssize column )
bool isNoData( qgssize row, qgssize column ) const
{
return isNoData( row * static_cast< qgssize >( mWidth ) + column );
}
@ -257,7 +257,7 @@ class CORE_EXPORT QgsRasterBlock
* \brief Check if value at position is no data
* \param index data matrix index (long type in Python)
* \returns true if value is no data */
bool isNoData( qgssize index )
bool isNoData( qgssize index ) const
{
if ( !mHasNoDataValue && !mNoDataBitmap )
return false;