Rust and cuda clang support

Cargo.toml

@@ -1,12 +1,13 @@
 [workspace]
 members = [
-        "rust/libceed",
-        "rust/libceed-sys",
-        "examples/rust/ex1-volume",
-        "examples/rust/ex1-volume-vector",
-        "examples/rust/ex2-surface",
-        "examples/rust/ex2-surface-vector",
-        "examples/rust/ex3-volume",
-        "examples/rust/ex3-volume-vector",
-        "examples/rust/mesh",
+    "rust/libceed",
+    "rust/libceed-sys",
+    "examples/rust/ex1-volume",
+    "examples/rust/ex1-volume-vector",
+    "examples/rust/ex2-surface",
+    "examples/rust/ex2-surface-vector",
+    "examples/rust/ex3-volume",
+    "examples/rust/ex3-volume-vector",
+    "examples/rust/mesh",
 ]
+exclude = ["examples/ceed/ex4-volume-rs"]

Makefile

@@ -134,17 +134,17 @@ OPT.clang               := $(OPT.gcc)
 OPT.icc                 := $(OPT.gcc)
 OPT.oneAPI              := $(OPT.clang)
 OPT.emcc                :=
-CFLAGS.gcc              := $(if $(STATIC),,-fPIC) -std=c11 -Wall -Wextra -Wno-unused-parameter -MMD -MP
+CFLAGS.gcc              := $(if $(STATIC),,-fPIC) -std=c17 -Wall -Wextra -Wno-unused-parameter -MMD -MP
 CFLAGS.clang            := $(CFLAGS.gcc)
 CFLAGS.icc              := $(CFLAGS.gcc)
 CFLAGS.oneAPI           := $(CFLAGS.clang)
 CFLAGS.XL               := $(if $(STATIC),,-qpic) -MMD
 CFLAGS.emcc             := $(CFLAGS.clang)
-CXXFLAGS.gcc            := $(if $(STATIC),,-fPIC) -std=c++11 -Wall -Wextra -Wno-unused-parameter -MMD -MP
+CXXFLAGS.gcc            := $(if $(STATIC),,-fPIC) -std=c++17 -Wall -Wextra -Wno-unused-parameter -MMD -MP
 CXXFLAGS.clang          := $(CXXFLAGS.gcc)
 CXXFLAGS.icc            := $(CXXFLAGS.gcc)
 CXXFLAGS.oneAPI         := $(CXXFLAGS.clang)
-CXXFLAGS.XL             := $(if $(STATIC),,-qpic) -std=c++11 -MMD
+CXXFLAGS.XL             := $(if $(STATIC),,-qpic) -std=c++17 -MMD
 CXXFLAGS.emcc           := $(CXXFLAGS.clang)
 FFLAGS.GNU              := $(if $(STATIC),,-fPIC) -cpp -Wall -Wextra -Wno-unused-parameter -Wno-unused-dummy-argument -MMD -MP
 FFLAGS.ifort            := $(if $(STATIC),,-fPIC) -cpp

backends/cuda/ceed-cuda-compile.cpp

@@ -15,7 +15,10 @@
 #include <stdarg.h>
 #include <string.h>
 
+#include <filesystem>
 #include <sstream>
+#include <iostream>
+#include <fstream>
 
 #include "ceed-cuda-common.h"
 
@@ -34,6 +37,10 @@
 //------------------------------------------------------------------------------
 // Compile CUDA kernel
 //------------------------------------------------------------------------------
