You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@mxnet.apache.org by GitBox <gi...@apache.org> on 2019/11/13 16:05:43 UTC
[GitHub] [incubator-mxnet] nickguletskii commented on a change in pull
request #16790: migrate cudaMemcpy to cudaMemcpyAsync+cudaStreamSynchronize
nickguletskii commented on a change in pull request #16790: migrate cudaMemcpy to cudaMemcpyAsync+cudaStreamSynchronize
URL: https://github.com/apache/incubator-mxnet/pull/16790#discussion_r345848641
##########
File path: src/operator/contrib/proposal.cu
##########
@@ -456,9 +459,10 @@ class ProposalGPUOp : public Operator{
float* workspace_proposals_ptr = NULL;
FRCNN_CUDA_CHECK(cudaMalloc(&workspace_proposals_ptr, sizeof(float) * count * 5));
Tensor<xpu, 2> workspace_proposals(workspace_proposals_ptr, Shape2(count, 5));
- FRCNN_CUDA_CHECK(cudaMemcpy(workspace_proposals.dptr_,
- &anchors[0], sizeof(float) * anchors.size(),
- cudaMemcpyHostToDevice));
+ cudaStream_t stream = mshadow::Stream<gpu>::GetStream(s);
+ FRCNN_CUDA_CHECK(cudaMemcpyAsync(workspace_proposals.dptr_,
Review comment:
Here, memcpy is being scheduled to run on `stream`, but [line 471](https://github.com/apache/incubator-mxnet/blob/743c10c3d931b84de71a3731a69c6494921dd43d/src/operator/contrib/proposal.cu#L471) runs `ProposalGridKernel` on the default context. Unless I am missing something, this introduces a potential race condition: when `stream` does not match the default stream, `ProposalGridKernel` may run before `cudaMemcpyAsync` completes.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services