DPCT1097

メッセージ

関数 <backward function name> には、関数 <forward function name> からの中間結果を保存するためにワークスペースが必要になる場合があります。デフォルトでは、engine_ext のワークスペースがソースデータのポインターに従って選択されますが、これは正しくない可能性があり、ワークスペースのデータ競合を引き起こす可能性があります。コードを修正する必要があります。

詳細な説明

フォワード関数で生成された dnnl::memory オブジェクトを、バックエンド関数に手動で渡すことができます。

例えば、以下のオリジナル CUDA* コードについて考えてみます。

 void test(cudnnHandle_t handle, cudnnTensorDescriptor_t dataTensor, 
  cudnnTensorDescriptor_t outTensor, 
  cudnnTensorDescriptor_t diffdataTensor, 
  cudnnTensorDescriptor_t diffoutTensor, float *data, float *out, 
  float *diffdata, float *diffout, float alpha, float beta, 
  cudnnLRNDescriptor_t desc) { 
  ...
  cudnnLRNCrossChannelForward(handle, desc, CUDNN_LRN_CROSS_CHANNEL_DIM1, 
  &alpha, dataTensor, data, &beta, outTensor, out); 
 ... 
 cudnnLRNCrossChannelBackward(handle, desc, CUDNN_LRN_CROSS_CHANNEL_DIM1, 
 &alpha, outTensor, out, diffoutTensor, diffout, 
 dataTensor, data, &beta, diffdataTensor, 
 diffdata); 
 ... 
}

このコードは、以下の SYCL* コードに移行されます。

  void test(dpct::dnnl::engine_ext handle, dpct::dnnl::memory_desc_ext dataTensor, 
   dpct::dnnl::memory_desc_ext outTensor, 
   dpct::dnnl::memory_desc_ext diffdataTensor, 
   dpct::dnnl::memory_desc_ext diffoutTensor, float *data, float *out, 
   float *diffdata, float *diffout, float alpha, float beta, 
   dpct::dnnl::lrn_desc desc) { 
   ... 
   handle.async_lrn_forward(desc, alpha, dataTensor, data, beta, outTensor, out); 
   ... 
  /* 
  DPCT1097:0: The function "async_lrn_backward" may require the workspace used 
  to save intermediate results from function "async_lrn_forward".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.
  */ 
  handle.async_lrn_backward(desc, alpha, outTensor, out, diffoutTensor, diffout, 
  dataTensor, data, beta, diffdataTensor, diffdata); 
  ... 
}

このコードを以下のように手動で調整します。

 void test(dpct::dnnl::engine_ext handle, dpct::dnnl::memory_desc_ext dataTensor, 
  dpct::dnnl::memory_desc_ext outTensor, 
  dpct::dnnl::memory_desc_ext diffdataTensor, 
  dpct::dnnl::memory_desc_ext diffoutTensor, float *data, float *out, 
  float *diffdata, float *diffout, float alpha, float beta, 
  dpct::dnnl::lrn_desc desc) { 
  ... 
  dnnl::memory workspace; 
  handle.async_lrn_forward(desc, alpha, dataTensor, data, beta, outTensor, out, 
 &workspace); 
 ... 
 handle.async_lrn_backward(desc, alpha, outTensor, out, diffoutTensor, diffout, 
 dataTensor, data, beta, diffdataTensor, diffdata, 
 &workspace); 
 ... 
}

修正方法の提案

オリジナルコードを調整する必要があります。