[SYCLomatic][DOC] 2025.1 Options & Warnings (#2697)

Author: lapel22

docs/_include_files/options_def.rst

@@ -27,8 +27,6 @@ to be migrated. Default: the value of ``--in-root``.
 
 
 
-
-
 .. _opt-assume-nd-range-dim:
 
 ``--assume-nd-range-dim=<value>``
@@ -587,7 +585,8 @@ Specify the type of migration report. Values are:
 
 .. _desc-rule-file:
 
-Specify the rule file path that contains rules used for migration.
+Specify the rule file for migration. Also, reference the predefined rules in the
+``extensions``  directory in the root folder of the tool.
 
 .. _end-rule-file:
 
@@ -612,7 +611,7 @@ Stop migration and generation of reports if parsing errors happened. Default: ``
 .. _desc-suppress-warnings:
 
 A comma-separated list of migration warnings to suppress. Valid warning IDs
-range from 1000 to 1132. Hyphen-separated ranges are also allowed. For
+range from 1000 to 1136. Hyphen-separated ranges are also allowed. For
 example: ``-suppress-warnings=1000-1010,1011``.
 
 .. _end-suppress-warnings:
@@ -676,8 +675,10 @@ By default, experimental features will not be used in migrated code.
 
 The values are:
 
-- ``=bfloat16_math_functions``: Experimental extension that allows use of bfloat16 math functions. `See more details <https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_bfloat16_math_functions.asciidoc>`__.
-- ``=bindless_images``: Experimental extension that allows use of bindless images APIs. `See more details <https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc>`__.
+- ``=bfloat16_math_functions``: Experimental extension that allows use of bfloat16 math
+  functions. `See more details <https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_bfloat16_math_functions.asciidoc>`__.
+- ``=bindless_images``: Experimental extension that allows use of bindless images APIs.
+  `See more details <https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc>`__.
 - ``=dpl-experimental-api``: Experimental extension that allows use of experimental
   oneDPL APIs. `See more details <https://github.com/oneapi-src/oneDPL/tree/main/include/oneapi/dpl/pstl/experimental>`__.
 - ``=free-function-queries``: Experimental extension that allows getting
@@ -689,16 +690,29 @@ The values are:
   group work-items. See more details in ``dpct::experimental::logical_group`` in header file ``util.hpp``.
 - ``=masked-sub-group-operation``: Experimental helper function used to execute
   sub-group operation with mask. See more details in ``dpct::experimental::select_from_sub_group``, ``dpct::experimental::shift_sub_group_left``, ``dpct::experimental::shift_sub_group_right`` and ``dpct::experimental::shift_sub_group_right`` in header file ``util.hpp``.
-- ``=matrix``: Experimental extension that allows use of matrix extension like class ``joint_matrix``. `See more details <https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_matrix/sycl_ext_oneapi_matrix.asciidoc>`__.
-- ``=nd_range_barrier``: DEPRECATED. Experimental helper function used to help cross-group synchronization during migration. Please use the following option instead: ``--use-experimental-features=root-group``
+- ``=matrix``: Experimental extension that allows use of matrix extension like class
+  ``joint_matrix``. `See more details <https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_matrix/sycl_ext_oneapi_matrix.asciidoc>`__.
+- ``=nd_range_barrier``: DEPRECATED. Experimental helper function used to help cross-group
+  synchronization during migration. Please use the following option instead: ``--use-experimental-features=root-group``
 - ``=root-group``: Experimental extension that allows use of root group class and relative API.
 - ``=graph``: Experimental extension that allows use of SYCL Graph APIs.
-- ``=occupancy-calculation``: Experimental helper function used to calculate occupancy. See more details in ``dpct::experimental::calculate_max_active_wg_per_xecore`` and ``dpct::experimental::calculate_max_potential_wg`` in header file ``util.hpp``.
+- ``=occupancy-calculation``: Experimental helper function used to calculate occupancy. See more
+  details in ``dpct::experimental::calculate_max_active_wg_per_xecore`` and ``dpct::experimental::calculate_max_potential_wg`` in header file ``util.hpp``.
 - ``=user-defined-reductions``: Experimental extension that allows user-defined
   reductions. `See more details <https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_user_defined_reductions.asciidoc>`__.
-- ``=non-uniform-groups``: Experimental extension that allows use of non-uniform groups. `See more details <https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_non_uniform_groups.asciidoc>`__.
-- ``=device_global``: Experimental extension that allows device scoped memory allocations into SYCL that can
+- ``=non-uniform-groups``: Experimental extension that allows use of non-uniform groups.
+  `See more details <https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_non_uniform_groups.asciidoc>`__.
+- ``=device_global``: Experimental extension that allows device scoped memory allocations into
+  SYCL that can
   be accessed within a kernel using syntax similar to C++ global variables. `See more details <https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc>`__.
+- ``=virtual_mem``: Experimental extension that allows for mapping of an address range
+  onto multiple allocations of physical memory.
+- ``=in_order_queue_events``: Experimental extension that allows getting the event from
+  the last command submission into the queue and setting an external event as an implicit
+  dependence on the next command submitted to the queue.
+- ``=non-standard-sycl-builtins``: Experimental extension that allows use of non standard
+  SYCL builtin functions.
+- ``=prefetch``: Experimental extension that allows use of SYCL prefetch APIs.
 - ``=all``: Enable all experimental extensions listed in this option.
 
 .. _end-use-experimental-features:
@@ -882,6 +896,7 @@ Intercept build tool to generate a compilation database.
 EXPERIMENTAL: Migrate build script(s).
 
 - ``=CMake``: Migrate the CMake file(s).
+- ``=Python``: Migrate the Python build script file(s) of PyTorch based project.
 
 .. _end-migrate-build-script:
 

docs/dev_guide/reference/diagnostic_ref/dpct1009.rst

@@ -8,9 +8,8 @@ Message
 
 .. _msg-1009-start:
 
-SYCL uses exceptions to report errors and does not use the error codes. The
-original code was commented out and a warning string was inserted. You need to
-rewrite this code.
+SYCL reports errors using exceptions and does not use error codes. Please replace the
+\"get_error_string_dummy(...)\" with a real error-handling function.
 
 .. _msg-1009-end:
 
@@ -46,29 +45,34 @@ results in the following migrated SYCL code:
 .. code-block:: cpp
    :linenos:
 
-   void foo() {
+   void foo() try {
      float *f;
-     int err = (f = (float *)sycl::malloc_device(4, dpct::get_default_queue()), 0);
+     dpct::err0 err = DPCT_CHECK_ERROR(
+         f = (float *)sycl::malloc_device(4, dpct::get_in_order_queue()));
      /*
      DPCT1009:1: SYCL uses exceptions to report errors and does not use the error
      codes. The original code was commented out and a warning string was inserted.
      You need to rewrite this code.
      */
-     printf("%s\n",
-            "cudaGetErrorString is not supported" /*cudaGetErrorString(err)*/);
+     printf("%s\n", dpct::get_error_string_dummy(err));
+   }
+   catch (sycl::exception const &exc) {
+     std::cerr <<exc.what() << "Exception caught at file." << __FILE__
+               << ", line:" << __LINE__ << std:endl;
+     std::exit(1);
    }
 
 .. _example_dpct_end:
 
-which is rewritten to:
+which needs to be rewritten to:
 
 .. code-block:: cpp
    :linenos:
 
    void foo() {
      float *f;
      try {
-       f = (float *)sycl::malloc_device(4, dpct::get_default_queue())
+       f = (float *)sycl::malloc_device(4, dpct::get_in_order_queue())
      } catch (sycl::exception const &e) {
        std::cerr << e.what() << std::endl;
      }

docs/dev_guide/reference/diagnostic_ref/dpct1010.rst

@@ -8,23 +8,23 @@ Message
 
 .. _msg-1010-start:
 
-SYCL uses exceptions to report errors and does not use the error codes. The call
-was replaced with 0. You need to rewrite this code.
+SYCL uses exceptions to report errors and does not
+use the error codes. The <function name> function
+call was replaced with 0. You need to rewrite this
+code.
 
 .. _msg-1010-end:
 
 Detailed Help
 -------------
 
-SYCL\* uses exceptions to report errors and does not use error codes. The original
-code tries to get an error code, while SYCL does not require such functionality.
-The call was replaced with 0.
+SYCL\* does not support launching kernel in device
+code. The user needs to merge the parent kernel and
+child kernel together.
 
 Suggestions to Fix
 ------------------
 
-You may need to rewrite this code.
-
 For example, this original CUDA\* code:
 
 .. code-block:: cpp
@@ -51,21 +51,23 @@ results in the following migrated SYCL code:
    }
 
    void foo() {
-     dpct::get_default_queue().parallel_for(
+     dpct::get_in_order_queue().parallel_for(
          sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
          [=](sycl::nd_item<3> item_ct1) {
            kernel();
          });
      dpct::get_current_device().queues_wait_and_throw();
      /*
-     DPCT1010:0: SYCL uses exceptions to report errors and does not use the error
-     codes. The call was replaced with 0. You need to rewrite this code.
+     DPCT1010:0: SYCL uses exceptions to report
+     errors and does not use the error
+     codes. The cudaGetLastError function call was
+     replaced with 0. You need to rewrite this code.
      */
-     int err = 0;
+     dpct::err0 err = 0;
      printf("%d\n", err);
    }
 
-which is rewritten to:
+which needs to be rewritten to:
 
 .. code-block:: cpp
    :linenos:
@@ -76,7 +78,7 @@ which is rewritten to:
 
    void foo() {
      try {
-       dpct::get_default_queue().parallel_for(
+       dpct::get_in_order_queue().parallel_for(
            sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
            [=](sycl::nd_item<3> item_ct1) {
              kernel();
@@ -86,3 +88,4 @@ which is rewritten to:
        std::cerr << e.what() << std::endl;
      }
    }
+

docs/dev_guide/reference/diagnostic_ref/dpct1014.rst

@@ -8,23 +8,28 @@ Message
 
 .. _msg-1014-start:
 
-The flag and priority options are not supported for SYCL queues. The output
+The flag and priority options are not supported
+for SYCL queues. The output
 parameter(s) are set to 0.
 
+Or
+
+The <CUDA event property name> is not supported
+for SYCL event. The output parameter(s) is set to
+0.
+
 .. _msg-1014-end:
 
 Detailed Help
 -------------
 
-The flag and priority options are not supported in the SYCL\* queue. You may need
+The flag, priority options, or CUDA event property
+is not supported in the SYCL* queue. You may need
 to rewrite the generated code.
 
-
 Suggestions to Fix
 ------------------
 
-Review the logic and adjust it.
-
 For example, this original CUDA\* code:
 
 .. code-block:: cpp
@@ -35,13 +40,19 @@ For example, this original CUDA\* code:
      cudaStreamGetFlags(old_stream, &flag);
      cudaStreamCreateWithFlags(new_stream, flag);
    }
+   unsigned bar(bool f) {
+     unsigned flags = CU_EVENT_DISABLE_TIMING;
+     if (f)
+       flags |= CU_EVENT_BLOCKING_SYNC;
+     return flags;
+   }
 
 results in the following migrated SYCL code:
 
 .. code-block:: cpp
    :linenos:
 
-   void foo(sycl::queue *old_stream, sycl::queue **new_stream) {
+   void foo(dpct::queue_ptr old_stream, dpct::queue_ptr *new_stream) {
      unsigned int flag;
      /*
      DPCT1014:0: The flag and priority options are not supported for SYCL queues.
@@ -53,13 +64,32 @@ results in the following migrated SYCL code:
      */
      *(new_stream) = dpct::get_current_device().create_queue();
    }
-
-which is rewritten to:
+   unsigned bar(bool f) {
+     /*
+     DPCT1014:2: The CU_EVENT_DISABLE_TIMING is
+     not supported for SYCL event. The
+     output parameter(s) is set to 0.
+     */
+     unsigned flags = 0;
+     if (f)
+       /*
+       DPCT1014:3: The CU_EVENT_BLOCKING_SYNC is
+       not supported for SYCL event. The
+       output parameter(s) is set to 0.
+       */
+       flags |= 0;
+   } return flags;
+
+which needs to be rewritten to:
 
 .. code-block:: cpp
    :linenos:
 
-   void foo(sycl::queue *old_stream, sycl::queue **new_stream) {
+   void foo(dpct::queue_ptr old_stream, dpct::queue_ptr *new_stream) {
      *(new_stream) = dpct::get_current_device().create_queue();
    }
 
+   unsigned bar(bool f){
+     return 0;
+   }
+

docs/dev_guide/reference/diagnostic_ref/dpct1083.rst

@@ -8,22 +8,24 @@ Message
 
 .. _msg-1083-start:
 
-The size of *<placeholder>* in the migrated code may be different
+The size of local memory in the migrated code may be different
 from the original code. Check that the allocated memory size in the migrated code
 is correct.
 
+Or
+
+The size of <type name> in the migrated code may be different from the original code. You may need to adjust the code.
+
 .. _msg-1083-end:
 
 Detailed Help
 -------------
 
 Some types have a different size in the migrated code than in the original code,
-for example ``sycl::float3`` compared to ``float3``. Check if the allocated size
-of memory is correct in the migrated code.
+for example ``sycl::float3`` compared to ``float3``. Check if the memory size is correct in the migrated code.
 
-In the example below, ``3*sizeof(float)`` is used to represent the size of
-``float3`` in the original code. In the migrated code the size of ``sycl::float3``
-is different, so the allocated size needs adjustment.
+Suggestions to Fix
+------------------
 
 For example, this original CUDA\* code:
 
@@ -40,6 +42,13 @@ For example, this original CUDA\* code:
        kernel<<<1, 1, shared_size>>>();
      }
 
+     void bar(char *data_handle, int width, int height) {
+       std::vector<uchar3> data;
+       data.assign(reinterpret_cast<uchar3 *>(data_handle),
+                  reinterpret_cast<uchar3 *>(data_handle) + width * height);
+       ...
+     }
+
 results in the following migrated SYCL\* code:
 
 .. code-block:: cpp
@@ -64,12 +73,23 @@ results in the following migrated SYCL\* code:
          cgh.parallel_for(
              sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
              [=](sycl::nd_item<3> item_ct1) {
-               kernel(dpct_local_acc_ct1.get_pointer());
+               kernel(dpct_local_acc_ct1.get__multi_ptr<sycl::access::decorated::no>()
+                          .get());
              });
        });
      }
 
-which is manually adjusted to:
+     void bar(char *data_handle, int width, int height) {
+       std::vector<sycl::uchar3> data;
+       /*
+       DPCT1083:1: The size of uchar3 in the migrated code may be different from the original code. You may need to adjust the code.
+       */
+       data.assign (reinterpret_cast<sycl::uchar3 *>(data_handle),
+                    reinterpret_cast<sycl::uchar3 *>(data_handle) + width * height);
+       ...
+     }
+
+which needs to be rewritten to:
 
 .. code-block:: cpp
    :linenos:
@@ -88,13 +108,18 @@ which is manually adjusted to:
          cgh.parallel_for(
              sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
              [=](sycl::nd_item<3> item_ct1) {
-               kernel(dpct_local_acc_ct1.get_pointer());
+               kernel(dpct_local_acc_ct1.get_multi_ptr<sycl::access::decorated::no>()
+                   .get());
              });
        });
      }
-     
-Suggestions to Fix
-------------------
 
-Check the allocated size of memory and replace it with the correct size if
-necessary.
+     void bar(char *data_handle, int width, int height) {
+       std::vector<sycl::uchar3> data;
+       for (int Index = 0; Index < width * height; Index++) {
+         data.emplace_back(data_handle[Index * 3], data_handle[Index * 3 = 1],
+                           data_handle[Index * 3+2]);
+       }
+       ...
+     }
+