[SYCLomatic] Add the cooperative_groups::this_thread migration. (#2783)

Signed-off-by: Chen, Sheng S <sheng.s.chen@intel.com>

Author: ShengchenJ

clang/lib/DPCT/RulesLang/APINamesCooperativeGroups.inc

@@ -192,6 +192,19 @@ CONDITIONAL_FACTORY_ENTRY(
     UNSUPPORT_FACTORY_ENTRY("tiled_partition", Diagnostics::API_NOT_MIGRATED,
                             ARG("tiled_partition")))
 
+CONDITIONAL_FACTORY_ENTRY(
+    UseLogicalGroup,
+    FEATURE_REQUEST_FACTORY(
+        HelperFeatureEnum::device_ext,
+        CALL_FACTORY_ENTRY("this_thread",
+                           CALL(MapNames::getDpctNamespace() +
+                                    "experimental::logical_group",
+                                NDITEM, GROUP, ARG("1")))),
+    UNSUPPORT_FACTORY_ENTRY("this_thread",
+                            Diagnostics::TRY_EXPERIMENTAL_FEATURE,
+                            ARG("this_thread"),
+                            ARG("--use-experimental-features=logical-group")))
+
 CONDITIONAL_FACTORY_ENTRY(
     UseNonUniformGroups,
     CALL_FACTORY_ENTRY(

clang/lib/DPCT/RulesLang/MapNamesLang.cpp

@@ -308,6 +308,7 @@ const std::vector<std::string> MemoryDataTypeRule::RemoveMember{
 const std::unordered_set<std::string> MapNamesLang::CooperativeGroupsAPISet{
     "this_thread_block",
     "this_grid",
+    "this_thread",
     "sync",
     "tiled_partition",
     "thread_rank",

clang/lib/DPCT/RulesLang/RulesLangCooperativeGroups.cpp

@@ -119,7 +119,7 @@ void CooperativeGroupsFunctionRule::runRule(
       FuncName == "num_threads" || FuncName == "inclusive_scan" ||
       FuncName == "exclusive_scan" || FuncName == "coalesced_threads" ||
       FuncName == "this_grid" || FuncName == "num_blocks" ||
-      FuncName == "block_rank") {
+      FuncName == "block_rank" || FuncName == "this_thread") {
     // There are 3 usages of cooperative groups APIs.
     // 1. cg::thread_block tb; tb.sync(); // member function
     // 2. cg::thread_block tb; cg::sync(tb); // free function

clang/lib/DPCT/SrcAPI/APINames.inc

@@ -2136,7 +2136,7 @@ ENTRY(cuGraphicsResourceGetMappedEglFrame, cuGraphicsResourceGetMappedEglFrame,
 // Coopeartive_group migration.
 ENTRY(cooperative_groups::coalesced_threads, cooperative_groups::__v1::coalesced_threads, true, NO_FLAG, P4, "Successful/DPCT1119")
 ENTRY(cooperative_groups::thread_rank, cooperative_groups::__v1::thread_rank, true, NO_FLAG, P4, "Successful")
-ENTRY(cooperative_groups::this_thread, cooperative_groups::__v1::this_thread, false, NO_FLAG, P4, "comment")
+ENTRY(cooperative_groups::this_thread, cooperative_groups::__v1::this_thread, true, NO_FLAG, P4, "Successful/DPCT1119")
 ENTRY(cooperative_groups::this_grid, cooperative_groups::__v1::this_grid, true, NO_FLAG, P4, "Successful")
 ENTRY(cooperative_groups::this_multi_grid, cooperative_groups::__v1::this_multi_grid, false, NO_FLAG, P4, "comment")
 ENTRY(cooperative_groups::this_thread_block, cooperative_groups::__v1::this_thread_block, true, NO_FLAG, P4, "Successful")

clang/test/dpct/sync_api_nd_range_barrier.cu

@@ -205,3 +205,8 @@ int foo3() {
 // CHECK-NEXT:           const sycl::group<3> *cptb);
 __device__ void foo4(const cg::thread_block &crtb,
                      const cg::thread_block *cptb);
+
+__device__ void foo5() {
+  // CHECK:  auto thread = dpct::experimental::logical_group(item_ct1, item_ct1.get_group(), 1);
+  auto thread = cg::this_thread();
+}