[clang-repl][CUDA] Add an unit test for interactive CUDA

[argentite]

Summary

This tests interactive CUDA support in clang interpreter and should pass
in case CUDA runtime is not available or no GPUs are available.

[tra]

clang/unittests/Interpreter/InteractiveCudaTest.cpp

93

This could be a bit of a problem.

There may be multiple CUDA SDK versions that may be installed on a system at any given time and the libcudart.so you pick here may not be the one you want.
E.g it may be from a recent CUDA version which is not supported by NVIDIA drivers yet.

I think you may need a way to let the user override CUDA SDK (or libcudart.so) location explicitly. I guess they could do that via LD_LIBRARY_PATH, but for the CUDA compilation in general, knowing CUDA SDK path is essential, as it does affect various compilation options set by the driver.

[argentite]

clang/unittests/Interpreter/InteractiveCudaTest.cpp

93

Yes, this probably would be an issue. It is currently possible to override the CUDA path with a command line argument in clang-repl. But I am not sure what we can do inside a test.

[tra]

clang/unittests/Interpreter/InteractiveCudaTest.cpp

93

To me it looks like CUDA location should be detected/set at the configuration time and then propagated to the individual tests that need that info.
CMake has cuda detection mechanisms that could be used for that purpose.
They are a bit of a pain to use in practice (I'm still not sure what's the reliable way to do it), but it's as close to the 'standard' way of doing it as we have at the moment.
I believe libc and mlir subtrees in LLVM are already using this mechanism. E.g https://github.com/llvm/llvm-project/blob/main/libc/utils/gpu/loader/CMakeLists.txt#L16

[v.g.vassilev]

@argentite what is the fate of this patch? Should we move forward with it?

[v.g.vassilev]

@argentite ping.

[argentite]

clang/unittests/Interpreter/InteractiveCudaTest.cpp

93

This version uses CMake to detect the paths as in the example above. I guess this would tie the test to the build system but it should probably be fine for most cases. Unfortunately I don't think we can use the same trick for determining the path in ClangREPL because the interpreter is likely to run on a different system.

clang/unittests/Interpreter/CMakeLists.txt

 set(LLVM_LINK_COMPONENTS
   ${LLVM_TARGETS_TO_BUILD}
   Core
   OrcJIT
   Support
   TargetParser
   )
 
+find_package(CUDAToolkit)
+if(CUDAToolkit_FOUND)
+  get_property(LIBCUDART_PATH TARGET CUDA::cudart PROPERTY LOCATION)
+  add_compile_definitions(LIBCUDART_PATH="${LIBCUDART_PATH}")
+endif()
+
 add_clang_unittest(ClangReplInterpreterTests
   IncrementalProcessingTest.cpp
+  InteractiveCudaTest.cpp
   InterpreterTest.cpp
   )
 target_link_libraries(ClangReplInterpreterTests PUBLIC
   clangAST
   clangBasic
   clangInterpreter
   clangFrontend
   clangSema
   )
 
 # Exceptions on Windows are not yet supported.
 if(NOT WIN32)
   add_subdirectory(ExceptionTests)
 endif()
 
 export_executable_symbols(ClangReplInterpreterTests)

clang/unittests/Interpreter/InteractiveCudaTest.cpp

