Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Fixed an issue when switching streams on device_buffer and host_buffer #358

Merged
merged 5 commits into from
Apr 4, 2019
Merged

Fixed an issue when switching streams on device_buffer and host_buffer #358

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
f4a8ac5
Fixed an issue when switching streams on MLCommon::device_buffer and …
jirikraus Mar 17, 2019
d0dd791
Avoided unnecessary cudaEventRecord + cudaStreamWaitEvent
jirikraus Mar 18, 2019
8245982
Merge branch 'branch-0.7' into bug-ext-buffer_stream
jirikraus Apr 2, 2019
553e35a
Fixed CHANGELOG.md
jirikraus Apr 2, 2019
1734f37
Merge branch 'branch-0.7' into bug-ext-buffer_stream
jirikraus Apr 4, 2019
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion CHANGELOG.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,6 +21,7 @@
- PR #422: Issue in the PCA tests was solved and CI can run with driver 418
- PR #409: Add entry to gitmodules to ignore build artifacts
- PR #412: Fix for svdQR function in ml-prims
- PR #358: Fixed an issue when switching streams on MLCommon::device_buffer and MLCommon::host_buffer
- PR #434: Fixing bug in CSR tests

# cuML 0.6.0 (22 Mar 2019)
Expand Down Expand Up @@ -92,7 +93,6 @@
- PR #380: Allow arbitrary data size on ingress for numba_utils.row_matrix
- PR #385: Fix for long import cuml time in containers and fix for setup_pip


# cuML 0.5.1 (05 Feb 2019)

## Bug Fixes
Expand Down
17 changes: 17 additions & 0 deletions ml-prims/src/common/buffer_base.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -88,6 +88,23 @@ class buffer_base
size_type _size;
size_type _capacity;
value_type* _data;
void set_stream( cudaStream_t stream )
{
if ( _stream != stream )
{
cudaEvent_t event;
CUDA_CHECK( cudaEventCreateWithFlags( &event, cudaEventDisableTiming ) );
CUDA_CHECK( cudaEventRecord( event, _stream ) );
CUDA_CHECK( cudaStreamWaitEvent ( stream, event, 0 ) );
_stream = stream;
CUDA_CHECK( cudaEventDestroy( event ) );
}
}
cudaStream_t get_stream() const
{
return _stream;
}
private:
cudaStream_t _stream;
};

Expand Down
19 changes: 9 additions & 10 deletions ml-prims/src/common/device_buffer.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -64,30 +64,30 @@ class device_buffer : public buffer_base<T>
{
if ( _size > 0 )
{
_data = static_cast<value_type*>(_allocator->allocate( _capacity*sizeof(value_type), _stream ));
CUDA_CHECK( cudaStreamSynchronize( _stream ) );
_data = static_cast<value_type*>(_allocator->allocate( _capacity*sizeof(value_type), get_stream() ));
CUDA_CHECK( cudaStreamSynchronize( get_stream() ) );
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@jirikraus Sorry, I realize that this question is not related to this PR! Do we need this synchronize step? I suppose, as long as the caller adheres to the stream semantics, we should be fine without it, no?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

IIRC we already touched on this in a discussion on the PR which introduced this data structure. From my point of view the expectation of a user of host_buffer or device_buffer is that it is usable after constructions. Thus the synchronization to avoid surprises. However I agree that given we pass in a stream that surprise of some users is something we can probably live with. Btw. if you want a fully asynchronous construction you can do it with:

device_buffer tmp( allocator, stream, 0 )
tmp.resize(n,stream);

I am undecided what the better option is. Any other thoughts on this?

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

host_buffer can also potentially be directly accessed from CPU code. So, it does make sense to put a sync call there to avoid surprises.

However, after a device_buffer is constructed, the only way it can be touched is by using a cuda api or a cuda kernel, both of which adhere to stream semantics. Thus, as long as the underlying allocator adheres to the stream semantics, we could live without a sync here. What say?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yes that is a good argument. Only concern I have with that it would be confusing if host_buffer and device_buffer have different semantics.

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Agree that this can be confusing to some folks. May be we could document this difference somewhere so that our devs are aware?

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@teju85 that's not what I meant. As I wrote: "I like the guarantee that the memory is immediately valid and available to all streams. I think it probably leads to fewer bugs."

Copy link
Member

@teju85 teju85 Apr 1, 2019 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@harrism I was referring to your following statement about RMM: In fact, even though RMM takes a stream in rmmAlloc(), the returned memory is immediately available for use on any stream!

I was wishing for any custom device allocators people might want to use with device_buffer, we could have this a guarantee, atleast for the specific stream in the ctor.

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

You could have a parameter that you set with the allocator. Call cudaStreamSynchronize() (or cudaDeviceSynchronize()) if the allocator is asynchronous, and don't call it if not.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks for all the inputs. I don't think we have come to a conclusion yet. @teju85 what do you think if we move the discussion if the constructor of device_buffer should synchronize or not to an issue and move on with this pull request (which as you state initially is unrelated to the open discussion) as it is.

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Agreed. Doesn't make sense to block this PR for this discussion. Have filed issue #425 and let's migrate this discussion over there.

}
}