+using std::ifstream;
+using std::ofstream;
+using std::ostringstream;
+
 static int CeedCompileCore_Cuda(Ceed ceed, const char *source, const bool throw_error, bool *is_compile_good, CUmodule *module,
                                 const CeedInt num_defines, va_list args) {
   size_t                ptx_size;
@@ -45,10 +52,23 @@ static int CeedCompileCore_Cuda(Ceed ceed, const char *source, const bool throw_
   struct cudaDeviceProp prop;
   Ceed_Cuda            *ceed_data;
 
+
   cudaFree(0);  // Make sure a Context exists for nvrtc
 
   std::ostringstream code;
 
+  bool usingClang;
+  CeedCallBackend(CeedGetIsClang(ceed, &usingClang));
+
+  if(CeedDebugFlag(ceed)){
+      if(usingClang){
+          CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "Compiling cuda with Clang backend (for Rust Qfunction support)");
+      } else {
+          CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "Compiling cuda with NVRTC backend (Does not support Rust Qfunctions). To use Clang, set the environmental variable CUDA_CLANG to 1");
+      }
+  }
+
+
   // Get kernel specific options, such as kernel constants
   if (num_defines > 0) {
     char *name;
@@ -63,7 +83,6 @@ static int CeedCompileCore_Cuda(Ceed ceed, const char *source, const bool throw_
 
   // Standard libCEED definitions for CUDA backends
   code << "#include <ceed/jit-source/cuda/cuda-jit.h>\n\n";
-
   // Non-macro options
   CeedCallBackend(CeedCalloc(num_opts, &opts));
   opts[0] = "-default-device";
@@ -116,67 +135,178 @@ static int CeedCompileCore_Cuda(Ceed ceed, const char *source, const bool throw_
   code << source;
 
   // Create Program
-  CeedCallNvrtc(ceed, nvrtcCreateProgram(&prog, code.str().c_str(), NULL, 0, NULL, NULL));
 
   // Compile kernel
   CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "---------- ATTEMPTING TO COMPILE JIT SOURCE ----------\n");
   CeedDebug(ceed, "Source:\n%s\n", code.str().c_str());
-  CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "---------- END OF JIT SOURCE ----------\n");
-  if (CeedDebugFlag(ceed)) {
-    // LCOV_EXCL_START
-    CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "---------- JiT COMPILER OPTIONS ----------\n");
-    for (CeedInt i = 0; i < num_opts + num_jit_source_dirs + num_jit_defines; i++) {
-      CeedDebug(ceed, "Option %d: %s", i, opts[i]);
-    }
-    CeedDebug(ceed, "");
-    CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "---------- END OF JiT COMPILER OPTIONS ----------\n");
-    // LCOV_EXCL_STOP
-  }
-  nvrtcResult result = nvrtcCompileProgram(prog, num_opts + num_jit_source_dirs + num_jit_defines, opts);
+  CeedDebug256(ceed, CEED_DEBUG_COLOR_ERROR, "---------- END OF JIT SOURCE ----------\n");
 
