[HIPIFY][#207][fix] Translate all preprocessor's conditional blocks

+ Start to translate preprocessor's false conditional blocks too:
  based on clang's https://reviews.llvm.org/D66597;
  available only starting from LLVM 10.0 or trunk.
+ Option -skip-excluded-preprocessor-conditional-blocks for skipping excluded conditional blocks:
  the default behavior for hipify-clang built with LLVM < 10.0;
  false by default for hipify-clang built with LLVM 10 or trunk.
+ Add 4 preprocessor unit tests, 2 of which are LLVM 10.0 only
+ Update couple of existing tests by setting -skip-excluded-preprocessor-conditional-blocks option:
  update lit testing accordingly


[ROCm/clr commit: 32f22a0117]

projects/clr/hipamd/hipify-clang/src/ArgParse.cpp

@@ -138,4 +138,9 @@ cl::list<std::string> MacroNames("D",
   cl::Prefix,
   cl::cat(ToolTemplateCategory));
 
+cl::opt<bool> SkipExcludedPPConditionalBlocks("skip-excluded-preprocessor-conditional-blocks",
+  cl::desc("Enable default preprocessor behaviour by skipping undefined conditional blocks"),
+  cl::value_desc("skip-excluded-preprocessor-conditional-blocks"),
+  cl::cat(ToolTemplateCategory));
+
 cl::extrahelp CommonHelp(ct::CommonOptionsParser::HelpMessage);

projects/clr/hipamd/hipify-clang/src/ArgParse.h

@@ -52,3 +52,4 @@ extern cl::opt<bool> Examine;
 extern cl::extrahelp CommonHelp;
 extern cl::opt<bool> TranslateToRoc;
 extern cl::opt<bool> DashDash;
+extern cl::opt<bool> SkipExcludedPPConditionalBlocks;

projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp

@@ -484,6 +484,11 @@ public:
 
 }
 
