Skip to content

Commit

Permalink
IFastWindingNumber refactoring (#3937)
Browse files Browse the repository at this point in the history
  • Loading branch information
@Fedr
Fedr authored Jan 3, 2025
1 parent 4096696 commit c5a84c9
Show file tree
Hide file tree
Showing 6 changed files with 134 additions and 104 deletions.
32 changes: 21 additions & 11 deletions source/MRCuda/MRCudaBasic.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -17,11 +17,21 @@ namespace Cuda
/// spdlog::error the information about some CUDA error including optional filename and line number
MRCUDA_API cudaError_t logError( cudaError_t code, const char * file = nullptr, int line = 0 );

/// executes given CUDA function and checks the error code after
#define CUDA_EXEC( func ) logError( func, __FILE__ , __LINE__ )
/// evaluates given expression and logs the error if any
#define CUDA_LOGE( expr ) MR::Cuda::logError( expr, __FILE__ , __LINE__ )
// deprecated
#define CUDA_EXEC( expr ) CUDA_LOGE( expr )

/// executes given CUDA function and returns error if it fails
#define CUDA_EXEC_RETURN( func ) if ( auto code = CUDA_EXEC( func ); code != cudaError::cudaSuccess ) return code
/// evaluates given expression, logs if it fails and returns error code
#define CUDA_LOGE_RETURN( expr ) if ( auto code = CUDA_LOGE( expr ); code != cudaError::cudaSuccess ) return code
// deprecated
#define CUDA_EXEC_RETURN( expr ) CUDA_LOGE_RETURN( expr )

/// if given expression evaluates to not cudaError::cudaSuccess, then returns MR::unexpected with the error string without logging
#define CUDA_RETURN_UNEXPECTED( expr ) if ( auto code = ( expr ); code != cudaError::cudaSuccess ) return MR::unexpected( MR::Cuda::getError( code ) )

/// evaluates given expression, logs if it fails and returns MR::unexpected with the error string
#define CUDA_LOGE_RETURN_UNEXPECTED( expr ) if ( auto code = CUDA_LOGE( expr ); code != cudaError::cudaSuccess ) return MR::unexpected( MR::Cuda::getError( code ) )

template<typename T>
DynamicArray<T>::DynamicArray( size_t size )
Expand Down Expand Up @@ -49,7 +59,7 @@ inline cudaError_t DynamicArray<T>::fromVector( const std::vector<U>& vec )
static_assert ( sizeof( T ) == sizeof( U ) );
if ( auto code = resize( vec.size() ) )
return code;
return CUDA_EXEC( cudaMemcpy( data_, vec.data(), size_ * sizeof( T ), cudaMemcpyHostToDevice ) );
return CUDA_LOGE( cudaMemcpy( data_, vec.data(), size_ * sizeof( T ), cudaMemcpyHostToDevice ) );
}


Expand All @@ -58,13 +68,13 @@ inline cudaError_t DynamicArray<T>::fromBytes( const uint8_t* data, size_t numBy
{
assert( numBytes % sizeof( T ) == 0 );
resize( numBytes / sizeof( T ) );
return CUDA_EXEC( cudaMemcpy( data_, data, numBytes, cudaMemcpyHostToDevice ) );
return CUDA_LOGE( cudaMemcpy( data_, data, numBytes, cudaMemcpyHostToDevice ) );
}

template <typename T>
inline cudaError_t DynamicArray<T>::toBytes( uint8_t* data )
{
return CUDA_EXEC( cudaMemcpy( data, data_, size_ * sizeof( T ), cudaMemcpyDeviceToHost ) );
return CUDA_LOGE( cudaMemcpy( data, data_, size_ * sizeof( T ), cudaMemcpyDeviceToHost ) );
}

template<typename T>
Expand All @@ -74,14 +84,14 @@ cudaError_t DynamicArray<T>::resize( size_t size )
return cudaSuccess;
if ( size_ != 0 )
{
if ( auto code = CUDA_EXEC( cudaFree( data_ ) ) )
if ( auto code = CUDA_LOGE( cudaFree( data_ ) ) )
return code;
}

size_ = size;
if ( size_ != 0 )
{
if ( auto code = CUDA_EXEC( cudaMalloc( ( void** )&data_, size_ * sizeof( T ) ) ) )
if ( auto code = CUDA_LOGE( cudaMalloc( ( void** )&data_, size_ * sizeof( T ) ) ) )
return code;
}
return cudaSuccess;
Expand All @@ -93,14 +103,14 @@ cudaError_t DynamicArray<T>::toVector( std::vector<U>& vec ) const
{
static_assert ( sizeof( T ) == sizeof( U ) );
vec.resize( size_ );
return CUDA_EXEC( cudaMemcpy( vec.data(), data_, size_ * sizeof( T ), cudaMemcpyDeviceToHost ) );
return CUDA_LOGE( cudaMemcpy( vec.data(), data_, size_ * sizeof( T ), cudaMemcpyDeviceToHost ) );
}

inline cudaError_t setToZero( DynamicArrayF& devArray )
{
if ( devArray.size() == 0 )
return cudaSuccess;
return CUDA_EXEC( cudaMemset( devArray.data(), 0, devArray.size() * sizeof( float ) ) );
return CUDA_LOGE( cudaMemset( devArray.data(), 0, devArray.size() * sizeof( float ) ) );
}

} // namespace Cuda
Expand Down
152 changes: 85 additions & 67 deletions source/MRCuda/MRCudaFastWindingNumber.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -26,90 +26,109 @@ FastWindingNumber::FastWindingNumber( const Mesh& mesh ) : mesh_( mesh )
{
}

