DPCT1027#

Message#

The call to <API name> was replaced with 0 because <reason>.

Detailed Help#

API calls from the original application, which do not have functionally compatible SYCL* API calls are replaced with 0 if SYCLomatic determines that this call removal should not affect the program logic and at the same time the return value of that call is used for error handling.

Possible reasons for replacement:

  • SYCL currently does not support setting resource limits on devices.

  • SYCL currently does not support associating USM with a specific queue.

  • SYCL currently does not support query operations on queues.

  • SYCL currently does not support capture operations on queues.

  • SYCL currently does not support configuring shared memory on devices.

  • SYCL currently does not support setting cache config on devices.

  • SYCL currently does not support registering of existing host memory for use by device. Use USM to allocate memory for use by host and device.

  • SYCL currently does not support setting flags for devices.

  • SYCL currently does not support memory access across peer devices.

  • There is no corresponding API in SYCL.

  • The call is redundant in SYCL.

Suggestions to Fix#

Verify the code correctness.

For example, this original CUDA* code:

 1__global__ void kernel() {
 2  return;
 3}
 4
 5void foo() {
 6  int err = 0;
 7  cudaLimit limit;
 8  cudaStream_t stream;
 9  unsigned int flags;
10  cudaSharedMemConfig config;
11
12  err = cudaDeviceSetLimit(limit, 0);
13  err = cudaStreamAttachMemAsync(stream, nullptr);
14  err = cudaStreamQuery(stream);
15  err = cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
16  err = cudaDeviceSetSharedMemConfig(config);
17  err = cudaDeviceSetCacheConfig(cudaFuncCachePreferNone);
18  err = cuMemHostRegister(ptr, count, flags);
19  err = cudaSetDeviceFlags(flags);
20  err = cudaDeviceEnablePeerAccess(peer_device, flags);
21  err = cuInit(0);
22  float *aptr;
23  cudaStreamAttrValue stream_attribute = {};
24  stream_attribute.accessPolicyWindow.base_ptr = aptr;
25  stream_attribute.accessPolicyWindow.num_bytes = 8 * 8 * 8;
26  stream_attribute.accessPolicyWindow.hitRatio = 1.0f;
27  stream_attribute.accessPolicyWindow.hitProp = cudaAccessPropertyPersisting;
28  stream_attribute.accessPolicyWindow.missProp = cudaAccessPropertyStreaming;
29  err = cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &stream_attribute);
30  err = cudaStreamGetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &stream_attribute);
31  err = cudaCtxResetPersistingL2Cache();
32  err = cuCtxResetPersistingL2Cache();
33  err = cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, 8 * 8 * 8);
34}

results in the following migrated SYCL code:

 1void kernel() {
 2  return;
 3}
 4
 5void foo() try {
 6  int err = 0;
 7  cudaLimit limit;
 8  dpct::queue_ptr stream;
 9  unsigned int flags;
10  int config;
11
12  /*
13  DPCT1027:0: The call to cudaDeviceSetLimit was replaced with 0 because SYCL
14  currently does not support setting resource limits on devices.
15  */
16  err = 0;
17  /*
18  DPCT1027:1: The call to cudaStreamAttachMemAsync was replaced with 0 because
19  SYCL currently does not support associating USM with a specific queue.
20  */
21  err = 0;
22  /*
23  DPCT1027:2: The call to cudaStreamQuery was replaced with 0 because SYCL
24  currently does not support query operations on queues.
25  */
26  err = 0;
27  /*
28  DPCT1027:3: The call to cudaDeviceSetSharedMemConfig was replaced with 0
29  because SYCL currently does not support configuring shared memory on devices.
30  */
31  err = 0;
32  /*
33  DPCT1027:4: The call to cudaDeviceSetCacheConfig was replaced with 0 because
34  SYCL currently does not support setting cache config on devices.
35  */
36  err = 0;
37  /*
38  DPCT1027:5: The call to cudaSetDeviceFlags was replaced with 0 because SYCL
39  currently does not support setting flags for devices.
40  */
41  err = 0;
42  /*
43  DPCT1027:6: The call to cuInit was replaced with 0 because this call is
44  redundant in SYCL.
45  */
46  err = 0;
47  float *aptr;
48  /*
49  DPCT1007:7: Migration of cudaStreamAttributeValue is not supported.
50  */
51  cudaStreamAttrValue stream_attribute = {};
52  /*
53  DPCT1027:8: The call to cudaStreamSetAttribute was removed because SYCL currently
54  does not support setting cache config on devices.
55  */
56  err = 0;
57  /*
58  DPCT1027:9: The call to cudaStreamGetAttribute was removed because SYCL currently
59  does not support setting cache config on devices.
60  */
61  err = 0;
62  /*
63  DPCT1027:10: The call to cudaCtxResetPersistingL2Cache was removed because SYCL
64  currently does not support setting cache config on devices.
65  */
66  err = 0;
67  /*
68  DPCT1027:11: The call to cuCtxResetPersistingL2Cache was removed because SYCL
69  currently does not support setting cache config on devices.
70  */
71  err = 0;
72  /*
73  DPCT1027:12: The call to cudaFuncSetAttribute was removed because SYCL currently
74  does not support corresponding setting.
75  */
76  err = 0;
77  /*
78  DPCT1027:13: The call to cudaFuncSetAttribute was removed because SYCL currently
79  does not support corresponding setting.
80  */
81  err = 0;
82}
83catch (sycl::exception const &exc) {
84  std::cerr << exc.what() << "Exception caught at file:" << __FILE__
85            << ", line:" << __LINE__ << std::endl;
86  std::exit(1);
87}

which is rewritten to:

1void foo() {
2  int err = 0;
3  cudaLimit limit;
4  dpct::queue_ptr stream;
5  unsigned int flags;
6  int config;
7  float *aptr;
8}