[SYCLomatic][DOC] 2025.0 new options and warning messages (#2394)

* Updating DPCT Warning Messages and Option Descriptions for 2025.0

* Adding new DPCT warning messages for 2025.0

Author: lapel22

docs/_include_files/options_def.rst

@@ -612,7 +612,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 1127. Hyphen-separated ranges are also allowed. For
+range from 1000 to 1132. Hyphen-separated ranges are also allowed. For
 example: ``-suppress-warnings=1000-1010,1011``.
 
 .. _end-suppress-warnings:
@@ -690,7 +690,9 @@ The values are:
 - ``=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``: Experimental helper function used to help cross-group synchronization during migration. See more details in ``dpct::experimental::nd_range_barrier`` in header file ``util.hpp``.
+- ``=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``.
 - ``=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>`__.
@@ -710,11 +712,12 @@ The values are:
 .. _desc-use-explicit-namespace:
 
 Define the namespaces to use explicitly in generated code. The value is
-a comma-separated list. Default: ``dpct, sycl``.
+a comma-separated list. Default: ``dpct/syclcompat, sycl``.
 
 Possible values are:
 
 - ``=dpct``: Generate code with ``dpct::`` namespace.
+- ``=syclcompat``: Generate code with ``syclcompat::`` namespace.
 - ``=none``: Generate code without any namespaces. Cannot be used with other values.
 - ``=sycl``: Generate code with ``sycl::`` namespace. Cannot be used with ``cl``
   or ``sycl-math`` values.
@@ -725,6 +728,18 @@ Possible values are:
 
 
 
+.. _opt-use-syclcompat:
+
+``--use-syclcompat``
+
+.. _desc-use-syclcompat:
+
+Use SYCLcompat header-only library (``syclcompat::`` namespace) to assist the migration of input source code. Default: ``off``.
+
+.. _end-use-syclcompat:
+
+
+
 .. _opt-usm-level:
 
 ``--usm-level=<value>``
@@ -739,6 +754,8 @@ Set the Unified Shared Memory (USM) level to use in source code generation:
 
 .. _end-usm-level:
 
+
+
 .. _opt-vcxprojfile:
 
 ``--vcxprojfile=<file>``
@@ -783,6 +800,7 @@ Same as -p.
 .. _end-compilation-db:
 
 
+
 .. _opt-gen-helper-func:
 
 ``--gen-helper-function``
@@ -794,6 +812,7 @@ Generate helper function files in the ``--out-root`` directory. Default: ``off``
 .. _end-gen-helper-func:
 
 
+
 .. _opt-analysis-mode:
 
 ``--analysis-mode``
@@ -805,6 +824,7 @@ Only generate a report for porting effort. Default: ``off``.
 .. _end-analysis-mode:
 
 
+
 .. _opt-analysis-mode-output-file:
 
 ``--analysis-mode-output-file``
@@ -816,6 +836,7 @@ Specify the file where the analysis mode report is saved. Default: Output to ``s
 .. _end-analysis-mode-output-file:
 
 
+
 .. _opt-codepin-report:
 
 ``--codepin-report``
@@ -826,6 +847,8 @@ Call ``codepin-report.py`` to generate CodePin report by parsing execution log f
 
 .. _end-codepin-report:
 
+
+
 .. _opt-enable-codepin:
 
 ``--enable-codepin``
@@ -837,6 +860,7 @@ EXPERIMENTAL: Generate instrumented CUDA and SYCL code for debug and verificatio
 .. _end-enable-codepin:
 
 
+
 .. _opt-intercept-build:
 
 ``--intercept-build``
@@ -848,6 +872,7 @@ Intercept build tool to generate a compilation database.
 .. _end-intercept-build:
 
 
+
 .. _opt-migrate-build-script:
 
 ``--migrate-build-script=<value>``
@@ -861,6 +886,7 @@ EXPERIMENTAL: Migrate build script(s).
 .. _end-migrate-build-script:
 
 
+
 .. _opt-migrate-build-script-only:
 
 ``--migrate-build-script-only``
@@ -872,6 +898,7 @@ EXPERIMENTAL: Only migrate the build script(s). Default: ``off``.
 .. _end-migrate-build-script-only:
 
 
+
 .. _opt-sycl-file-extension:
 
 ``--sycl-file-extension=<value>``
@@ -888,6 +915,7 @@ The values are:
 .. _end-sycl-file-extension:
 
 
+
 .. _opt-intercept-build-block:
 
 intercept-build Options

docs/dev_guide/reference/command-line-options-ref/alpha-list.rst

@@ -318,6 +318,12 @@ The following table lists |tool_name| command line options in alphabetical order
      - .. include:: /_include_files/options_def.rst
           :start-after: desc-use-explicit-namespace:
           :end-before: end-use-explicit-namespace:
+   * - .. include:: /_include_files/options_def.rst
+          :start-after: opt-use-syclcompat:
+          :end-before: desc-use-syclcompat:
+     - .. include:: /_include_files/options_def.rst
+          :start-after: desc-use-syclcompat:
+          :end-before: end-use-syclcompat:
    * - .. include:: /_include_files/options_def.rst
           :start-after: opt-usm-level:
           :end-before: desc-usm-level:

docs/dev_guide/reference/diagnostic_ref/dpct1087.rst

@@ -1,33 +1,31 @@
 .. _DPCT1087:
 
-DPCT1087
-========
+DPCT1087 [UPDATE]
+=================
 
 Message
 -------
 
 .. _msg-1087-start:
 
-SYCL currently does not support cross group synchronization. You can specify
-``--use-experimental-features=nd_range_barrier`` to use the dpct helper function
-``nd_range_barrier`` to migrate *<synchronization API call>*.
+SYCL currently does not support cross group synchronization. You can
+specify ``--use-experimental-features=root-group`` to use the root-group to migrate
+*<synchronization API call>*.
 
 .. _msg-1087-end:
 
 Detailed Help
 -------------
 
-By default, the dpct helper function ``nd_range_barrier`` is not used to migrate
-CUDA\* grid level synchronization. To use ``nd_range_barrier`` to migrate CUDA
-grid level synchronization, specify ``--use-experimental-features=nd_range_barrier``
-in the migration command.
+By default, the SYCL root group extension is not used to migrate CUDA\* grid level
+synchronization. To use ``root-group`` to migrate CUDA grid level synchronization,
+specify ``--use-experimental-features=root-group`` in the migration command.
 
 Suggestions to Fix
 ------------------
 
-Specify ``--use-experimental-features=nd_range_barrier`` in the migration command
-to use dpct helper function ``nd_range_barrier`` to migrate CUDA grid level
-synchronization.
+Specify ``--use-experimental-features= root-group`` in the migration command to use
+the root-group to migrate CUDA grid level synchronization.
 
 For example, this original CUDA\* code:
 
@@ -52,14 +50,19 @@ results in the following migrated SYCL code:
    void kernel() {
    
      /*
-     DPCT1087:1: SYCL currently does not support cross group synchronization. You
-     can specify "--use-experimental-features=nd_range_barrier" to use the dpct
-     helper function nd_range_barrier to migrate this_grid().
+     DPCT1119:1: Migration of cooperative_groups::__v1::grid_group::this_grid is not
+     supported, please try to remigrate with option:
+     --use-experimental-features=root-group.
+     */
+     /*
+     DPCT1119:2: Migration of cooperative_groups::__v1::grid_group is not
+     supported, please try to remigrate with option:
+     --use-experimental-features=root-group.
      */
      cg::grid_group grid = cg::this_grid();
      /*
      DPCT1087:0: SYCL currently does not support cross group synchronization. You
-     can specify "--use-experimental-features=nd_range_barrier" to use the dpct
+     can specify "--use-experimental-features=root-group" to use the dpct
      helper function nd_range_barrier to migrate grid.sync().
      */
      grid.sync();
@@ -79,25 +82,18 @@ which is rewritten to:
    :linenos:
 
    void kernel(const sycl::nd_item<3> &item_ct1,
-               sycl::atomic_ref<unsigned int, sycl::memory_order::seq_cst, sycl::memory_scope::device, sycl::access::address_space::global_space> &sync_ct1) {
-   
-     dpct::experimental::nd_range_barrier(item_ct1, sync_ct1);
+     sycl::ext::oneapi::experimental::root_group grid =
+         item_ct1.ext_oneapi_get_root_group();
+     sycl::group_barrier(grid);
    }
    
    void foo() {
-     dpct::global_memory<unsigned int, 0> d_sync_ct1(0);
-     unsigned *sync_ct1 = d_sync_ct1.get_ptr(dpct::get_in_order_queue());
-     dpct::get_in_order_queue().memset(sync_ct1, 0, sizeof(int)).wait();
-     dpct::get_in_order_queue()
-         .parallel_for(
-             sycl::nd_range<3>(sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)),
+     auto exp_props = sycl::ext::oneapi::experimental::properties{
+         sycl::ext::oneapi::experimental::use_root_sync};
+     dpct::get_in_order_queue().parallel_for(
+         sycl::nd_range<3>(sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)),
+         exp_props, [=](sycl::nd_item<3> item_ct1) {
+           kernel(item_ct1);
              [=](sycl::nd_item<3> item_ct1) {
-               auto atm_sync_ct1 =
-                   sycl::atomic_ref<unsigned int, sycl::memory_order::seq_cst,
-                                    sycl::memory_scope::device,
-                                    sycl::access::address_space::global_space>(
-                       sync_ct1[0]);
-               kernel(item_ct1, atm_sync_ct1);
-             })
-         .wait();
+         });
    }

docs/dev_guide/reference/diagnostic_ref/dpct1130.rst

@@ -0,0 +1,98 @@
+.. _DPCT1130:
+
+DPCT1130
+========
+
+Message
+-------
+
+.. _msg-1130-start:
+
+SYCL 2020 standard does not support dynamic parallelism (launching kernel in device
+code). Please rewrite the code.
+
+.. _msg-1130-end:
+
+Detailed Help
+-------------
+
+SYCL\* does not support launching kernel in device code. The user needs to merge the
+parent kernel and child kernel together.
+
+Suggestions to Fix
+------------------
+
+For example, this original CUDA\* code:
+
+.. code-block:: cpp
+   :linenos:
+
+   __global__ void childKernel() {
+	 ...
+   }
+   __global__ void parentKernel() {
+	 ...
+	 childKernel<<<4, 4>>>();
+	 ...
+   }
+   void foo() {
+	 ...
+	 parentKernel<<<8, 8>>>();
+	 ...
+   }
+
+results in the following migrated SYCL code:
+
+.. code-block:: cpp
+   :linenos:
+
+   void childKernel() {
+     ...
+   }
+   void parentKernel() {
+     ...
+     /*
+     DPCT1130:0: SYCL 2020 standard does not support dynamic parallelism (launching
+     kernel in device code). Please rewrite the code.
+     */
+     childKernel<<<4, 4>>>();
+     ...
+   }
+  void foo() {
+    ...
+    dpct::get_in_order_queue().parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, 8) * sycl::range<3>(1, 1, 8),
+                          sycl::range<3>(1, 1, 8)),
+        [=](sycl::nd_item<3> item_ct1) {
+          parentKernel();
+        });
+    ...
+  }
+
+which is rewritten to:
+
+.. code-block:: cpp
+   :linenos:
+
+   void childKernel() {
+     ...
+   }
+   void parentKernel() {
+     ...
+     childKernel(); // call childKernel() as a device function, need to adjust the work
+     for each work item.
+     ...
+   }
+   void foo() {
+     ...
+     dpct::get_in_order_queue().parallel_for(
+         sycl::nd_range<3>(sycl::range<3>(1, 1, placeholder /*Adjust the global range
+         based on the thread model between parentKernel and childKernel*/),
+                           sycl::range<3>(1, 1, placeholder /*Adjust the local range
+                           based on the thread model between parentKernel and
+                           childKernel */)),
+         [=](sycl::nd_item<3> item_ct1) {
+           parentKernel();
+         });
+     ...
+   }

docs/dev_guide/reference/diagnostic_ref/dpct1131.rst

@@ -0,0 +1,26 @@
+.. _DPCT1131:
+
+DPCT1131
+========
+
+Message
+-------
+
+.. _msg-1131-start:
+
+The migration of "*<API name>*" is not currently supported with SYCLcompat. Please
+adjust the code manually.
+
+.. _msg-1131-end:
+
+Detailed Help
+-------------
+
+When using syclcompat header-only library to assist the migration, this API migration is
+not supported.
+
+Suggestions to Fix
+------------------
+
+This kind of code requires manual migration to SYCL\*. Please refer to the migration
+result without specifying “--use-syclcompat”.