bool FastWindingNumber::prepareData_( ProgressCallback cb )
Expected<void> FastWindingNumber::prepareData_( ProgressCallback cb )
{
CUDA_EXEC( cudaSetDevice( 0 ) );
CUDA_LOGE_RETURN_UNEXPECTED( cudaSetDevice( 0 ) );
if ( data_ )
return reportProgress( cb, 1.0f );
{
if ( !reportProgress( cb, 1.0f ) )
return unexpectedOperationCanceled();
return {};
}
MR_TIMER

auto data = std::make_shared<FastWindingNumberData>();

if ( !reportProgress( cb, 0.0f ) )
return false;
return unexpectedOperationCanceled();

data->cudaMeshPoints.fromVector( mesh_.points.vec_ );
CUDA_LOGE_RETURN_UNEXPECTED( data->cudaMeshPoints.fromVector( mesh_.points.vec_ ) );
if ( !reportProgress( cb, 0.1f ) )
return false;
return unexpectedOperationCanceled();

data->cudaFaces.fromVector( mesh_.topology.getTriangulation().vec_ );
CUDA_LOGE_RETURN_UNEXPECTED( data->cudaFaces.fromVector( mesh_.topology.getTriangulation().vec_ ) );
if ( !reportProgress( cb, 0.3f ) )
return false;
return unexpectedOperationCanceled();

const AABBTree& tree = mesh_.getAABBTree();
if ( !reportProgress( cb, 0.5f ) )
return false;
return unexpectedOperationCanceled();

const auto& nodes = tree.nodes();
data->cudaNodes.fromVector( nodes.vec_ );
CUDA_LOGE_RETURN_UNEXPECTED( data->cudaNodes.fromVector( nodes.vec_ ) );
if ( !reportProgress( cb, 0.6f ) )
return false;
return unexpectedOperationCanceled();

data->dipoles.fromVector( mesh_.getDipoles().vec_ );
CUDA_LOGE_RETURN_UNEXPECTED( data->dipoles.fromVector( mesh_.getDipoles().vec_ ) );
if ( !reportProgress( cb, 1.0f ) )
return false;
return unexpectedOperationCanceled();

data_ = std::move( data );
return true;
return {};
}