-  for (CeedInt i = 0; i < num_jit_source_dirs; i++) {
-    CeedCallBackend(CeedFree(&opts[num_opts + i]));
-  }
-  for (CeedInt i = 0; i < num_jit_defines; i++) {
-    CeedCallBackend(CeedFree(&opts[num_opts + num_jit_source_dirs + i]));
-  }
-  CeedCallBackend(CeedFree(&opts));
-  *is_compile_good = result == NVRTC_SUCCESS;
-  if (!*is_compile_good) {
-    char  *log;
-    size_t log_size;
-
-    CeedCallNvrtc(ceed, nvrtcGetProgramLogSize(prog, &log_size));
-    CeedCallBackend(CeedMalloc(log_size, &log));
-    CeedCallNvrtc(ceed, nvrtcGetProgramLog(prog, log));
-    if (throw_error) {
-      return CeedError(ceed, CEED_ERROR_BACKEND, "%s\n%s", nvrtcGetErrorString(result), log);
-    } else {
-      // LCOV_EXCL_START
-      CeedDebug256(ceed, CEED_DEBUG_COLOR_ERROR, "---------- COMPILE ERROR DETECTED ----------\n");
-      CeedDebug(ceed, "Error: %s\nCompile log:\n%s\n", nvrtcGetErrorString(result), log);
-      CeedDebug256(ceed, CEED_DEBUG_COLOR_WARNING, "---------- BACKEND MAY FALLBACK ----------\n");
-      CeedCallBackend(CeedFree(&log));
-      CeedCallNvrtc(ceed, nvrtcDestroyProgram(&prog));
-      return CEED_ERROR_SUCCESS;
-      // LCOV_EXCL_STOP
+  if(!usingClang){
+    CeedCallNvrtc(ceed, nvrtcCreateProgram(&prog, code.str().c_str(), NULL, 0, NULL, NULL));
+
+    if (CeedDebugFlag(ceed)) {
+        // LCOV_EXCL_START
+        CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "---------- JiT COMPILER OPTIONS ----------\n");
+        for (CeedInt i = 0; i < num_opts + num_jit_source_dirs + num_jit_defines; i++) {
+          CeedDebug(ceed, "Option %d: %s", i, opts[i]);
+        }
+        CeedDebug(ceed, "");
+        CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "---------- END OF JiT COMPILER OPTIONS ----------\n");
+        // LCOV_EXCL_STOP
+      }
+
+    nvrtcResult result = nvrtcCompileProgram(prog, num_opts + num_jit_source_dirs + num_jit_defines, opts);
+
+
+    for (CeedInt i = 0; i < num_jit_source_dirs; i++) {
+        CeedCallBackend(CeedFree(&opts[num_opts + i]));
+    }
+    for (CeedInt i = 0; i < num_jit_defines; i++) {
+        CeedCallBackend(CeedFree(&opts[num_opts + num_jit_source_dirs + i]));
+    }
+    CeedCallBackend(CeedFree(&opts));
+    *is_compile_good = result == NVRTC_SUCCESS;
+    if (!*is_compile_good) {
+        char  *log;
+        size_t log_size;
+
+        CeedCallNvrtc(ceed, nvrtcGetProgramLogSize(prog, &log_size));
+        CeedCallBackend(CeedMalloc(log_size, &log));
+        CeedCallNvrtc(ceed, nvrtcGetProgramLog(prog, log));
+        if (throw_error) {
+        return CeedError(ceed, CEED_ERROR_BACKEND, "%s\n%s", nvrtcGetErrorString(result), log);
+        } else {
+        // LCOV_EXCL_START
+        CeedDebug256(ceed, CEED_DEBUG_COLOR_ERROR, "---------- COMPILE ERROR DETECTED ----------\n");
+        CeedDebug(ceed, "Error: %s\nCompile log:\n%s\n", nvrtcGetErrorString(result), log);
+        CeedDebug256(ceed, CEED_DEBUG_COLOR_ERROR, "---------- BACKEND MAY FALLBACK ----------\n");
+        CeedCallBackend(CeedFree(&log));
+        CeedCallNvrtc(ceed, nvrtcDestroyProgram(&prog));
+        return CEED_ERROR_SUCCESS;
+        // LCOV_EXCL_STOP
+        }
     }
-  }
 
