GitHub

Builds TinyCC Cli and Library For C Scripting in R

R-CMD-check CRAN status Rtinycc status badge

Abstract

Rtinycc is an R interface to TinyCC, providing both CLI access and a libtcc-backed in-memory compiler. It includes an FFI inspired by Bun’s FFI for binding C symbols with predictable type conversions and pointer utilities. The package works on unix-alikes and Windows and focuses on embedding TinyCC and enabling JIT-compiled bindings directly from R. Combined with treesitter.c, which provides C header parsers, it can be used to rapidly generate declarative bindings.

How it works

When you call tcc_compile(), Rtinycc generates C wrapper functions whose signature follows the .Call convention (SEXP in, SEXP out). These wrappers convert R types to C, call the target function, and convert the result back. TCC compiles them in-memory – no shared library is written to disk and no R_init_* registration is needed.

After tcc_relocate(), wrapper pointers are retrieved via tcc_get_symbol(), which internally calls RC_libtcc_get_symbol(). That function converts TCC’s raw void* into a DL_FUNC wrapped with R_MakeExternalPtrFn (tagged "native symbol"). The symbol pointer retains its owning tcc_state, so the relocated code remains alive while the symbol is reachable. On the R side, make_callable() creates a closure that passes this external pointer to .Call (aliased as .RtinyccCall to keep R CMD check happy). Non-memory states instead produce one-shot artifacts with tcc_output_file() and cannot expose callable symbol pointers.

The design follows CFFI’s API-mode pattern: instead of computing struct layouts and calling conventions in R (ABI-mode, like Python’s ctypes), the generated C code lets TCC handle sizeof, offsetof, and argument passing. Rtinycc never replicates platform-specific layout rules. The wrappers can also link against external shared libraries whose symbols TCC resolves at relocation time. For background on how this compares to a libffi approach, see the RSimpleFFI README.

On macOS the configure script strips -flat_namespace from TCC’s build to avoid BUS ERROR issues. Without it, TCC cannot resolve host symbols (e.g. RC_free_finalizer) through the dynamic linker. Rtinycc works around this with RC_libtcc_add_host_symbols(), which registers package-internal C functions via tcc_add_symbol() before relocation. Any new C function referenced by generated TCC code must be added there.

On Windows, the configure.win script generates a UCRT-backed msvcrt.def so TinyCC resolves CRT symbols against ucrtbase.dll (R 4.2+ uses UCRT).

Ownership semantics are explicit. Pointers from tcc_malloc() are tagged rtinycc_owned and can be released with tcc_free() (or by their R finalizer). Generated struct constructors use a struct-specific tag (struct_<name>) and host-allocated storage with RC_owned_native_finalizer; free them with struct_<name>_free(), not tcc_free(). Pointers from tcc_data_ptr() are tagged rtinycc_borrowed and are never freed by Rtinycc. Array returns are copied into a fresh R vector; set free = TRUE only when the C function returns a malloc-owned buffer.

Installation

install.packages(
      'Rtinycc',
        repos = c('https://sounkou-bioinfo.r-universe.dev',
                  'https://cloud.r-project.org')
        )

Usage

CLI

The CLI interface compiles C source files to standalone executables using the bundled TinyCC toolchain.

library(Rtinycc)
src <- system.file("c_examples", "forty_two.c", package = "Rtinycc")
exe <- tempfile()
tcc_run_cli(c(
  "-B", tcc_prefix(),
  paste0("-I", tcc_include_paths()),
  paste0("-L", tcc_lib_paths()),
  src, "-o", exe
))
#> [1] 0
Sys.chmod(exe, mode = "0755")
system2(exe, stdout = TRUE)
#> [1] "42"

For in-memory workflows, prefer libtcc instead.

In-memory compilation with libtcc

We can compile and call C functions entirely in memory. This is the simplest path for quick JIT compilation.

state <- tcc_state(output = "memory")
tcc_compile_string(state, "int forty_two(){ return 42; }")
#> [1] 0
tcc_relocate(state)
#> [1] 0
tcc_call_symbol(state, "forty_two", return = "int")
#> [1] 42

For low-level pointer-style calls, tcc_call_symbol() can also use an R .C()-like convention: call a void C routine with pointers to copied argument buffers, then return a list of the modified values.

state <- tcc_state(output = "memory")
tcc_compile_string(state, "void add_one(int *x) { x[0] += 1; }")
#> [1] 0
tcc_relocate(state)
#> [1] 0
tcc_call_symbol(state, "add_one", x = as.integer(41))
#> $x
#> [1] 42