+//===- unittests/Interpreter/CudaTest.cpp --- Interactive CUDA tests ------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// Unit tests for interactive CUDA in Clang interpreter
+//
+//===----------------------------------------------------------------------===//
+
+#include "clang/Basic/Version.h"
+#include "clang/Config/config.h"
+
+#include "clang/Frontend/CompilerInstance.h"
+#include "clang/Interpreter/Interpreter.h"
+
+#include "llvm/Support/TargetSelect.h"
+
+#include "gmock/gmock.h"
+#include "gtest/gtest.h"
+
+namespace {
+
+std::string MakeResourcesPath() {
+  using namespace llvm;
+#ifdef LLVM_BINARY_DIR
+  StringRef Dir = LLVM_BINARY_DIR;
+#else
+  // Dir is bin/ or lib/, depending on where BinaryPath is.
+  void *MainAddr = (void *)(intptr_t)MakeResourcesPath;
+  std::string BinaryPath =
+      llvm::sys::fs::getMainExecutable(/*Argv0=*/nullptr, MainAddr);
+
+  // build/tools/clang/unittests/Interpreter/Executable -> build/
+  StringRef Dir = sys::path::parent_path(BinaryPath);
+
+  Dir = sys::path::parent_path(Dir);
+  Dir = sys::path::parent_path(Dir);
+  Dir = sys::path::parent_path(Dir);
+  Dir = sys::path::parent_path(Dir);
+#endif // LLVM_BINARY_DIR
+  SmallString<128> P(Dir);
+  sys::path::append(P, CLANG_INSTALL_LIBDIR_BASENAME, "clang",
+                    CLANG_VERSION_MAJOR_STRING);
+  return P.str().str();
+}
+
+static std::unique_ptr<clang::Interpreter>
+createInterpreter(const std::vector<const char *> &ExtraArgs = {}) {
+  static bool firstrun = true;
+  if (firstrun) {
+    llvm::InitializeAllTargetInfos();
+    llvm::InitializeAllTargets();
+    llvm::InitializeAllTargetMCs();
+    llvm::InitializeAllAsmPrinters();
+
+    firstrun = false;
+  }
+
+  clang::IncrementalCompilerBuilder CB;
+
+  // Help find cuda's runtime headers.
+  std::string ResourceDir = MakeResourcesPath();
+
+  std::vector Args = {"-resource-dir", ResourceDir.c_str(), "-std=c++20"};
+  Args.insert(Args.end(), ExtraArgs.begin(), ExtraArgs.end());
+  CB.SetCompilerArgs(Args);
+
+  // Create the device code compiler
+  std::unique_ptr<clang::CompilerInstance> DeviceCI;
+  CB.SetOffloadArch("sm_35");
+  DeviceCI = cantFail(CB.CreateCudaDevice());
+
+  std::unique_ptr<clang::CompilerInstance> CI;
+  CI = cantFail(CB.CreateCudaHost());
+
+  auto Interp = cantFail(
+      clang::Interpreter::createWithCUDA(std::move(CI), std::move(DeviceCI)));
+
+  return Interp;
+}
+
+enum {
+  // Defined in CUDA Runtime API
+  cudaErrorNoDevice = 100,
+};
+
+TEST(InteractiveCudaTest, Sanity) {
+  std::unique_ptr<clang::Interpreter> Interp = createInterpreter();
+
+#ifdef LIBCUDART_PATH
+  auto Err = Interp->LoadDynamicLibrary(LIBCUDART_PATH);
+  if (Err) { // CUDA runtime is not usable, cannot continue testing
+    consumeError(std::move(Err));
+    return;
+  }
+#else
+  return;
+#endif
+
+  // Check if we have any GPU for test
+  int CudaError = 0;
+  auto GpuCheckCommand = std::string(R"(
+    int device_id = -1;
+    int *error = (int *))" + std::to_string((uintptr_t)&CudaError) +
+                                     R"(;
+    *error = cudaGetDevice(&device_id);
+  )");
+  cantFail(Interp->ParseAndExecute(GpuCheckCommand));
+  if (CudaError == cudaErrorNoDevice) {
+    // No GPU is available on this machine, cannot continue testing
+    return;
+  }
+  ASSERT_EQ(CudaError, 0);
+
+  int HostSum = 0;
+  auto Command1 = std::string(R"(
+    __host__ __device__ inline int sum(int a, int b){ return a + b; }
+    __global__ void kernel(int * output){ *output = sum(40,2); }
+    int *hostsum = (int *) )") +
+                  std::to_string((uintptr_t)&HostSum) +
+                  R"(;
+    *hostsum = sum(41,1);)";
+  cantFail(Interp->ParseAndExecute(Command1));
+
+  int DeviceSum = 0;
+  auto Command2 = std::string(R"(
+    int *devicesum = (int *))" +
+                              std::to_string((uintptr_t)&DeviceSum) +
+                              R"(;
+    int *deviceVar;
+    *error |= cudaMalloc((void **) &deviceVar, sizeof(int));
+    kernel<<<1,1>>>(deviceVar);
+    *error |= cudaGetLastError();
+    *error |= cudaMemcpy(devicesum, deviceVar, sizeof(int), cudaMemcpyDeviceToHost);
+    *error |= cudaGetLastError();
+)");
+  cantFail(Interp->ParseAndExecute(Command2));
+
+  ASSERT_EQ(HostSum, 42);
+  ASSERT_EQ(DeviceSum, 42);
+  ASSERT_EQ(CudaError, 0);
+}
+
+} // end anonymous namespace