-#if CUDA_VERSION >= 11010
-  CeedCallNvrtc(ceed, nvrtcGetCUBINSize(prog, &ptx_size));
-  CeedCallBackend(CeedMalloc(ptx_size, &ptx));
-  CeedCallNvrtc(ceed, nvrtcGetCUBIN(prog, ptx));
-#else
-  CeedCallNvrtc(ceed, nvrtcGetPTXSize(prog, &ptx_size));
-  CeedCallBackend(CeedMalloc(ptx_size, &ptx));
-  CeedCallNvrtc(ceed, nvrtcGetPTX(prog, ptx));
-#endif
-  CeedCallNvrtc(ceed, nvrtcDestroyProgram(&prog));
+    #if CUDA_VERSION >= 11010
+    CeedCallNvrtc(ceed, nvrtcGetCUBINSize(prog, &ptx_size));
+    CeedCallBackend(CeedMalloc(ptx_size, &ptx));
+    CeedCallNvrtc(ceed, nvrtcGetCUBIN(prog, ptx));
+    #else
+    CeedCallNvrtc(ceed, nvrtcGetPTXSize(prog, &ptx_size));
+    CeedCallBackend(CeedMalloc(ptx_size, &ptx));
+    CeedCallNvrtc(ceed, nvrtcGetPTX(prog, ptx));
+    #endif
+    CeedCallNvrtc(ceed, nvrtcDestroyProgram(&prog));
+
+    CeedCallCuda(ceed, cuModuleLoadData(module, ptx));
+    CeedCallBackend(CeedFree(&ptx));
+    return CEED_ERROR_SUCCESS;
+} else{
+        const char* full_filename = "temp-jit.cu";
+        FILE* file = fopen(full_filename, "w");
+        if (!file) {
+            CeedDebug256(ceed, CEED_DEBUG_COLOR_ERROR, "Failed to create file. Write access is required for cuda-clang\n");
+            return 1;
+        }
+        fputs(code.str().c_str(), file);
+        fclose(file);
+
+        // Get rust crate directories
+
+        const char **rust_source_dirs = nullptr;
+        int num_rust_source_dirs = 0;
+
+
+        CeedCallBackend(CeedGetRustSourceRoots(ceed, &num_rust_source_dirs, &rust_source_dirs));
+
+        std::string rust_dirs[10];
+
+        if(num_rust_source_dirs > 0){
+            CeedDebug(ceed, "There are %d source dirs, including %s\n", num_rust_source_dirs, rust_source_dirs[0]);
+        }
+
+        for(int i = 0; i < num_rust_source_dirs; i++){
+            rust_dirs[i] = std::string(rust_source_dirs[i]);
+        }
+
+
+        CeedCallBackend(CeedRestoreRustSourceRoots(ceed, &rust_source_dirs));
+
+        // Compile with rust
+
+        int err;
+        std::string cmd;
+
+        for(int i = 0; i < num_rust_source_dirs; i++){
+            cmd = "cargo +nightly build --release --target nvptx64-nvidia-cuda --manifest-path " + rust_dirs[i] + "/Cargo.toml";
+            err = system(cmd.c_str());
+            if(err){
+                CeedDebug256(ceed, CEED_DEBUG_COLOR_ERROR, "Failed gpu jit task 0 (build rust crates)\nBuilding rust crate %d with command: %s\n", i, cmd.c_str());
+                abort();
+            }
+        }
+
+        cmd = "clang++ -flto=thin --cuda-gpu-arch=sm_" + std::to_string(prop.major) + std::to_string(prop.minor) + " --cuda-device-only -emit-llvm -S temp-jit.cu -o kern.ll ";
+        cmd += opts[4];
+        err = system(cmd.c_str());
+
+        if(err){
+            perror("Failed gpu jit task 1 (compile qfunction source to llvm)\n");
+            abort();
+        }
+
+        cmd = "llvm-link kern.ll --ignore-non-bitcode --internalize --only-needed -S -o kern2.ll ";
+
+        // Searches for .rlib files in rust directoy
+        // Note: this is necessary because rust crate names may not match the folder they are in
+        for(int i = 0; i < num_rust_source_dirs; i++){
+            std::string dir = rust_dirs[i] + "/target/nvptx64-nvidia-cuda/release";
+
+            for(auto p : std::filesystem::directory_iterator(dir)){
+                if (p.path().extension() == ".rlib"){
+                    cmd += p.path().string() + " ";
+                }
+            }
+
+        }
+
+        err = system(cmd.c_str());
+
+        CeedDebug(ceed, "llvm-link command was %s\n", cmd.c_str());
+
+
+        if(err){
+            printf("Failed gpu jit task 2 (link c and rust sources with llvm)\n");
+            abort();
+        }
+
+        err = system("opt --passes internalize,inline kern2.ll -o kern3.bc");
+
+        if(err){
+            printf("Failed gpu jit task 3 (optimize qfunction)\n");
+            abort();
+        }
+
+        err = system(("llc -O0 -mcpu=sm_" + std::to_string(prop.major) + std::to_string(prop.minor) + " kern3.bc -o kern.ptx").c_str());
+
+        if(err){
+            printf("Failed gpu jit task 4 (compile llvm)\n");
+            abort();
+        }
+
+        ifstream           ptxfile("kern.ptx");
+
+        ostringstream      sstr;
+        sstr << ptxfile.rdbuf();
+        auto ptx_data = sstr.str();
+
+        ptx_size = ptx_data.length();
+
+        CeedCallCuda(ceed, cuModuleLoadData(module, ptx_data.c_str()));
+    CeedCallBackend(CeedFree(&ptx_data));
+    }
+    return CEED_ERROR_SUCCESS;
 
