DPCT1097#

Message#

The function <backward function name> may require the workspace used to save intermediate results from function <forward function name>. By default, a workspace from engine_ext is selected according to the source data pointer, but this may be incorrect and cause a workspace data race. You may need to rewrite this code.

Detailed Help#

You can manually pass a dnnl::memory object generated from the forward function to the backward function.

For example, this original CUDA* code:

 1void test(cudnnHandle_t handle, cudnnTensorDescriptor_t dataTensor,
 2          cudnnTensorDescriptor_t outTensor,
 3          cudnnTensorDescriptor_t diffdataTensor,
 4          cudnnTensorDescriptor_t diffoutTensor, float *data, float *out,
 5          float *diffdata, float *diffout, float alpha, float beta,
 6          cudnnLRNDescriptor_t desc) {
 7  ...
 8  cudnnLRNCrossChannelForward(handle, desc, CUDNN_LRN_CROSS_CHANNEL_DIM1,
 9                              &alpha, dataTensor, data, &beta, outTensor, out);
10  ...
11  cudnnLRNCrossChannelBackward(handle, desc, CUDNN_LRN_CROSS_CHANNEL_DIM1,
12                               &alpha, outTensor, out, diffoutTensor, diffout,
13                               dataTensor, data, &beta, diffdataTensor,
14                               diffdata);
15  ...
16}

results in the following migrated SYCL* code:

 1void test(dpct::dnnl::engine_ext handle, dpct::dnnl::memory_desc_ext dataTensor,
 2          dpct::dnnl::memory_desc_ext outTensor,
 3          dpct::dnnl::memory_desc_ext diffdataTensor,
 4          dpct::dnnl::memory_desc_ext diffoutTensor, float *data, float *out,
 5          float *diffdata, float *diffout, float alpha, float beta,
 6          dpct::dnnl::lrn_desc desc) {
 7  ...
 8  handle.async_lrn_forward(desc, alpha, dataTensor, data, beta, outTensor, out);
 9  ...
10  /*
11  DPCT1097:0: The function "async_lrn_backward" may require the workspace used
12  to save intermediate results from function "async_lrn_forward". By default, a
13  workspace from engine_ext is selected according to the source data pointer,
14  but this may be incorrect and cause a workspace data race. You may need to
15  rewrite this code.
16  */
17  handle.async_lrn_backward(desc, alpha, outTensor, out, diffoutTensor, diffout,
18                            dataTensor, data, beta, diffdataTensor, diffdata);
19  ...
20}

which is manually adjusted to:

 1void test(dpct::dnnl::engine_ext handle, dpct::dnnl::memory_desc_ext dataTensor,
 2          dpct::dnnl::memory_desc_ext outTensor,
 3          dpct::dnnl::memory_desc_ext diffdataTensor,
 4          dpct::dnnl::memory_desc_ext diffoutTensor, float *data, float *out,
 5          float *diffdata, float *diffout, float alpha, float beta,
 6          dpct::dnnl::lrn_desc desc) {
 7  ...
 8  dnnl::memory workspace;
 9  handle.async_lrn_forward(desc, alpha, dataTensor, data, beta, outTensor, out,
10                           &workspace);
11  ...
12  handle.async_lrn_backward(desc, alpha, outTensor, out, diffoutTensor, diffout,
13                            dataTensor, data, beta, diffdataTensor, diffdata,
14                            &workspace);
15  ...
16}

Suggestions to Fix#

You may need to adjust the original code.