[SYCLomatic] Add 12.9 header file migration support

README.md

@@ -27,7 +27,7 @@ in daily releases. None of the branches in the project are stable or rigorously
 tested for production quality control, so the quality of these releases is
 expected to be similar to the daily releases.
 
-SYCLomatic supports migrating programs implemented with CUDA versions 8.0, 9.x, 10.x, 11.x, 12.0-12.8. The list of supported languages and versions may be extended in the future.
+SYCLomatic supports migrating programs implemented with CUDA versions 8.0, 9.x, 10.x, 11.x, 12.0-12.9. The list of supported languages and versions may be extended in the future.
 
 ## Build from source code
 ### Prerequisites

clang/include/clang/Basic/Cuda.h

@@ -50,13 +50,16 @@ enum class CudaVersion {
   CUDA_126,
   CUDA_128,
 #ifdef SYCLomatic_CUSTOMIZATION
-  FULLY_SUPPORTED = CUDA_126,
+  CUDA_129,
+  FULLY_SUPPORTED = CUDA_123,
+  PARTIALLY_SUPPORTED =
+      CUDA_129, // Partially supported. Proceed with a warning.
 #else
   FULLY_SUPPORTED = CUDA_123,
-#endif
   PARTIALLY_SUPPORTED =
       CUDA_128, // Partially supported. Proceed with a warning.
-  NEW = 10000,  // Too new. Issue a warning, but allow using it.
+#endif
+  NEW = 10000, // Too new. Issue a warning, but allow using it.
 };
 const char *CudaVersionToString(CudaVersion V);
 #ifdef SYCLomatic_CUSTOMIZATION

clang/lib/Basic/Cuda.cpp

@@ -45,6 +45,9 @@ static const CudaVersionMapEntry CudaNameVersionMap[] = {
     CUDA_ENTRY(12, 5),
     CUDA_ENTRY(12, 6),
     CUDA_ENTRY(12, 8),
+#ifdef SYCLomatic_CUSTOMIZATION
+    CUDA_ENTRY(12, 9),
+#endif
     {"", CudaVersion::NEW, llvm::VersionTuple(std::numeric_limits<int>::max())},
     {"unknown", CudaVersion::UNKNOWN, {}} // End of list tombstone.
 };

clang/lib/DPCT/RulesLang/RulesLang.cpp

@@ -53,6 +53,7 @@ using namespace clang::tooling;
 extern clang::tooling::UnifiedPath DpctInstallPath; // Installation directory for this tool
 extern DpctOption<opt, bool> ProcessAll;
 extern DpctOption<opt, bool> AsyncHandler;