~device_buffer()
{
if ( nullptr != _data )
{
_allocator->deallocate( _data, _capacity*sizeof(value_type), _stream );
_allocator->deallocate( _data, _capacity*sizeof(value_type), get_stream() );
}
}

void reserve( const size_type new_capacity, cudaStream_t stream )
{
_stream = stream;
set_stream( stream );
if ( new_capacity > _capacity )
{
value_type* new_data = static_cast<value_type*>(_allocator->allocate( new_capacity*sizeof(value_type), _stream ));
value_type* new_data = static_cast<value_type*>(_allocator->allocate( new_capacity*sizeof(value_type), get_stream() ));
if ( _size > 0 ) {
CUDA_CHECK( cudaMemcpyAsync( new_data, _data, _size*sizeof(value_type), cudaMemcpyDeviceToDevice, _stream ) );
CUDA_CHECK( cudaMemcpyAsync( new_data, _data, _size*sizeof(value_type), cudaMemcpyDeviceToDevice, get_stream() ) );
}
if ( nullptr != _data ) {
_allocator->deallocate( _data, _capacity*sizeof(value_type), _stream );
_allocator->deallocate( _data, _capacity*sizeof(value_type), get_stream() );
}
_data = new_data;
_capacity = new_capacity;
Expand All @@ -102,9 +102,9 @@ class device_buffer : public buffer_base<T>

void release( cudaStream_t stream )
{
_stream = stream;
set_stream( stream );
if ( nullptr != _data ) {
_allocator->deallocate( _data, _capacity*sizeof(value_type), _stream );
_allocator->deallocate( _data, _capacity*sizeof(value_type), get_stream() );
}
_data = nullptr;
_capacity = 0;
Expand All @@ -121,7 +121,6 @@ class device_buffer : public buffer_base<T>
using buffer_base<T>::_size;
using buffer_base<T>::_capacity;
using buffer_base<T>::_data;
using buffer_base<T>::_stream;
};

} // end namespace ML
19 changes: 9 additions & 10 deletions ml-prims/src/common/host_buffer.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -65,16 +65,16 @@ class host_buffer : public buffer_base<T>
{
if ( _capacity > 0 )
{
_data = static_cast<value_type*>(_allocator->allocate( _capacity*sizeof(value_type), _stream ));
CUDA_CHECK( cudaStreamSynchronize( _stream ) );
_data = static_cast<value_type*>(_allocator->allocate( _capacity*sizeof(value_type), get_stream() ));
CUDA_CHECK( cudaStreamSynchronize( get_stream() ) );
}
}

~host_buffer()
{
if ( nullptr != _data )
{
_allocator->deallocate( _data, _capacity*sizeof(value_type), _stream );
_allocator->deallocate( _data, _capacity*sizeof(value_type), get_stream() );
}
}

Expand All @@ -90,15 +90,15 @@ class host_buffer : public buffer_base<T>

void reserve( const size_type new_capacity, cudaStream_t stream )
{
_stream = stream;
set_stream( stream );
if ( new_capacity > _capacity )
{
value_type* new_data = static_cast<value_type*>(_allocator->allocate( new_capacity*sizeof(value_type), _stream ));
value_type* new_data = static_cast<value_type*>(_allocator->allocate( new_capacity*sizeof(value_type), get_stream() ));
if ( _size > 0 ) {
CUDA_CHECK( cudaMemcpyAsync( new_data, _data, _size*sizeof(value_type), cudaMemcpyHostToHost, _stream ) );
CUDA_CHECK( cudaMemcpyAsync( new_data, _data, _size*sizeof(value_type), cudaMemcpyHostToHost, get_stream() ) );
}
if ( nullptr != _data ) {
_allocator->deallocate( _data, _capacity*sizeof(value_type), _stream );
_allocator->deallocate( _data, _capacity*sizeof(value_type), get_stream() );
}
_data = new_data;
_capacity = new_capacity;
Expand All @@ -113,9 +113,9 @@ class host_buffer : public buffer_base<T>

void release( cudaStream_t stream )
{
_stream = stream;
set_stream( stream );
if ( nullptr != _data ) {
_allocator->deallocate( _data, _capacity*sizeof(value_type), _stream );
_allocator->deallocate( _data, _capacity*sizeof(value_type), get_stream() );
}
_data = nullptr;
_capacity = 0;
Expand All @@ -132,7 +132,6 @@ class host_buffer : public buffer_base<T>
using buffer_base<T>::_size;
using buffer_base<T>::_capacity;
using buffer_base<T>::_data;
using buffer_base<T>::_stream;
};

} // end namespace ML