+bool HipifyAction::BeginInvocation(clang::CompilerInstance &CI) {
+  llcompat::RetainExcludedConditionalBlocks(CI);
+  return true;
+}
+
 void HipifyAction::ExecuteAction() {
   clang::Preprocessor& PP = getCompilerInstance().getPreprocessor();
   clang::SourceManager& SM = getCompilerInstance().getSourceManager();

projects/clr/hipamd/hipify-clang/src/HipifyAction.h

@@ -91,6 +91,8 @@ protected:
   void insertReplacement(const ct::Replacement& rep, const clang::FullSourceLoc& fullSL);
   // FrontendAction entry point.
   void ExecuteAction() override;
+  // Callback before starting processing a single input; used by hipify-clang for setting Preprocessor options.
+  bool BeginInvocation(clang::CompilerInstance &CI) override;
   // Called at the start of each new file to process.
   void EndSourceFileAction() override;
   // MatchCallback API entry point. Called by the AST visitor while searching the AST for things we registered an interest for.

projects/clr/hipamd/hipify-clang/src/LLVMCompat.cpp

@@ -20,8 +20,13 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
 
+#include "ArgParse.h"
 #include "LLVMCompat.h"
 #include "llvm/Support/Path.h"
+#include "clang/Lex/PreprocessorOptions.h"
+#include "clang/Frontend/CompilerInstance.h"
+
+const std::string sHipify = "[HIPIFY] ", sConflict = "conflict: ", sError = "error: ", sWarning = "warning: ";
 
 namespace llcompat {
 
@@ -122,4 +127,20 @@ bool pragma_once_outside_header() {
 #endif
 }
 
+void RetainExcludedConditionalBlocks(clang::CompilerInstance &CI) {
+#if LLVM_VERSION_MAJOR > 9
+  clang::PreprocessorOptions &PPOpts = CI.getPreprocessorOpts();
+  PPOpts.RetainExcludedConditionalBlocks = !SkipExcludedPPConditionalBlocks;
+#endif
+}
+
+bool CheckCompatibility() {
+#if LLVM_VERSION_MAJOR < 10
+  if (SkipExcludedPPConditionalBlocks) {
+    llvm::errs() << "\n" << sHipify << sWarning << "Option '" << SkipExcludedPPConditionalBlocks.ArgStr.str() << "' is supported starting from LLVM version 10.0\n";
+  }
+#endif
+  return true;
+}
+
 } // namespace llcompat

projects/clr/hipamd/hipify-clang/src/LLVMCompat.h

@@ -30,6 +30,8 @@ THE SOFTWARE.
 
 namespace ct = clang::tooling;
 
+extern const std::string sHipify, sConflict, sError, sWarning;
+
 // Things for papering over the differences between different LLVM versions.
 
 namespace llcompat {
@@ -83,4 +85,8 @@ std::error_code real_path(const Twine &path, SmallVectorImpl<char> &output,
 
 bool pragma_once_outside_header();
 
+void RetainExcludedConditionalBlocks(clang::CompilerInstance &CI);
+
+bool CheckCompatibility();
+
 } // namespace llcompat

projects/clr/hipamd/hipify-clang/src/main.cpp

@@ -37,7 +37,6 @@ THE SOFTWARE.
 
 #define DEBUG_TYPE "cuda2hip"
 
-std::string sHipify = "[HIPIFY] ", sConflict = "conflict: ", sError = "error: ";
 namespace ct = clang::tooling;
 
 std::string getAbsoluteFilePath(const std::string& sFile, std::error_code& EC) {
@@ -240,13 +239,16 @@ int main(int argc, const char **argv) {
   }
   llcompat::PrintStackTraceOnErrorSignal();
   ct::CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::ZeroOrMore);
+  if (!llcompat::CheckCompatibility()) {
+    return 1;
+  }
   std::vector<std::string> fileSources = OptionsParser.getSourcePathList();
   if (fileSources.empty() && !GeneratePerl && !GeneratePython) {
-    llvm::errs() << "\n" << sHipify << sError << "Must specify at least 1 positional argument for source file." << "\n";
+    llvm::errs() << "\n" << sHipify << sError << "Must specify at least 1 positional argument for source file" << "\n";
     return 1;
   }
   if (!generatePerl(GeneratePerl)) {
-    llvm::errs() << "\n" << sHipify << sError << "hipify-perl generating failed." << "\n";
+    llvm::errs() << "\n" << sHipify << sError << "hipify-perl generating failed" << "\n";
     return 1;
   }
   bool bToRoc = TranslateToRoc;
@@ -254,7 +256,7 @@ int main(int argc, const char **argv) {
   bool bToPython = generatePython(GeneratePython);
   TranslateToRoc = bToRoc;
   if (!bToPython) {
-    llvm::errs() << "\n" << sHipify << sError << "hipify-python generating failed." << "\n";
+    llvm::errs() << "\n" << sHipify << sError << "hipify-python generating failed" << "\n";
     return 1;
   }
   if (fileSources.empty()) {
@@ -268,15 +270,15 @@ int main(int argc, const char **argv) {
   }
   if (!dst.empty()) {
     if (fileSources.size() > 1) {
-      llvm::errs() << sHipify << sConflict << "-o and multiple source files are specified.\n";
+      llvm::errs() << sHipify << sConflict << "-o and multiple source files are specified\n";
       return 1;
     }
     if (Inplace) {
-      llvm::errs() << sHipify << sConflict << "both -o and -inplace options are specified.\n";
+      llvm::errs() << sHipify << sConflict << "both -o and -inplace options are specified\n";
       return 1;
     }
     if (NoOutput) {
-      llvm::errs() << sHipify << sConflict << "both -no-output and -o options are specified.\n";
+      llvm::errs() << sHipify << sConflict << "both -no-output and -o options are specified\n";
       return 1;
     }
     if (!dstDir.empty()) {
@@ -284,11 +286,11 @@ int main(int argc, const char **argv) {
     }
   }
   if (NoOutput && Inplace) {
-    llvm::errs() << sHipify << sConflict << "both -no-output and -inplace options are specified.\n";
+    llvm::errs() << sHipify << sConflict << "both -no-output and -inplace options are specified\n";
     return 1;
   }
   if (!dstDir.empty() && Inplace) {
-    llvm::errs() << sHipify << sConflict << "both -o-dir and -inplace options are specified.\n";
+    llvm::errs() << sHipify << sConflict << "both -o-dir and -inplace options are specified\n";
     return 1;
   }
   if (Examine) {
@@ -425,8 +427,8 @@ int main(int argc, const char **argv) {
     // Hipify _all_ the things!
     if (Tool.runAndSave(&actionFactory)) {
       currentStat.hasErrors = true;
-      LLVM_DEBUG(llvm::dbgs() << "Skipped some replacements.\n");
       Result = 1;
+      LLVM_DEBUG(llvm::dbgs() << "Skipped some replacements.\n");
     }
     // Copy the tmpfile to the output
     if (!NoOutput && !currentStat.hasErrors) {

projects/clr/hipamd/tests/hipify-clang/lit.cfg

@@ -42,6 +42,10 @@ if config.cuda_version_major < 10:
     config.excludes.append('cuSPARSE_10.cu')
     config.excludes.append('cuSPARSE_11.cu')
 
+if config.llvm_version_major < 10:
+    config.excludes.append('pp_if_else_conditionals_LLVM_10.cu')
+    config.excludes.append('pp_if_else_conditionals_01_LLVM_10.cu')
+
 # name: The name of this test suite.
 config.name = 'hipify'
 

projects/clr/hipamd/tests/hipify-clang/lit.site.cfg.in

@@ -3,6 +3,7 @@ import os
 
 config.pointer_size = @CMAKE_SIZEOF_VOID_P@
 config.llvm_version = "@LLVM_PACKAGE_VERSION@"
+config.llvm_version_major = int("@LLVM_VERSION_MAJOR@")
 config.llvm_tools_dir = "@LLVM_TOOLS_BINARY_DIR@"
 config.obj_root = "@CMAKE_CURRENT_BINARY_DIR@"
 config.cuda_root = "@CUDA_TOOLKIT_ROOT_DIR@"

projects/clr/hipamd/tests/hipify-clang/unit_tests/libraries/cuComplex/cuComplex_Julia.cu

@@ -1,4 +1,4 @@
-// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
+// RUN: %run_test hipify "%s" "%t" %hipify_args "--skip-excluded-preprocessor-conditional-blocks" %clang_args
 
 // CHECK: #include <hip/hip_runtime.h>
 // CHECK: #include "hip/hip_complex.h"

projects/clr/hipamd/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu

@@ -1,4 +1,4 @@
-// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
+// RUN: %run_test hipify "%s" "%t" %hipify_args "--skip-excluded-preprocessor-conditional-blocks" %clang_args
 #include <stdio.h>
 #include <stdlib.h>
 #include <assert.h>

projects/clr/hipamd/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals.cu

@@ -0,0 +1,30 @@
+// RUN: %run_test hipify "%s" "%t" %hipify_args "--skip-excluded-preprocessor-conditional-blocks" %clang_args
+// CHECK: #include <hip/hip_runtime.h>
+
+#include <cuda.h>
+
+__global__ void axpy_kernel(float a, float* x, float* y) {
+  y[threadIdx.x] = a * x[threadIdx.x];
+}
+
+void axpy(float a, float* x, float* y) {
+
+#ifdef SOME_MACRO
+  // CHECK: axpy_kernel <<<1, 1>>> (a, y, x);
+  axpy_kernel <<<1, 1>>> (a, y, x);
+#endif
+
+#ifndef SOME_MACRO
+  // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(2), 0, 0, a, y, x);
+  axpy_kernel <<<1, 2>>> (a, y, x);
+#endif
+
+#ifdef SOME_MACRO
+  // CHECK: axpy_kernel <<<1, 3>>> (a, y, x);
+  axpy_kernel <<<1, 3>>> (a, y, x);
+#else
+  // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(4), 0, 0, a, x, y);
+  axpy_kernel <<<1, 4>>> (a, x, y);
+#endif
+
+}

projects/clr/hipamd/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals_01.cu

@@ -0,0 +1,52 @@
+// RUN: %run_test hipify "%s" "%t" %hipify_args "--skip-excluded-preprocessor-conditional-blocks" %clang_args
+// CHECK: #include <hip/hip_runtime.h>
+
+__global__ void axpy_kernel(float a, float* x, float* y) {
+  y[threadIdx.x] = a * x[threadIdx.x];
+}
+
+void axpy(float a, float* x, float* y) {
+float* y_new = nullptr;
+#ifdef SOME_MACRO
+  y_new = x;
+  // CHECK: axpy_kernel <<<1, 1>>> (a, y_new, x);
+  axpy_kernel <<<1, 1>>> (a, y_new, x);
+#endif
+
+#ifndef SOME_MACRO
+  y_new = y;
+  // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(2), 0, 0, a, y_new, x);
+  axpy_kernel <<<1, 2>>> (a, y_new, x);
+#endif
+
+#ifdef SOME_MACRO
+  // CHECK: axpy_kernel <<<1, 3>>> (a, y, x);
+  axpy_kernel <<<1, 3>>> (a, y, x);
+#else
+  // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(4), 0, 0, a, x, y);
+  axpy_kernel <<<1, 4>>> (a, x, y);
+#endif
+
+#ifdef SOME_MACRO
+  // CHECK: axpy_kernel <<<1, 5>>> (a, y, x);
+  axpy_kernel <<<1, 5>>> (a, y, x);
+#elif defined SOME_MACRO_1
+  // CHECK: axpy_kernel <<<1, 6>>> (a, x, y);
+  axpy_kernel <<<1, 6>>> (a, x, y);
+#else
+  // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(7), 0, 0, a, x, y);
+  axpy_kernel <<<1, 7>>> (a, x, y);
+#endif
+
+#ifndef SOME_MACRO
+  // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(8), 0, 0, a, y, x);
+  axpy_kernel <<<1, 8>>> (a, y, x);
+#elif !defined(SOME_MACRO_1)
+  // CHECK: axpy_kernel <<<1, 9>>> (a, x, y);
+  axpy_kernel <<<1, 9>>> (a, x, y);
+#else
+  // CHECK: axpy_kernel <<<1, 10>>> (a, x, y);
+  axpy_kernel <<<1, 10>>> (a, x, y);
+#endif
+
+}

projects/clr/hipamd/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals_01_LLVM_10.cu

@@ -0,0 +1,52 @@
+// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
+// CHECK: #include <hip/hip_runtime.h>
+
+__global__ void axpy_kernel(float a, float* x, float* y) {
+  y[threadIdx.x] = a * x[threadIdx.x];
+}
+
+void axpy(float a, float* x, float* y) {
+float* y_new = nullptr;
+#ifdef SOME_MACRO
+  y_new = x;
+  // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(1), 0, 0, a, y_new, x);
+  axpy_kernel <<<1, 1>>> (a, y_new, x);
+#endif
+
+#ifndef SOME_MACRO
+  y_new = y;
+  // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(2), 0, 0, a, y_new, x);
+  axpy_kernel <<<1, 2>>> (a, y_new, x);
+#endif
+
+#ifdef SOME_MACRO
+  // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(3), 0, 0, a, y, x);
+  axpy_kernel <<<1, 3>>> (a, y, x);
+#else
+  // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(4), 0, 0, a, x, y);
+  axpy_kernel <<<1, 4>>> (a, x, y);
+#endif
+
+#ifdef SOME_MACRO
+  // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(5), 0, 0, a, y, x);
+  axpy_kernel <<<1, 5>>> (a, y, x);
+#elif defined SOME_MACRO_1
+  // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(6), 0, 0, a, x, y);
+  axpy_kernel <<<1, 6>>> (a, x, y);
+#else
+  // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(7), 0, 0, a, x, y);
+  axpy_kernel <<<1, 7>>> (a, x, y);
+#endif
+
+#ifndef SOME_MACRO
+  // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(8), 0, 0, a, y, x);
+  axpy_kernel <<<1, 8>>> (a, y, x);
+#elif !defined(SOME_MACRO_1)
+  // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(9), 0, 0, a, x, y);
+  axpy_kernel <<<1, 9>>> (a, x, y);
+#else
+  // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(10), 0, 0, a, x, y);
+  axpy_kernel <<<1, 10>>> (a, x, y);
+#endif
+
+}

projects/clr/hipamd/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals_LLVM_10.cu

@@ -0,0 +1,30 @@
+// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
+// CHECK: #include <hip/hip_runtime.h>
+
+#include <cuda.h>
+
+__global__ void axpy_kernel(float a, float* x, float* y) {
+  y[threadIdx.x] = a * x[threadIdx.x];
+}
+
+void axpy(float a, float* x, float* y) {
+
+#ifdef SOME_MACRO
+  // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(1), 0, 0, a, y, x);
+  axpy_kernel <<<1, 1>>> (a, y, x);
+#endif
+
+#ifndef SOME_MACRO
+  // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(2), 0, 0, a, y, x);
+  axpy_kernel <<<1, 2>>> (a, y, x);
+#endif
+
+#ifdef SOME_MACRO
+  // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(3), 0, 0, a, y, x);
+  axpy_kernel <<<1, 3>>> (a, y, x);
+#else
+  // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(4), 0, 0, a, x, y);
+  axpy_kernel <<<1, 4>>> (a, x, y);
+#endif
+
+}