+extern int ThrustVersion;
 
 namespace clang{
 namespace dpct{
@@ -280,6 +281,17 @@ void MiscAPIRule::runRule(const MatchFinder::MatchResult &Result) {
 
 // Rule for types migration in var declarations and field declarations
 void TypeInDeclRule::registerMatcher(MatchFinder &MF) {
+  int ThrustMajorVersion = ThrustVersion / 100000;
+  int ThrustMinorVersion = ThrustVersion / 100 % 1000;
+
+  auto thrustNamespace = [=]() -> std::string {
+    if (ThrustMajorVersion >= 2 && ThrustMinorVersion >= 8) {
+      // For CUDA-12.9 or later
+      return "cuda::std::";
+    }
+    return "thrust::";
+  };
+
   MF.addMatcher(
       typeLoc(
           loc(qualType(hasDeclaration(namedDecl(hasAnyName(
@@ -294,7 +306,7 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) {
               "thrust::permutation_iterator", "thrust::iterator_difference",
               "cusolverDnHandle_t", "cusolverDnParams_t", "gesvdjInfo_t",
               "syevjInfo_t", "thrust::device_malloc_allocator",
-              "thrust::divides", "thrust::tuple", "thrust::maximum",
+              "thrust::divides", thrustNamespace() + "tuple", "thrust::maximum",
               "thrust::multiplies", "thrust::plus", "cudaDataType_t",
               "cudaError_t", "CUresult", "CUdevice", "cudaEvent_t",
               "cublasStatus_t", "cuComplex", "cuFloatComplex",
@@ -316,15 +328,16 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) {
               "curandRngType_t", "curandOrdering_t", "cufftHandle", "cufftReal",
               "cufftDoubleReal", "cufftComplex", "cufftDoubleComplex",
               "cufftResult_t", "cufftResult", "cufftType_t", "cufftType",
-              "thrust::pair", "CUdeviceptr", "cudaDeviceAttr", "CUmodule",
-              "CUjit_option", "CUfunction", "cudaMemcpyKind", "cudaComputeMode",
-              "__nv_bfloat16", "cooperative_groups::__v1::thread_group",
+              thrustNamespace() + "pair", "CUdeviceptr", "cudaDeviceAttr",
+              "CUmodule", "CUjit_option", "CUfunction", "cudaMemcpyKind",
+              "cudaComputeMode", "__nv_bfloat16",
+              "cooperative_groups::__v1::thread_group",
               "cooperative_groups::__v1::thread_block", "libraryPropertyType_t",
               "libraryPropertyType", "cudaDataType_t", "cudaDataType",
               "cublasComputeType_t", "cublasAtomicsMode_t", "cublasMath_t",
               "CUmem_advise_enum", "CUmem_advise", "CUmemorytype",
-              "CUmemorytype_enum", "thrust::tuple_element",
-              "thrust::tuple_size", "thrust::zip_iterator",
+              "CUmemorytype_enum", thrustNamespace() + "tuple_element",
+              thrustNamespace() + "tuple_size", "thrust::zip_iterator",
               "cudaPointerAttributes", "CUpointer_attribute",
               "cusolverEigRange_t", "cudaUUID_t", "cusolverDnFunction_t",
               "cusolverAlgMode_t", "cusparseIndexType_t", "cusparseFormat_t",

clang/lib/DPCT/RulesLangLib/CUBAPIMigration.cpp

@@ -129,15 +129,9 @@ void CubTypeRule::runRule(
 
 bool CubTypeRule::CanMappingToSyclNativeBinaryOp(StringRef OpTypeName) {
   return OpTypeName == "cub::Sum" || OpTypeName == "cub::Max" ||
-         OpTypeName == "cub::Min";
-}
-
-bool CubTypeRule::CanMappingToSyclType(StringRef OpTypeName) {
-  return CanMappingToSyclNativeBinaryOp(OpTypeName) ||
-         OpTypeName == "cub::Equality" || OpTypeName == "cub::NullType" ||
-
-         // Ignore template arguments, .e.g cub::KeyValuePair<int, int>
-         OpTypeName.starts_with("cub::KeyValuePair");
+         OpTypeName == "cub::Min" || OpTypeName == "cuda::std::plus<void>" ||
+         OpTypeName == "cuda::maximum<void>" ||
+         OpTypeName == "cuda::minimum<void>";
 }
 
 void CubDeviceLevelRule::registerMatcher(ast_matchers::MatchFinder &MF) {
@@ -854,9 +848,9 @@ std::string CubRule::getOpRepl(const Expr *Operator) {
         Obj->getType().getCanonicalType());
     if (OpType == "cub::Sum" || OpType == "cuda::std::plus<void>") {
       OpRepl = MapNames::getClNamespace() + "plus<>()";
-    } else if (OpType == "cub::Max") {
+    } else if (OpType == "cub::Max" || OpType == "cuda::maximum<void>") {
       OpRepl = MapNames::getClNamespace() + "maximum<>()";
-    } else if (OpType == "cub::Min") {
+    } else if (OpType == "cub::Min" || OpType == "cuda::minimum<void>") {
       OpRepl = MapNames::getClNamespace() + "minimum<>()";
     }
   };
@@ -873,7 +867,9 @@ std::string CubRule::getOpRepl(const Expr *Operator) {
         std::string OpType = DpctGlobalInfo::getUnqualifiedTypeName(
             D->getType().getCanonicalType());
         if (OpType == "cub::Sum" || OpType == "cub::Max" ||
-            OpType == "cub::Min" || OpType == "cuda::std::plus<void>") {
+            OpType == "cub::Min" || OpType == "cuda::std::plus<void>" ||
+            OpType == "cuda::maximum<void>" ||
+            OpType == "cuda::minimum<void>") {
           ExprAnalysis EA(Operator);
           OpRepl = EA.getReplacedString();
         }

clang/lib/DPCT/RulesLangLib/CUBAPIMigration.h

@@ -21,7 +21,6 @@ class CubTypeRule : public NamedMigrationRule<CubTypeRule> {
   void runRule(const ast_matchers::MatchFinder::MatchResult &Result);
 
   static bool CanMappingToSyclNativeBinaryOp(StringRef OpTypeName);
-  static bool CanMappingToSyclType(StringRef OpTypeName);
 };
 
 class CubDeviceLevelRule : public NamedMigrationRule<CubDeviceLevelRule> {

clang/lib/DPCT/RulesLangLib/ThrustAPIMigration.cpp

@@ -22,116 +22,44 @@ void ThrustAPIRule::registerMatcher(ast_matchers::MatchFinder &MF) {
   // API register
   auto functionName = [&]() { return hasAnyName("on"); };
 
-  // THRUST_200302___CUDA_ARCH_LIST___NS is newly imported inline
-  // namespace by thrust library in CUDA header file 12.4.
-  auto thrustFuncNameCuda124 = [&]() {
-    return hasAnyName("THRUST_200302___CUDA_ARCH_LIST___NS",
-                      "THRUST_200302___CUDA_ARCH_LIST___NS::detail",
-                      "THRUST_200302___CUDA_ARCH_LIST___NS::system");
-  };
-
-  // THRUST_200400___CUDA_ARCH_LIST___NS is newly imported inline
-  // namespace by thrust library in CUDA header file 12.5.
-  auto thrustFuncNameCuda125 = [&]() {
-    return hasAnyName("THRUST_200400___CUDA_ARCH_LIST___NS",
-                      "THRUST_200400___CUDA_ARCH_LIST___NS::detail",
-                      "THRUST_200400___CUDA_ARCH_LIST___NS::system");
-  };
-
-  // THRUST_200500___CUDA_ARCH_LIST___NS is newly imported inline
-  // namespace by thrust library in CUDA header file 12.6.
-  auto thrustFuncNameCuda126 = [&]() {
-    return hasAnyName("THRUST_200500___CUDA_ARCH_LIST___NS",
-                      "THRUST_200500___CUDA_ARCH_LIST___NS::detail",
-                      "THRUST_200500___CUDA_ARCH_LIST___NS::system");
-  };
-
-  // THRUST_200700___CUDA_ARCH_LIST___NS is newly imported inline
-  // namespace by thrust library in CUDA header file 12.6.
-  auto thrustFuncNameCuda128 = [&]() {
-    return hasAnyName("THRUST_200700___CUDA_ARCH_LIST___NS",
-                      "THRUST_200700___CUDA_ARCH_LIST___NS::detail",
-                      "THRUST_200700___CUDA_ARCH_LIST___NS::system");
-  };
-
-  auto thrustFuncNameCudaCommon = [&]() {
-    return hasAnyName("thrust", "thrust::detail", "thrust::system", "__4");
-  };
-
   int ThrustMajorVersion = ThrustVersion / 100000;
   int ThrustMinorVersion = ThrustVersion / 100 % 1000;
 
-  if (ThrustMajorVersion == 2 && ThrustMinorVersion == 3) {
-    // For CUDA-12.4
-    MF.addMatcher(
-        callExpr(
-            anyOf(callee(functionDecl(anyOf(
-                      hasDeclContext(namespaceDecl(thrustFuncNameCuda124())),
-                      hasDeclContext(namespaceDecl(thrustFuncNameCudaCommon())),
-                      functionName()))),
-                  callee(unresolvedLookupExpr(
-                      hasAnyDeclaration(namedDecl(hasDeclContext(namespaceDecl(
-                          anyOf(thrustFuncNameCuda124(),
-                                thrustFuncNameCudaCommon())))))))))
-            .bind("thrustFuncCall"),
-        this);
-
-  } else if (ThrustMajorVersion == 2 && ThrustMinorVersion == 4) {
-    // For CUDA-12.5
-    MF.addMatcher(
-        callExpr(
-            anyOf(callee(functionDecl(anyOf(
-                      hasDeclContext(namespaceDecl(thrustFuncNameCuda125())),
-                      hasDeclContext(namespaceDecl(thrustFuncNameCudaCommon())),
-                      functionName()))),
-                  callee(unresolvedLookupExpr(
-                      hasAnyDeclaration(namedDecl(hasDeclContext(namespaceDecl(
-                          anyOf(thrustFuncNameCuda125(),
-                                thrustFuncNameCudaCommon())))))))))
-            .bind("thrustFuncCall"),
-        this);
-  } else if (ThrustMajorVersion == 2 && ThrustMinorVersion == 5) {
-    // For CUDA-12.6
-    MF.addMatcher(
-        callExpr(
-            anyOf(callee(functionDecl(anyOf(
-                      hasDeclContext(namespaceDecl(thrustFuncNameCuda126())),
-                      hasDeclContext(namespaceDecl(thrustFuncNameCudaCommon())),
-                      functionName()))),
-                  callee(unresolvedLookupExpr(
-                      hasAnyDeclaration(namedDecl(hasDeclContext(namespaceDecl(
-                          anyOf(thrustFuncNameCuda126(),
-                                thrustFuncNameCudaCommon())))))))))
-            .bind("thrustFuncCall"),
-        this);
-  } else if (ThrustMajorVersion == 2 && ThrustMinorVersion == 7) {
-    // For CUDA-12.8
-    MF.addMatcher(
-        callExpr(
-            anyOf(callee(functionDecl(anyOf(
-                      hasDeclContext(namespaceDecl(thrustFuncNameCuda128())),
-                      hasDeclContext(namespaceDecl(thrustFuncNameCudaCommon())),
-                      functionName()))),
-                  callee(unresolvedLookupExpr(
-                      hasAnyDeclaration(namedDecl(hasDeclContext(namespaceDecl(
-                          anyOf(thrustFuncNameCuda128(),
-                                thrustFuncNameCudaCommon())))))))))
-            .bind("thrustFuncCall"),
-        this);
-  } else {
-    // For CUDA SDK versions before CUDA-12.4
-    MF.addMatcher(
-        callExpr(
-            anyOf(callee(functionDecl(anyOf(
-                      hasDeclContext(namespaceDecl(thrustFuncNameCudaCommon())),
-
-                      functionName()))),
-                  callee(unresolvedLookupExpr(
-                      hasAnyDeclaration(namedDecl(hasDeclContext(
-                          namespaceDecl(thrustFuncNameCudaCommon()))))))))
-            .bind("thrustFuncCall"),
-        this);
-  }
+  auto thrustFuncName = [&]() {
+#define COMMON "thrust", "thrust::detail", "thrust::system", "__4"
+    if (ThrustMajorVersion == 2 && ThrustMinorVersion == 3)
+      return hasAnyName("THRUST_200302___CUDA_ARCH_LIST___NS",
+                        "THRUST_200302___CUDA_ARCH_LIST___NS::detail",
+                        "THRUST_200302___CUDA_ARCH_LIST___NS::system", COMMON);
+    if (ThrustMajorVersion == 2 && ThrustMinorVersion == 4)
+      return hasAnyName("THRUST_200400___CUDA_ARCH_LIST___NS",
+                        "THRUST_200400___CUDA_ARCH_LIST___NS::detail",
+                        "THRUST_200400___CUDA_ARCH_LIST___NS::system", COMMON);
+    if (ThrustMajorVersion == 2 && ThrustMinorVersion == 5)
+      return hasAnyName("THRUST_200500___CUDA_ARCH_LIST___NS",
+                        "THRUST_200500___CUDA_ARCH_LIST___NS::detail",
+                        "THRUST_200500___CUDA_ARCH_LIST___NS::system", COMMON);
+    if (ThrustMajorVersion == 2 && ThrustMinorVersion == 7)
+      return hasAnyName("THRUST_200700___CUDA_ARCH_LIST___NS",
+                        "THRUST_200700___CUDA_ARCH_LIST___NS::detail",
+                        "THRUST_200700___CUDA_ARCH_LIST___NS::system", COMMON);
+    if (ThrustMajorVersion == 2 && ThrustMinorVersion == 8)
+      return hasAnyName("THRUST_200802_SM___CUDA_ARCH_LIST___NS",
+                        "THRUST_200802_SM___CUDA_ARCH_LIST___NS::detail",
+                        "THRUST_200802_SM___CUDA_ARCH_LIST___NS::system",
+                        COMMON);
+    return hasAnyName(COMMON);
+#undef COMMON
+  };
+
+  MF.addMatcher(
+      callExpr(anyOf(callee(functionDecl(
+                         anyOf(hasDeclContext(namespaceDecl(thrustFuncName())),
+                               functionName()))),
+                     callee(unresolvedLookupExpr(hasAnyDeclaration(namedDecl(
+                         hasDeclContext(namespaceDecl(thrustFuncName()))))))))
+          .bind("thrustFuncCall"),
+      this);
 
   // THRUST_STATIC_ASSERT macro register
   MF.addMatcher(staticAssertDecl(isExpandedFromMacro("THRUST_STATIC_ASSERT"))

clang/lib/Driver/ToolChains/Cuda.cpp

@@ -199,9 +199,10 @@ bool CudaInstallationDetector::ParseCudaVersionFile(const std::string &FilePath)
     Version = CudaVersion::CUDA_126;
   } else if (Major == 12 && Minor == 8) {
     Version = CudaVersion::CUDA_128;
+  } else if (Major == 12 && Minor == 9) {
+    Version = CudaVersion::CUDA_129;
   }
 
-
   if (Version != CudaVersion::UNKNOWN) {
     IsVersionSupported = true;
     return true;
@@ -264,7 +265,7 @@ CudaVersion getCudaVersion(uint32_t raw_version) {
     return CudaVersion::CUDA_124;
   if (raw_version < 12060)
     return CudaVersion::CUDA_125;
-  if (raw_version < 12080)
+  if (raw_version < 12070)
     return CudaVersion::CUDA_126;
   if (raw_version < 12090)
     return CudaVersion::CUDA_128;
@@ -350,31 +351,10 @@ CudaInstallationDetector::CudaInstallationDetector(
 
   // In decreasing order so we prefer newer versions to older versions.
 #ifdef SYCLomatic_CUSTOMIZATION
-  std::initializer_list<const char *> Versions = {"12.8",
-                                                  "12.6",
-                                                  "12.5",
-                                                  "12.4",
-                                                  "12.3",
-                                                  "12.2",
-                                                  "12.1",
-                                                  "12.0",
-                                                  "11.8",
-                                                  "11.7",
-                                                  "11.6",
-                                                  "11.5",
-                                                  "11.4",
-                                                  "11.3",
-                                                  "11.2",
-                                                  "11.1",
-                                                  "10.2",
-                                                  "10.1",
-                                                  "10.0",
-                                                  "9.2",
-                                                  "9.1",
-                                                  "9.0",
-                                                  "8.0",
-                                                  "7.5",
-                                                  "7.0"};
+  std::initializer_list<const char *> Versions = {
+      "12.9", "12.8", "12.6", "12.5", "12.4", "12.3", "12.2", "12.1", "12.0",
+      "11.8", "11.7", "11.6", "11.5", "11.4", "11.3", "11.2", "11.1", "10.2",
+      "10.1", "10.0", "9.2",  "9.1",  "9.0",  "8.0",  "7.5",  "7.0"};
 #else
   std::initializer_list<const char *> Versions = {
       "11.4", "11.3", "11.2", "11.1", "10.2", "10.1", "10.0",

clang/test/dpct/NVTX/NVTX-linux.cu

@@ -1,4 +1,5 @@
 // UNSUPPORTED: system-windows
+// UNSUPPORTED: cuda-12.9
 // RUN: dpct --format-range=none -in-root %S -out-root %T %S/NVTX-linux.cu --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only
 // RUN: FileCheck --input-file %T/NVTX-linux.dp.cpp --match-full-lines %s
 

clang/test/dpct/allocator_syclcompat.cu

@@ -1,5 +1,5 @@
-// UNSUPPORTED: cuda-8.0, cuda-12.0, cuda-12.1, cuda-12.2, cuda-12.3, cuda-12.4, cuda-12.5, cuda-12.6, cuda-12.8
-// UNSUPPORTED: v8.0, v12.0, v12.1, v12.2, v12.3, v12.4, v12.5, v12.6, v12.8
+// UNSUPPORTED: cuda-8.0, cuda-12.0, cuda-12.1, cuda-12.2, cuda-12.3, cuda-12.4, cuda-12.5, cuda-12.6, cuda-12.8, cuda-12.9
+// UNSUPPORTED: v8.0, v12.0, v12.1, v12.2, v12.3, v12.4, v12.5, v12.6, v12.8, v12.9
 // RUN: dpct --format-range=none --use-syclcompat -out-root %T/allocator_syclcompat %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only
 // RUN: FileCheck --match-full-lines --input-file %T/allocator_syclcompat/allocator_syclcompat.dp.cpp %s
 // RUN: %if build_lit %{icpx -c -fsycl -DNO_BUILD_TEST %T/allocator_syclcompat/allocator_syclcompat.dp.cpp -o %T/allocator_syclcompat/allocator_syclcompat.dp.o %}

clang/test/dpct/comments.cu

@@ -1,5 +1,5 @@
-// UNSUPPORTED: cuda-12.0, cuda-12.1, cuda-12.2, cuda-12.3, cuda-12.4, cuda-12.5, cuda-12.6, cuda-12.8
-// UNSUPPORTED: v12.0, v12.1, v12.2, v12.3, v12.4, v12.5, v12.6, v12.8
+// UNSUPPORTED: cuda-12.0, cuda-12.1, cuda-12.2, cuda-12.3, cuda-12.4, cuda-12.5, cuda-12.6, cuda-12.8, cuda-12.9
+// UNSUPPORTED: v12.0, v12.1, v12.2, v12.3, v12.4, v12.5, v12.6, v12.8, v12.9
 // RUN: dpct --format-range=none -out-root %T/comments %s --cuda-include-path="%cuda-path/include" --comments -- -std=c++14 -x cuda --cuda-host-only
 // RUN: FileCheck %s --match-full-lines --input-file %T/comments/comments.dp.cpp
 // RUN: %if build_lit %{icpx -c -fsycl %T/comments/comments.dp.cpp -o %T/comments/comments.dp.o %}

clang/test/dpct/cooperative_groups_unsupport.cu

@@ -1,5 +1,5 @@
-// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1, cuda-10.2, cuda-11.0, cuda-12.0, cuda-12.1, cuda-12.2, cuda-12.3, cuda-12.4, cuda-12.5, cuda-12.6, cuda-12.8
-// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1, v10.2, v11.0, v12.0, v12.1, v12.2, v12.3, v12.4, v12.5, v12.6, v12.8
+// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1, cuda-10.2, cuda-11.0, cuda-12.0, cuda-12.1, cuda-12.2, cuda-12.3, cuda-12.4, cuda-12.5, cuda-12.6, cuda-12.8, cuda-12.9
+// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1, v10.2, v11.0, v12.0, v12.1, v12.2, v12.3, v12.4, v12.5, v12.6, v12.8, v12.9
 // RUN: dpct --format-range=none -out-root %T/cooperative_groups_unsupport %s --cuda-include-path="%cuda-path/include" --use-experimental-features=logical-group --extra-arg="-std=c++14"
 // RUN: FileCheck %s --match-full-lines --input-file %T/cooperative_groups_unsupport/cooperative_groups_unsupport.dp.cpp
 // RUN: %if build_lit %{icpx -c -fsycl %T/cooperative_groups_unsupport/cooperative_groups_unsupport.dp.cpp -o %T/cooperative_groups_unsupport/cooperative_groups_unsupport.dp.o %}