The lower-level API gives full control over include paths, libraries, and the R C API. Using #define _Complex as a workaround for TCC’s lack of complex type support, we can link against R’s headers and call into libR.

state <- tcc_state(output = "memory")
tcc_add_include_path(state, R.home("include"))
#> [1] 0
tcc_add_library_path(state, R.home("lib"))
#> [1] 0
code <- '
#define _Complex
#include <R.h>
#include <Rinternals.h>

double call_r_sqrt(void) {
  SEXP fn   = PROTECT(Rf_findFun(Rf_install("sqrt"), R_BaseEnv));
  SEXP val  = PROTECT(Rf_ScalarReal(16.0));
  SEXP call = PROTECT(Rf_lang2(fn, val));
  SEXP out  = PROTECT(Rf_eval(call, R_GlobalEnv));
  double res = REAL(out)[0];
  UNPROTECT(4);
  return res;
}
'
tcc_compile_string(state, code)
#> [1] 0
tcc_relocate(state)
#> [1] 0
tcc_call_symbol(state, "call_r_sqrt", return = "double")
#> [1] 4

Pointer utilities

Rtinycc ships a set of typed memory access functions similar to what the ctypesio package offers, but designed around our FFI pointer model. Every scalar C type has a corresponding tcc_read_* / tcc_write_* pair that operates at a byte offset into any external pointer, so you can walk structs, arrays, and output parameters without writing C helpers.

ptr <- tcc_cstring("hello")
tcc_read_cstring(ptr)
#> [1] "hello"
tcc_read_bytes(ptr, 5)
#> [1] 68 65 6c 6c 6f
tcc_ptr_addr(ptr, hex = TRUE)
#> [1] "0x6531efbdcb10"
tcc_ptr_is_null(ptr)
#> [1] FALSE
tcc_free(ptr)
#> NULL

Typed reads and writes cover the full scalar range (i8/u8, i16/u16, i32/u32, i64/u64, f32/f64) plus pointer dereferencing via tcc_read_ptr / tcc_write_ptr. All operations use a byte offset and memcpy internally for alignment safety.

buf <- tcc_malloc(32)
tcc_write_i32(buf, 0L, 42L)
tcc_write_f64(buf, 8L, pi)
tcc_read_i32(buf, offset = 0L)
#> [1] 42
tcc_read_f64(buf, offset = 8L)
#> [1] 3.141593
tcc_free(buf)
#> NULL

Pointer-to-pointer workflows are supported for C APIs that return values through output parameters.

ptr_ref <- tcc_malloc(.Machine$sizeof.pointer %||% 8L)
target <- tcc_malloc(8)
tcc_ptr_set(ptr_ref, target)
#> <pointer: 0x6531ef32bcb0>
tcc_data_ptr(ptr_ref)
#> <pointer: 0x6531efb6ac30>
tcc_ptr_set(ptr_ref, tcc_null_ptr())
#> <pointer: 0x6531ef32bcb0>
tcc_free(target)
#> NULL
tcc_free(ptr_ref)
#> NULL

Declarative FFI

A declarative interface inspired by Bun’s FFI sits on top of the lower-level API. We define types explicitly and Rtinycc generates the binding code, compiling it in memory with TCC.

Type system

The FFI exposes a small set of type mappings between R and C. Conversions are explicit and predictable so callers know when data is shared versus copied.

The scalar type names are C-facing, but the R-side carriers are not all one-to-one with those C widths:

  • i8, i16, i32, u8, and u16 are mediated through R integer scalars
  • u32, i64, u64, f32, and f64 are mediated through R numeric (double) coercion and boxing
  • bool uses R logical
  • cstring uses an R character scalar

This means u32 is routed through double to preserve the full unsigned 32-bit range, and i64 / u64 are only exact up to 2^53 on the R side.

Array arguments pass R vectors to C with zero copy: raw maps to uint8_t*, integer_array to int32_t*, numeric_array to double*.

Pointer types include ptr (opaque external pointer), sexp (pass a SEXP directly), and callback signatures like callback:double(double).

Variadic functions are supported in two forms: typed prefix tails (varargs) and bounded dynamic tails (varargs_types + varargs_min/varargs_max). Prefix mode is the cheaper runtime path because dispatch is by tail arity only; bounded dynamic mode adds per-call scalar type inference to select a compatible wrapper. For hot loops, prefer fixed arity first, then prefix variadics with a tight maximum tail size.