-  CeedCallCuda(ceed, cuModuleLoadData(module, ptx));
-  CeedCallBackend(CeedFree(&ptx));
-  return CEED_ERROR_SUCCESS;
 }
 
 int CeedCompile_Cuda(Ceed ceed, const char *source, CUmodule *module, const CeedInt num_defines, ...) {
@@ -206,7 +336,8 @@ int CeedTryCompile_Cuda(Ceed ceed, const char *source, bool *is_compile_good, CU
 // Get CUDA kernel
 //------------------------------------------------------------------------------
 int CeedGetKernel_Cuda(Ceed ceed, CUmodule module, const char *name, CUfunction *kernel) {
-  CeedCallCuda(ceed, cuModuleGetFunction(kernel, module, name));
+
+    CeedCallCuda(ceed, cuModuleGetFunction(kernel, module, name));
   return CEED_ERROR_SUCCESS;
 }
 

examples/ceed/Makefile

@@ -9,8 +9,8 @@ OPT ?= -O -g
 
 # Ceed directory
 CEED_DIR ?= ../..
-CEED_FLAGS ?= -I$(CEED_DIR)/include -std=c11 $(OPT)
-CEED_LIBS ?= -Wl,-rpath,$(abspath $(CEED_DIR)/lib) -L$(CEED_DIR)/lib -lceed -lm
+CEED_FLAGS ?= -I$(CEED_DIR)/include -std=c11  $(OPT)
+CEED_LIBS ?= -Wl,-rpath,$(abspath $(CEED_DIR)/lib) -L$(CEED_DIR)/lib -lceed -L$(CEED_DIR)/examples/ceed -L$(CEED_DIR)/examples/ceed/ex4-volume-rs/target/release -lex4_volume_rs -lm
 
 EXAMPLES.c = $(wildcard ex*.c)
 EXAMPLES = $(EXAMPLES.c:%.c=%)
@@ -24,6 +24,11 @@ all: $(EXAMPLES)
 # Remove built-in rules
 %: %.c
 
+# Special build rule for example 4 (rust)
+ex4-volume-rust: ex4-volume-rust.c
+	cargo +nightly build --release --manifest-path ex4-volume-rs/Cargo.toml
+	$(LINK.c) $(CEED_FLAGS) $(CEED_LDFLAGS) $< -o $@ $(CEED_LIBS)
+
 # Rules for building the examples
 %: %.c
 	$(LINK.c) $(CEED_FLAGS) $(CEED_LDFLAGS) $< -o $@ $(CEED_LIBS)

examples/ceed/ex4-volume-rs/.cargo/config.toml

@@ -0,0 +1,10 @@
+#[build]
+#target = "nvptx64-nvidia-cuda"  # default compilation target
+
+
+[target.nvptx64-nvidia-cuda]
+rustflags = [
+  "-C", "embed-bitcode=yes",
+  "-C", "target-cpu=sm_80",
+  "-Z", "dylib-lto",
+]

examples/ceed/ex4-volume-rs/Cargo.toml

@@ -0,0 +1,16 @@
+[package]
+name = "ex4-volume-rs"
+version = "0.1.0"
+edition = "2021"
+
+[profile.dev]
+panic = "abort"
+
+[profile.release]
+panic = "abort"
+
+# Compiles the crate as a lib (for GPU) and staticlib (for CPU)
+[lib]
+crate-type = ["lib", "staticlib"]
+
+[dependencies]

examples/ceed/ex4-volume-rs/rust-toolchain.toml

@@ -0,0 +1,2 @@
+[toolchain]
+channel = "nightly"