void FastWindingNumber::calcFromVector( std::vector<float>& res, const std::vector<Vector3f>& points, float beta, FaceId skipFace )
Expected<void> FastWindingNumber::calcFromVector( std::vector<float>& res, const std::vector<Vector3f>& points, float beta, FaceId skipFace, const ProgressCallback& cb )
{
MR_TIMER
prepareData_( {} );

const size_t size = points.size();
res.resize( size );
data_->cudaPoints.fromVector( points );
DynamicArrayF cudaResult( size );

fastWindingNumberFromVector( data_->cudaPoints.data(), data_->dipoles.data(), data_->cudaNodes.data(), data_->cudaMeshPoints.data(), data_->cudaFaces.data(), cudaResult.data(), beta, int( skipFace ), size );
CUDA_EXEC( cudaGetLastError() );

CUDA_EXEC( cudaResult.toVector( res ) );
return prepareData_( subprogress( cb, 0.0, 0.5f ) ).and_then( [&]() -> Expected<void>
{
const size_t size = points.size();
res.resize( size );
CUDA_LOGE_RETURN_UNEXPECTED( data_->cudaPoints.fromVector( points ) );
if ( !reportProgress( cb, 0.6f ) )
return unexpectedOperationCanceled();

DynamicArrayF cudaResult;
CUDA_LOGE_RETURN_UNEXPECTED( cudaResult.resize( size ) );
fastWindingNumberFromVector( data_->cudaPoints.data(), data_->dipoles.data(), data_->cudaNodes.data(), data_->cudaMeshPoints.data(), data_->cudaFaces.data(), cudaResult.data(), beta, int( skipFace ), size );
CUDA_LOGE_RETURN_UNEXPECTED( cudaGetLastError() );
if ( !reportProgress( cb, 0.7f ) )
return unexpectedOperationCanceled();

CUDA_LOGE_RETURN_UNEXPECTED( cudaResult.toVector( res ) );
if ( !reportProgress( cb, 1.0f ) )
return unexpectedOperationCanceled();
return {};
} );
}

bool FastWindingNumber::calcSelfIntersections( FaceBitSet& res, float beta, ProgressCallback cb )
Expected<void> FastWindingNumber::calcSelfIntersections( FaceBitSet& res, float beta, const ProgressCallback& cb )
{
MR_TIMER
if ( !prepareData_( subprogress( cb, 0.0f, 0.5f ) ) )
return false;

const size_t size = mesh_.topology.faceSize();
DynamicArrayF cudaResult( size );

fastWindingNumberFromMesh(data_->dipoles.data(), data_->cudaNodes.data(), data_->cudaMeshPoints.data(), data_->cudaFaces.data(), cudaResult.data(), beta, size);
if ( CUDA_EXEC( cudaGetLastError() ) )
return false;

std::vector<float> wns;
if ( CUDA_EXEC( cudaResult.toVector( wns ) ) )
return false;
if ( !reportProgress( cb, 0.9f ) )
return false;

res.resize( size );
return BitSetParallelForAll( res, [&] (FaceId f)
return prepareData_( subprogress( cb, 0.0, 0.5f ) ).and_then( [&]() -> Expected<void>
{
if ( wns[f] < 0 || wns[f] > 1 )
res.set( f );
}, subprogress( cb, 0.9f, 1.0f ) );
const size_t size = mesh_.topology.faceSize();
DynamicArrayF cudaResult;
CUDA_LOGE_RETURN_UNEXPECTED( cudaResult.resize( size ) );
if ( !reportProgress( cb, 0.6f ) )
return unexpectedOperationCanceled();

fastWindingNumberFromMesh(data_->dipoles.data(), data_->cudaNodes.data(), data_->cudaMeshPoints.data(), data_->cudaFaces.data(), cudaResult.data(), beta, size);
CUDA_LOGE_RETURN_UNEXPECTED( cudaGetLastError() );
if ( !reportProgress( cb, 0.7f ) )
return unexpectedOperationCanceled();

std::vector<float> wns;
CUDA_LOGE_RETURN_UNEXPECTED( cudaResult.toVector( wns ) );
if ( !reportProgress( cb, 0.9f ) )
return unexpectedOperationCanceled();

res.resize( size );
if ( !BitSetParallelForAll( res, [&] (FaceId f)
{
if ( wns[f] < 0 || wns[f] > 1 )
res.set( f );
}, subprogress( cb, 0.9f, 1.0f ) ) )
return unexpectedOperationCanceled();
return {};
} );
}