Array returns use returns = list(type = "integer_array", length_arg = 2, free = TRUE) to copy the result into a new R vector. The length_arg is the 1-based index of the C argument that carries the array length. Set free = TRUE when the C function returns a malloc-owned buffer.

Simple functions

ffi <- tcc_ffi() |>
  tcc_source("
    int add(int a, int b) { return a + b; }
  ") |>
  tcc_bind(add = list(args = list("i32", "i32"), returns = "i32")) |>
  tcc_compile()
ffi$add(5L, 3L)
#> [1] 8

Variadic calls (e.g. Rprintf style)

Rtinycc supports two ways to bind variadic tails. The legacy approach uses varargs as a typed prefix tail, while the bounded dynamic approach uses varargs_types together with varargs_min and varargs_max. In the bounded mode, wrappers are generated across the allowed arity and type combinations, and runtime dispatch selects the matching wrapper from the scalar tail values provided at call time.

ffi_var <- tcc_ffi() |>
  tcc_header("#include <R_ext/Print.h>") |>
  tcc_source('
    #include <stdarg.h>

    int sum_fmt(int n, ...) {
      va_list ap;
      va_start(ap, n);
      int s = 0;
      for (int i = 0; i < n; i++) s += va_arg(ap, int);
      va_end(ap);
      Rprintf("sum_fmt(%d) = %d\\n", n, s);
      return s;
    }
  ') |>
  tcc_bind(
    Rprintf = list(
      args = list("cstring"),
      variadic = TRUE,
      varargs_types = list("i32"),
      varargs_min = 0L,
      varargs_max = 4L,
      returns = "void"
    ),
    sum_fmt = list(
      args = list("i32"),
      variadic = TRUE,
      varargs_types = list("i32"),
      varargs_min = 0L,
      varargs_max = 4L,
      returns = "i32"
    )
  ) |>
  tcc_compile()
ffi_var$Rprintf("Rprintf via bind: %d + %d = %d\n", 2L, 3L, 5L)
#> Rprintf via bind: 2 + 3 = 5
#> NULL
ffi_var$sum_fmt(0L)
#> sum_fmt(0) = 0
#> [1] 0
ffi_var$sum_fmt(2L, 10L, 20L)
#> sum_fmt(2) = 30
#> [1] 30
ffi_var$sum_fmt(4L, 1L, 2L, 3L, 4L)
#> sum_fmt(4) = 10
#> [1] 10

Linking external libraries

We can bind directly to symbols in shared libraries. Here we link against libm.

math <- tcc_ffi() |>
  tcc_library("m") |>
  tcc_bind(
    sqrt  = list(args = list("f64"), returns = "f64"),
    sin   = list(args = list("f64"), returns = "f64"),
    floor = list(args = list("f64"), returns = "f64")
  ) |>
  tcc_compile()
math$sqrt(16.0)
#> [1] 4
math$sin(pi / 2)
#> [1] 1
math$floor(3.7)
#> [1] 3

CUDA through NVRTC on a remote GPU

TinyCC compiles the host adapter, not the CUDA kernel. The adapter links to the CUDA Driver API and NVRTC; NVRTC compiles CUDA C to PTX, and the driver loads and launches that PTX. A persistent mirai daemon on the GPU rig holds the Rtinycc state, compiled host adapter, and CUDA configuration across the steps below.

Set RTINYCC_RUN_SSH_DEMOS=true while rendering to launch the rig session and execute the example. Ordinary CRAN and CI renders show the code without opening an SSH connection.

First launch one persistent daemon through mirai’s SSH tunnel, load Rtinycc in that remote R process, and inspect the remote host.

library(mirai)
daemons(
  n = 1L,
  url = local_url(tcp = TRUE),
  remote = ssh_config(
    "ssh://sounkoutoure@localhost:2222",
    tunnel = TRUE
  ),
  .compute = "gpu"
)
everywhere(library(Rtinycc), .compute = "gpu")
gpu_session <- mirai(
  list(
    host = Sys.info()[["nodename"]],
    pid = Sys.getpid(),
    rtinycc = as.character(packageVersion("Rtinycc")),
    device = trimws(system2(
      "/usr/lib/wsl/lib/nvidia-smi",
      c("--query-gpu=name", "--format=csv,noheader"),
      stdout = TRUE
    )[[1L]])
  ),
  .compute = "gpu"
)
as.data.frame(gpu_session[], check.names = FALSE)

The long step is collapsed. It runs inside the existing daemon with everywhere(): detect the CUDA architecture, compile the host adapter with Rtinycc, link the CUDA Driver and NVRTC libraries, then retain the compiled FFI object in the daemon’s global environment.

Show the remote Rtinycc, NVRTC, and CUDA compilation step 
cuda_setup <- everywhere(
  {
    # This opt-in demonstration is run on an NVIDIA Linux/WSL host. Override the
    # paths and architecture when CUDA is installed elsewhere.
    cuda_path <- Sys.getenv(
      "RTINYCC_CUDA_DRIVER",
      "/usr/lib/wsl/lib/libcuda.so.1"
    )
    nvrtc_path <- Sys.getenv(
      "RTINYCC_NVRTC",
      file.path(Sys.getenv("CUDA_HOME", "/usr/local/cuda"), "lib64", "libnvrtc.so")
    )
    stopifnot(file.exists(cuda_path), file.exists(nvrtc_path))
    cuda_arch <- Sys.getenv("RTINYCC_CUDA_ARCH")
    if (!nzchar(cuda_arch)) {
      nvidia_smi <- Sys.getenv(
        "RTINYCC_NVIDIA_SMI",
        "/usr/lib/wsl/lib/nvidia-smi"
      )
      compute_capability <- trimws(system2(
        nvidia_smi,
        c("--query-gpu=compute_cap", "--format=csv,noheader"),
        stdout = TRUE
      )[[1L]])
      cuda_arch <- paste0("compute_", gsub("[^0-9]", "", compute_capability))
    }
    stopifnot(grepl("^compute_[0-9]+$", cuda_arch))
    code <- '
    #include <stdlib.h>
    #include <stdio.h>
    #include <string.h>

    typedef int CUresult;
    typedef int CUdevice;
    typedef unsigned long long CUdeviceptr;
    typedef struct CUctx_st *CUcontext;
    typedef struct CUmod_st *CUmodule;
    typedef struct CUfunc_st *CUfunction;
    typedef struct CUstream_st *CUstream;

    typedef int nvrtcResult;
    typedef struct _nvrtcProgram *nvrtcProgram;

    extern CUresult cuInit(unsigned int flags);
    extern CUresult cuDeviceGet(CUdevice *device, int ordinal);
    extern CUresult cuCtxCreate_v2(CUcontext *ctx, unsigned int flags,
                                   CUdevice device);
    extern CUresult cuCtxDestroy_v2(CUcontext ctx);
    extern CUresult cuMemAlloc_v2(CUdeviceptr *ptr, size_t bytes);
    extern CUresult cuMemFree_v2(CUdeviceptr ptr);
    extern CUresult cuMemcpyHtoD_v2(CUdeviceptr dst, const void *src,
                                    size_t bytes);
    extern CUresult cuMemcpyDtoH_v2(void *dst, CUdeviceptr src, size_t bytes);
    extern CUresult cuModuleLoadData(CUmodule *module, const void *image);
    extern CUresult cuModuleUnload(CUmodule module);
    extern CUresult cuModuleGetFunction(CUfunction *function, CUmodule module,
                                        const char *name);
    extern CUresult cuLaunchKernel(CUfunction function,
                                   unsigned int grid_x,
                                   unsigned int grid_y,
                                   unsigned int grid_z,
                                   unsigned int block_x,
                                   unsigned int block_y,
                                   unsigned int block_z,
                                   unsigned int shared_mem_bytes,
                                   CUstream stream,
                                   void **kernel_params,
                                   void **extra);
    extern CUresult cuCtxSynchronize(void);

    extern nvrtcResult nvrtcCreateProgram(nvrtcProgram *program,
                                          const char *source,
                                          const char *name,
                                          int n_headers,
                                          const char * const *headers,
                                          const char * const *include_names);
    extern nvrtcResult nvrtcCompileProgram(nvrtcProgram program,
                                           int n_options,
                                           const char * const *options);
    extern nvrtcResult nvrtcGetPTXSize(nvrtcProgram program, size_t *size);
    extern nvrtcResult nvrtcGetPTX(nvrtcProgram program, char *ptx);
    extern nvrtcResult nvrtcGetProgramLogSize(nvrtcProgram program,
                                              size_t *size);
    extern nvrtcResult nvrtcGetProgramLog(nvrtcProgram program, char *log);
    extern nvrtcResult nvrtcDestroyProgram(nvrtcProgram *program);

    static char rt_cuda_error[8192];

    static void rt_cuda_set_error(const char *stage, int code) {
      snprintf(rt_cuda_error, sizeof(rt_cuda_error), "%s failed: %d", stage, code);
    }

    char *rt_cuda_last_error(void) {
      return rt_cuda_error;
    }

    static void rt_cuda_set_nvrtc_log(nvrtcProgram program, int code) {
      char *log;
      size_t copy_size;
      size_t size = 0;

      snprintf(rt_cuda_error, sizeof(rt_cuda_error),
               "nvrtcCompileProgram failed: %d", code);
      if (nvrtcGetProgramLogSize(program, &size) != 0 || size <= 1)
        return;
      log = malloc(size);
      if (!log)
        return;
      if (nvrtcGetProgramLog(program, log) != 0) {
        free(log);
        rt_cuda_set_error("nvrtcGetProgramLog", code);
        return;
      }
      copy_size = size - 1;
      if (copy_size >= sizeof(rt_cuda_error))
        copy_size = sizeof(rt_cuda_error) - 1;
      memcpy(rt_cuda_error, log, copy_size);
      rt_cuda_error[copy_size] = 0;
      free(log);
    }

    int rt_cuda_vec_add(double *x, double *y, double *out, int n) {
      static const char *kernel_source =
        "extern \\"C\\" __global__ void vec_add(const double *x, "
        "const double *y, double *out, int n) {"
        "  int i = (int)(blockIdx.x * blockDim.x + threadIdx.x);"
        "  if (i < n) out[i] = x[i] + y[i];"
        "}";
      const char *options[] = {
        "--gpu-architecture=@CUDA_ARCH@",
        "--std=c++11"
      };
      nvrtcProgram program = 0;
      CUcontext context = 0;
      CUmodule module = 0;
      CUfunction function = 0;
      CUdevice device = 0;
      CUdeviceptr device_x = 0;
      CUdeviceptr device_y = 0;
      CUdeviceptr device_out = 0;
      char *ptx = 0;
      size_t ptx_size = 0;
      size_t bytes;
      void *params[4];
      unsigned int blocks;
      int rc;
      int status = -1;

      rt_cuda_error[0] = 0;
      if (n < 0) {
        strcpy(rt_cuda_error, "negative vector length");
        return -1;
      }
      if (n == 0)
        return 0;
      bytes = (size_t)n * sizeof(double);

      rc = nvrtcCreateProgram(&program, kernel_source, "vec_add.cu",
                              0, 0, 0);
      if (rc != 0) {
        rt_cuda_set_error("nvrtcCreateProgram", rc);
        goto cleanup;
      }
      rc = nvrtcCompileProgram(program, 2, options);
      if (rc != 0) {
        rt_cuda_set_nvrtc_log(program, rc);
        goto cleanup;
      }
      rc = nvrtcGetPTXSize(program, &ptx_size);
      if (rc != 0) {
        rt_cuda_set_error("nvrtcGetPTXSize", rc);
        goto cleanup;
      }
      ptx = malloc(ptx_size);
      if (!ptx) {
        strcpy(rt_cuda_error, "PTX allocation failed");
        goto cleanup;
      }
      rc = nvrtcGetPTX(program, ptx);
      if (rc != 0) {
        rt_cuda_set_error("nvrtcGetPTX", rc);
        goto cleanup;
      }

      rc = cuInit(0);
      if (rc != 0) {
        rt_cuda_set_error("cuInit", rc);
        goto cleanup;
      }
      rc = cuDeviceGet(&device, 0);
      if (rc != 0) {
        rt_cuda_set_error("cuDeviceGet", rc);
        goto cleanup;
      }
      rc = cuCtxCreate_v2(&context, 0, device);
      if (rc != 0) {
        rt_cuda_set_error("cuCtxCreate_v2", rc);
        goto cleanup;
      }
      rc = cuModuleLoadData(&module, ptx);
      if (rc != 0) {
        rt_cuda_set_error("cuModuleLoadData", rc);
        goto cleanup;
      }
      rc = cuModuleGetFunction(&function, module, "vec_add");
      if (rc != 0) {
        rt_cuda_set_error("cuModuleGetFunction", rc);
        goto cleanup;
      }
      rc = cuMemAlloc_v2(&device_x, bytes);
      if (rc != 0) {
        rt_cuda_set_error("cuMemAlloc(x)", rc);
        goto cleanup;
      }
      rc = cuMemAlloc_v2(&device_y, bytes);
      if (rc != 0) {
        rt_cuda_set_error("cuMemAlloc(y)", rc);
        goto cleanup;
      }
      rc = cuMemAlloc_v2(&device_out, bytes);
      if (rc != 0) {
        rt_cuda_set_error("cuMemAlloc(out)", rc);
        goto cleanup;
      }
      rc = cuMemcpyHtoD_v2(device_x, x, bytes);
      if (rc != 0) {
        rt_cuda_set_error("cuMemcpyHtoD(x)", rc);
        goto cleanup;
      }
      rc = cuMemcpyHtoD_v2(device_y, y, bytes);
      if (rc != 0) {
        rt_cuda_set_error("cuMemcpyHtoD(y)", rc);
        goto cleanup;
      }

      params[0] = &device_x;
      params[1] = &device_y;
      params[2] = &device_out;
      params[3] = &n;
      blocks = ((unsigned int)n + 255U) / 256U;
      rc = cuLaunchKernel(function, blocks, 1, 1, 256, 1, 1,
                          0, 0, params, 0);
      if (rc != 0) {
        rt_cuda_set_error("cuLaunchKernel", rc);
        goto cleanup;
      }
      rc = cuCtxSynchronize();
      if (rc != 0) {
        rt_cuda_set_error("cuCtxSynchronize", rc);
        goto cleanup;
      }
      rc = cuMemcpyDtoH_v2(out, device_out, bytes);
      if (rc != 0) {
        rt_cuda_set_error("cuMemcpyDtoH", rc);
        goto cleanup;
      }
      status = 0;

    cleanup:
      if (device_out)
        cuMemFree_v2(device_out);
      if (device_y)
        cuMemFree_v2(device_y);
      if (device_x)
        cuMemFree_v2(device_x);
      if (module)
        cuModuleUnload(module);
      if (context)
        cuCtxDestroy_v2(context);
      free(ptx);
      if (program)
        nvrtcDestroyProgram(&program);
      return status;
    }
    '
    code <- sub("@CUDA_ARCH@", cuda_arch, code, fixed = TRUE)
    cuda_compiled <- tcc_ffi() |>
      tcc_source(code) |>
      tcc_library(cuda_path) |>
      tcc_library(nvrtc_path) |>
      tcc_bind(
        rt_cuda_vec_add = list(
          args = list("numeric_array", "numeric_array", "numeric_array", "i32"),
          returns = "i32"
        ),
        rt_cuda_last_error = list(args = list(), returns = "cstring")
      ) |>
      tcc_compile()
    assign("cuda", cuda_compiled, envir = globalenv())
    assign("cuda_arch", cuda_arch, envir = globalenv())
    invisible(TRUE)
  },
  .compute = "gpu"
)
invisible(cuda_setup)

A second mirai call sees the state retained by the same daemon.

cuda_ready <- mirai(
  list(
    host = Sys.info()[["nodename"]],
    architecture = cuda_arch,
    ffi_class = class(cuda)[[1L]],
    bindings = paste(c("rt_cuda_vec_add", "rt_cuda_last_error"), collapse = ", ")
  ),
  .compute = "gpu"
)
as.data.frame(cuda_ready[], check.names = FALSE)

Now submit the actual vector addition to that daemon. This call uses the already-retained Rtinycc FFI object; it does not resend or parse an R script.

cuda_job <- mirai(
  {
    n <- as.integer(Sys.getenv("RTINYCC_CUDA_N", "1000000"))
    stopifnot(length(n) == 1L, !is.na(n), n > 0L)
    set.seed(1)
    x <- runif(n)
    y <- runif(n)
    out <- numeric(n)
    timing <- system.time({
      rc <- cuda$rt_cuda_vec_add(x, y, out, n)
    })
    if (rc != 0L) {
      stop("CUDA vector add failed: ", cuda$rt_cuda_last_error())
    }
    expected <- x + y
    max_error <- max(abs(out - expected))
    result <- list(
      host = Sys.info()[["nodename"]],
      rtinycc_version = as.character(packageVersion("Rtinycc")),
      device = trimws(system2(
        Sys.getenv("RTINYCC_NVIDIA_SMI", "/usr/lib/wsl/lib/nvidia-smi"),
        c("--query-gpu=name", "--format=csv,noheader"),
        stdout = TRUE
      )[[1L]]),
      architecture = cuda_arch,
      n = n,
      status = rc,
      elapsed_seconds_including_nvrtc = unname(timing[["elapsed"]]),
      max_error = max_error,
      first_values = head(out)
    )
    stopifnot(max_error == 0)
    invisible(result)
  },
  .compute = "gpu"
)
cuda_result <- cuda_job[]
stopifnot(
  is.list(cuda_result),
  identical(cuda_result$status, 0L),
  identical(cuda_result$max_error, 0)
)
data.frame(
  host = cuda_result$host,
  rtinycc = cuda_result$rtinycc_version,
  device = cuda_result$device,
  architecture = cuda_result$architecture,
  n = cuda_result$n,
  elapsed_seconds = cuda_result$elapsed_seconds_including_nvrtc,
  max_error = cuda_result$max_error,
  check.names = FALSE
)

Finally release the retained TCC state and stop the remote daemon.

everywhere(
  {
    rm(list = intersect(c("cuda", "cuda_arch"), ls(globalenv())),
       envir = globalenv())
    gc()
  },
  .compute = "gpu"
)
daemons(0, .compute = "gpu")

The displayed result is returned by the remote R process. Timing includes NVRTC compilation, CUDA context creation, device allocation, transfers, launch, and synchronization; this is a correctness demonstration rather than a steady-state kernel benchmark.

Compiler options

Use tcc_options() to pass raw TinyCC options in the high-level FFI pipeline. For low-level states, use tcc_set_options() directly.

ffi_opt_off <- tcc_ffi() |>
  tcc_options("-O0") |>
  tcc_source('
    int opt_macro() {
    #ifdef __OPTIMIZE__
      return 1;
    #else
      return 0;
    #endif
    }
  ') |>
  tcc_bind(opt_macro = list(args = list(), returns = "i32")) |>
  tcc_compile()
ffi_opt_on <- tcc_ffi() |>
  tcc_options(c("-Wall", "-O2")) |>
  tcc_source('
    int opt_macro() {
    #ifdef __OPTIMIZE__
      return 1;
    #else
      return 0;
    #endif
    }
  ') |>
  tcc_bind(opt_macro = list(args = list(), returns = "i32")) |>
  tcc_compile()
ffi_opt_off$opt_macro()
#> [1] 0
ffi_opt_on$opt_macro()
#> [1] 1

Working with arrays

R vectors are passed to C with zero copy. Mutations in C are visible in R.

ffi <- tcc_ffi() |>
  tcc_source("
    #include <stdlib.h>
    #include <string.h>

    int64_t sum_array(int32_t* arr, int32_t n) {
      int64_t s = 0;
      for (int i = 0; i < n; i++) s += arr[i];
      return s;
    }

    void bump_first(int32_t* arr) { arr[0] += 10; }

    int32_t* dup_array(int32_t* arr, int32_t n) {
      int32_t* out = malloc(sizeof(int32_t) * n);
      memcpy(out, arr, sizeof(int32_t) * n);
      return out;
    }
  ") |>
  tcc_bind(
    sum_array  = list(args = list("integer_array", "i32"), returns = "i64"),
    bump_first = list(args = list("integer_array"), returns = "void"),
    dup_array  = list(
      args = list("integer_array", "i32"),
      returns = list(type = "integer_array", length_arg = 2, free = TRUE)
    )
  ) |>
  tcc_compile()
x <- as.integer(1:100) # to avoid ALTREP
.Internal(inspect(x))
#> @6531f0dbe828 13 INTSXP g0c0 [REF(65535)]  1 : 100 (compact)
ffi$sum_array(x, length(x))
#> [1] 5050
# Zero-copy: C mutation reflects in R
ffi$bump_first(x)
#> NULL
x[1]
#> [1] 11
# Array return: copied into a new R vector, C buffer freed
y <- ffi$dup_array(x, length(x))
y[1]
#> [1] 11
.Internal(inspect(x))
#> @6531f0dbe828 13 INTSXP g0c0 [REF(65535)]  11 : 110 (expanded)

Advanced FFI features

Structs and unions

Complex C types are supported declaratively. Use tcc_struct() to generate allocation and accessor helpers. Free instances when done.

"ffi <- tcc_ffi() |> tcc_source(' #include

Read the original on github.com ↗