Merge remote-tracking branch 'origin/SYCLomatic' into shared

Author: zhiweij1

clang/lib/DPCT/DPCT.cpp

@@ -525,7 +525,8 @@ static void loadMainSrcFileInfo(clang::tooling::UnifiedPath OutRoot) {
       llvm::errs() << getLoadYamlFailWarning(YamlFilePath);
     }
 
-    if (MigrateBuildScriptOnly || DpctGlobalInfo::migrateCMakeScripts()) {
+    if (MigrateBuildScriptOnly && !DpctGlobalInfo::migratePythonScripts() ||
+        DpctGlobalInfo::migrateCMakeScripts()) {
       std::string Major = PreTU->SDKVersionMajor;
       std::string Minor = PreTU->SDKVersionMinor;
       if (!Major.empty() && !Minor.empty()) {

clang/lib/DPCT/RulesAsm/Parser/AsmParser.cpp

@@ -102,7 +102,7 @@ void InlineAsmParser::addBuiltinIdentifier() {
 
 InlineAsmBuiltinType *
 InlineAsmContext::getBuiltinType(InlineAsmBuiltinType::TypeKind Kind) {
-  assert(Kind > 0 && Kind < InlineAsmBuiltinType::NUM_TYPES && "Unknown Kind");
+  assert(Kind >= 0 && Kind < InlineAsmBuiltinType::NUM_TYPES && "Unknown Kind");
   if (AsmBuiltinTypes[Kind])
     return AsmBuiltinTypes[Kind];
 

clang/test/dpct/asm/optimize.cu

@@ -176,3 +176,14 @@ __device__ int test5() {
   MACRO(7, s32 == 0, asm("add.s32.sat %0, %1, %2;" : "=r"(s32) : "r"(s32x), "r"(0)));
   return 0;
 }
+
+// CHECK: inline void foo() {
+// CHECK-NEXT: #if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__)
+// CHECK-NEXT:   asm volatile(".reg .b8 byte0;\n");
+// CHECK-NEXT: #else
+// CHECK-NEXT:   uint8_t byte0;
+// CHECK-NEXT: #endif
+inline __device__ void foo() {
+  asm volatile(".reg .b8 byte0;\n");
+}
+

clang/test/dpct/lit.local.cfg

@@ -124,7 +124,7 @@ else:
 
 # Run the dpct sanity test.
 if not config.unsupported:
-    skipped_cases = ["cudnn_sanity.cu", "nccl_sanity.cu", "nvshmem.cu"]
+    skipped_cases = ["cudnn_sanity.cu", "nccl_sanity.cu", "nvshmem.cu", "cutensor.cu"]
     complete_process = run_sanity("cudnn_sanity.cu")
     err_message = ""
     if complete_process.returncode != 0:
@@ -147,6 +147,13 @@ if not config.unsupported:
                           "Please make sure install the header file of nvshmem and " + \
                           "export nvshmem.h in CPATH.\n"
             skipped_cases.extend(get_skipped_cases_with_string("<nccl.h>"))
+    complete_process = run_sanity("cutensor.cu")
+    if complete_process.returncode != 0:
+        if "'cutensor.h' file not found" in complete_process.stdout:
+            err_message += "'cutensor.h' header file not found in platform. " + \
+                          "Please make sure install the header file of cuTENSOR and " + \
+                          "export cutensor.h in CPATH.\n"
+            skipped_cases.extend(get_skipped_cases_with_string("<cutensor.h>"))
 
     if (err_message):
         sys.stderr.write(err_message)

clang/test/dpct/python_migration/case_006/MainSourceFiles.yaml

@@ -0,0 +1,103 @@
+---
+MainSourceFile:  '/path/to/MainSrcFiles_placehold'
+Replacements:
+  - FilePath:        '/path/to/src/foo.cpp'
+    Offset:          0
+    Length:          0
+    ReplacementText: ''
+    ConstantFlag:    ''
+    ConstantOffset:  0
+    InitStr:         ''
+    NewHostVarName:  ''
+    BlockLevelFormatFlag: false
+  - FilePath:        '/path/to/src/baz.cpp'
+    Offset:          0
+    Length:          0
+    ReplacementText: ''
+    ConstantFlag:    ''
+    ConstantOffset:  0
+    InitStr:         ''
+    NewHostVarName:  ''
+    BlockLevelFormatFlag: false
+  
+MainSourceFilesDigest:
+  - MainSourceFile:  '/path/to/src/foo.cpp'
+    Digest:          e2636fb8d174ac319083b0306294d3bd
+    HasCUDASyntax:   true
+  - MainSourceFile:  '/path/to/src/baz.cpp'
+    Digest:          991d7e4825fc597205bf68f2eda27acd
+    HasCUDASyntax:   true
+  - MainSourceFile:  '/path/to/src/layer.cudnn.cpp'
+    Digest:          991d7e4825fc597205bf68f2eda27ace
+    HasCUDASyntax:   true
+  - MainSourceFile:  '/path/to/chash/chash.c'
+    Digest:          9f681b5a2cafa7dff67d7e4ac13d1302
+    HasCUDASyntax:   true
+  - MainSourceFile:  '/path/to/cchash/cchash.cc'
+    Digest:          9f681b5a2cafa7dff67d7e4ac13d1303
+    HasCUDASyntax:   true
+  - MainSourceFile:  '/path/to/cxxhash/cxxhash.cxx'
+    Digest:          9f681b5a2cafa7dff67d7e4ac13d1304
+    HasCUDASyntax:   true
+  - MainSourceFile:  '/path/to/cxxhash/cpphash.cpp'
+    Digest:          9f681b5a2cafa7dff67d7e4ac13d1305
+    HasCUDASyntax:   true
+  - MainSourceFile:  '/path/to/chash/chash.c'
+    Digest:          9f681b5a2cafa7dff67d7e4ac13d1302
+    HasCUDASyntax:   false
+DpctVersion:     19.0.0
+SDKVersionMajor: '12'
+SDKVersionMinor: '0'
+MainHelperFileName: ''
+USMLevel:        ''
+FeatureMap:      {}
+CompileTargets:  {}
+OptionMap:
+  AnalysisScopePath:
+    Value:           '/path/to//target'
+    Specified:       false
+  AsyncHandler:
+    Value:           'false'
+    Specified:       false
+  BuildScript:
+    Value:           '1'
+    Specified:       true
+  CodePinEnabled:
+    Value:           'false'
+    Specified:       false
+  CommentsEnabled:
+    Value:           'false'
+    Specified:       false
+  CompilationsDir:
+    Value:           '/path/to/build'
+    Specified:       true
+  CtadEnabled:
+    Value:           'false'
+    Specified:       false
+  EnablepProfiling:
+    Value:           'true'
+    Specified:       true
+  ExperimentalFlag:
+    Value:           '16384'
+    Specified:       true
+  ExplicitNamespace:
+    Value:           '20'
+    Specified:       false
+  ExtensionDDFlag:
+    Value:           '0'
+    Specified:       false
+  ExtensionDEFlag:
+    Value:           '4294967295'
+    Specified:       false
+  RuleFile:
+    Value:           ''
+    ValueVec:
+      - '/path/to/cmake_script_migration_rule.yaml'
+    Specified:       false
+  SyclNamedLambda:
+    Value:           'false'
+    Specified:       false
+  UsmLevel:
+    Value:           '1'
+    Specified:       false
+...

clang/test/dpct/python_migration/case_006/case_006_torch_cuda.cpp

@@ -7,3 +7,9 @@
 
 // RUN: dpct -in-root ./ -out-root out_ipex ./input.py --migrate-build-script-only --migrate-build-script=Python --rule-file=%T/../../../../../../../extensions/python_rules/python_build_script_migration_rule_ipex.yaml
 // RUN: diff --strip-trailing-cr %S/expected.py %T/out_ipex/input.py >> %T/diff.txt
+
+// RUN: mkdir -p %T/out_pytorch1
+// RUN: cp %S/MainSourceFiles.yaml %T/out_pytorch1/MainSourceFiles.yaml
+// RUN: dpct -in-root ./ -out-root out_pytorch1 ./input.py --migrate-build-script-only --migrate-build-script=Python --rule-file=%T/../../../../../../../extensions/python_rules/python_build_script_migration_rule_pytorch.yaml
+// dpct.cmake should not be copied into %T/out_pytorch1 direcotrtory.
+// RUN: test ! -f %T/out_pytorch1/dpct.cmake

docs/dev_guide/api-mapping-status/ASM_API_migration_status.csv

@@ -15,14 +15,14 @@ bfind,NO,
 bmsk,NO,
 bra,NO,
 brev,YES,
-brkpt,NO,
+brkpt,YES,
 brx,NO,
 call,NO,
 clz,YES,
 cnot,YES,
 copysign,YES,
 cos,YES,
-cp,NO,
+cp,YES, Partial
 createpolicy,NO,
 cvt,YES, Partial
 cvta,NO,
@@ -69,7 +69,7 @@ popc,YES,
 prefetch,YES, Partial
 prmt,YES,
 rcp,YES,
-red,NO,
+red,YES, Partial
 redux,NO,
 rem,YES,
 ret,YES,

docs/dev_guide/api-mapping-status/CUB_API_migration_status.csv

@@ -31,8 +31,8 @@ cub::BlockDiscontinuity::FlagTails,NO,
 cub::BlockDiscontinuity::FlagHeadsAndTails,NO,
 cub::BlockExchange::StripedToBlocked,YES,
 cub::BlockExchange::BlockedToStriped,YES,
-cub::BlockExchange::WarpStripedToBlocked,NO,
-cub::BlockExchange::BlockedToWarpStriped,NO,
+cub::BlockExchange::WarpStripedToBlocked,YES,
+cub::BlockExchange::BlockedToWarpStriped,YES,
 cub::BlockExchange::ScatterToBlocked,YES,
 cub::BlockExchange::ScatterToStriped,YES,
 cub::BlockExchange::ScatterToStripedGuarded,NO,
@@ -160,8 +160,8 @@ cub::StoreDirectBlocked,YES,
 cub::StoreDirectBlockedVectorized,NO,
 cub::LoadDirectStriped,YES,
 cub::StoreDirectStriped,YES,
-cub::LoadDirectWarpStriped,NO,
-cub::StoreDirectWarpStriped,NO,
+cub::LoadDirectWarpStriped,YES,
+cub::StoreDirectWarpStriped,YES,
 cub::SHR_ADD,YES,
 cub::SHL_ADD,YES,
 cub::BFE,YES,