Expected<void> FastWindingNumber::calcFromGrid( std::vector<float>& res, const Vector3i& dims, const AffineXf3f& gridToMeshXf, float beta, ProgressCallback cb )
Expected<void> FastWindingNumber::calcFromGrid( std::vector<float>& res, const Vector3i& dims, const AffineXf3f& gridToMeshXf, float beta, const ProgressCallback& cb )
{
MR_TIMER
prepareData_( {} );
if ( auto maybe = prepareData_( subprogress( cb, 0.0, 0.5f ) ); !maybe )
return unexpected( std::move( maybe.error() ) );

const auto getCudaMatrix = [] ( const AffineXf3f& xf )
{
Expand All @@ -124,22 +143,21 @@ Expected<void> FastWindingNumber::calcFromGrid( std::vector<float>& res, const V

const Matrix4 cudaGridToMeshXf = ( gridToMeshXf == AffineXf3f{} ) ? Matrix4{} : getCudaMatrix( gridToMeshXf );
const size_t size = size_t( dims.x ) * dims.y * dims.z;
DynamicArrayF cudaResult( size );
if ( !reportProgress( cb, 0.0f ) )
DynamicArrayF cudaResult;
CUDA_LOGE_RETURN_UNEXPECTED( cudaResult.resize( size ) );
if ( !reportProgress( cb, 0.6f ) )
return unexpectedOperationCanceled();

fastWindingNumberFromGrid(
int3{ dims.x, dims.y, dims.z },
cudaGridToMeshXf,
data_->dipoles.data(), data_->cudaNodes.data(), data_->cudaMeshPoints.data(), data_->cudaFaces.data(),
cudaResult.data(), beta );

if ( auto code = CUDA_EXEC( cudaGetLastError() ) )
return unexpected( Cuda::getError( code ) );

if ( auto code = cudaResult.toVector( res ) )
return unexpected( Cuda::getError( code ) );
CUDA_LOGE_RETURN_UNEXPECTED( cudaGetLastError() );
if ( !reportProgress( cb, 0.7f ) )
return unexpectedOperationCanceled();

CUDA_LOGE_RETURN_UNEXPECTED( cudaResult.toVector( res ) );
if ( !reportProgress( cb, 1.0f ) )
return unexpectedOperationCanceled();
return {};
Expand All @@ -148,7 +166,8 @@ Expected<void> FastWindingNumber::calcFromGrid( std::vector<float>& res, const V
Expected<void> FastWindingNumber::calcFromGridWithDistances( std::vector<float>& res, const Vector3i& dims, const AffineXf3f& gridToMeshXf, const DistanceToMeshOptions& options, const ProgressCallback& cb )
{
MR_TIMER
prepareData_( {} );
if ( auto maybe = prepareData_( subprogress( cb, 0.0, 0.5f ) ); !maybe )
return unexpected( std::move( maybe.error() ) );

const auto getCudaMatrix = [] ( const AffineXf3f& xf )
{
Expand All @@ -163,22 +182,21 @@ Expected<void> FastWindingNumber::calcFromGridWithDistances( std::vector<float>&

const Matrix4 cudaGridToMeshXf = ( gridToMeshXf == AffineXf3f{} ) ? Matrix4{} : getCudaMatrix( gridToMeshXf );
const size_t size = size_t( dims.x ) * dims.y * dims.z;
DynamicArrayF cudaResult( size );
if ( !reportProgress( cb, 0.0f ) )
DynamicArrayF cudaResult;
CUDA_LOGE_RETURN_UNEXPECTED( cudaResult.resize( size ) );
if ( !reportProgress( cb, 0.6f ) )
return unexpectedOperationCanceled();

signedDistance(
int3{ dims.x, dims.y, dims.z },
cudaGridToMeshXf,
data_->dipoles.data(), data_->cudaNodes.data(), data_->cudaMeshPoints.data(), data_->cudaFaces.data(),
cudaResult.data(), options );
CUDA_LOGE_RETURN_UNEXPECTED( cudaGetLastError() );
if ( !reportProgress( cb, 0.7f ) )
return unexpectedOperationCanceled();

if ( auto code = CUDA_EXEC( cudaGetLastError() ) )
return unexpected( Cuda::getError( code ) );

if ( auto code = cudaResult.toVector( res ) )
return unexpected( Cuda::getError( code ) );

CUDA_LOGE_RETURN_UNEXPECTED( cudaResult.toVector( res ) );
if ( !reportProgress( cb, 1.0f ) )
return unexpectedOperationCanceled();
return {};
Expand Down
8 changes: 4 additions & 4 deletions source/MRCuda/MRCudaFastWindingNumber.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -24,13 +24,13 @@ class MRCUDA_CLASS FastWindingNumber : public IFastWindingNumber
MRCUDA_API FastWindingNumber( const Mesh& mesh );

// see methods' descriptions in IFastWindingNumber
MRCUDA_API void calcFromVector( std::vector<float>& res, const std::vector<Vector3f>& points, float beta, FaceId skipFace = {} ) override;
MRCUDA_API bool calcSelfIntersections( FaceBitSet& res, float beta, ProgressCallback cb ) override;
MRCUDA_API Expected<void> calcFromGrid( std::vector<float>& res, const Vector3i& dims, const AffineXf3f& gridToMeshXf, float beta, ProgressCallback cb ) override;
MRCUDA_API Expected<void> calcFromVector( std::vector<float>& res, const std::vector<Vector3f>& points, float beta, FaceId skipFace, const ProgressCallback& cb ) override;
MRCUDA_API Expected<void> calcSelfIntersections( FaceBitSet& res, float beta, const ProgressCallback& cb ) override;
MRCUDA_API Expected<void> calcFromGrid( std::vector<float>& res, const Vector3i& dims, const AffineXf3f& gridToMeshXf, float beta, const ProgressCallback& cb ) override;
MRCUDA_API Expected<void> calcFromGridWithDistances( std::vector<float>& res, const Vector3i& dims, const AffineXf3f& gridToMeshXf, const DistanceToMeshOptions& options, const ProgressCallback& cb ) override;

private:
bool prepareData_( ProgressCallback cb );
Expected<void> prepareData_( ProgressCallback cb );
};

} // namespace Cuda
Expand Down
13 changes: 5 additions & 8 deletions source/MRCuda/MRCudaTest.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -11,19 +11,16 @@ namespace Cuda
{
Expected<void> negatePicture( Image& image )
{
if ( auto code = CUDA_EXEC( cudaSetDevice( 0 ) ) )
return unexpected( Cuda::getError( code ) );
CUDA_LOGE_RETURN_UNEXPECTED( cudaSetDevice( 0 ) );

DynamicArray<Cuda::Color> cudaArray;
if ( auto code = cudaArray.fromVector( image.pixels ) )
return unexpected( Cuda::getError( code ) );
CUDA_LOGE_RETURN_UNEXPECTED( cudaArray.fromVector( image.pixels ) );

negatePictureKernel( cudaArray );
if ( auto code = CUDA_EXEC( cudaGetLastError() ) )
return unexpected( Cuda::getError( code ) );
CUDA_LOGE_RETURN_UNEXPECTED( cudaGetLastError() );

CUDA_LOGE_RETURN_UNEXPECTED( cudaArray.toVector( image.pixels ) );

if ( auto code = cudaArray.toVector( image.pixels ) )
return unexpected( Cuda::getError( code ) );
return {};
}

Expand Down
20 changes: 13 additions & 7 deletions source/MRMesh/MRFastWindingNumber.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -25,27 +25,33 @@ inline float FastWindingNumber::calc_( const Vector3f & q, float beta, FaceId sk
return calcFastWindingNumber( dipoles_, tree_, mesh_, q, beta, skipFace );
}

void FastWindingNumber::calcFromVector( std::vector<float>& res, const std::vector<Vector3f>& points, float beta, FaceId skipFace )
Expected<void> FastWindingNumber::calcFromVector( std::vector<float>& res, const std::vector<Vector3f>& points, float beta, FaceId skipFace, const ProgressCallback& cb )
{
MR_TIMER
res.resize( points.size() );
ParallelFor( points, [&]( size_t i )
if ( !ParallelFor( points, [&]( size_t i )
{
res[i] = calc_( points[i], beta, skipFace );
} );
}, cb ) )
return unexpectedOperationCanceled();
return {};
}

bool FastWindingNumber::calcSelfIntersections( FaceBitSet& res, float beta, ProgressCallback cb )
Expected<void> FastWindingNumber::calcSelfIntersections( FaceBitSet& res, float beta, const ProgressCallback& cb )
{
MR_TIMER
res.resize( mesh_.topology.faceSize() );
return BitSetParallelFor( mesh_.topology.getValidFaces(), [&] ( FaceId f )
if ( !BitSetParallelFor( mesh_.topology.getValidFaces(), [&] ( FaceId f )
{
auto wn = calc_( mesh_.triCenter( f ), beta, f );
if ( wn < 0 || wn > 1 )
res.set( f );
}, cb );
}, cb ) )
return unexpectedOperationCanceled();
return {};
}

Expected<void> FastWindingNumber::calcFromGrid( std::vector<float>& res, const Vector3i& dims, const AffineXf3f& gridToMeshXf, float beta, ProgressCallback cb )
Expected<void> FastWindingNumber::calcFromGrid( std::vector<float>& res, const Vector3i& dims, const AffineXf3f& gridToMeshXf, float beta, const ProgressCallback& cb )
{
MR_TIMER

Expand Down
Loading

0 comments on commit c5a84c9

Please sign in to comment.