Initial nvrtcc implementation

nvrtcc/CMakeLists.txt

@@ -0,0 +1,56 @@
+##===----------------------------------------------------------------------===##
+##
+## Part of nvrtcc in CUDA Core Compute Libraries,
+## 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
+## SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES.
+##
+##===----------------------------------------------------------------------===##
+
+cmake_minimum_required(VERSION 3.21)
+
+if (NOT CMAKE_BUILD_TYPE)
+  set(CMAKE_BUILD_TYPE Release)
+endif()
+
+project(nvrtcc LANGUAGES CXX)
+
+# Set output directory.
+set(nvrtcc_exec_output_dir ${nvrtcc_BINARY_DIR}/bin)
+
+find_package(CUDAToolkit REQUIRED)
+
+# Find libcufilt.
+find_library(
+  nvrtcc_cufilt_lib
+  "cufilt"
+  PATHS "${CUDAToolkit_LIBRARY_DIR}"
+  NO_DEFAULT_PATH
+)
+if (NOT nvrtcc_cufilt_lib)
+  message(FATAL_ERROR "nvrtcc: cu++filt library (libcufilt.a) not found.")
+endif()
+
+# Add nvrtcc_device_compiler target.
+add_executable(nvrtcc_device_compiler src/nvrtcc_device_compiler.cpp)
+target_compile_features(nvrtcc_device_compiler PRIVATE cxx_std_17)
+target_include_directories(
+  nvrtcc_device_compiler
+  PRIVATE "${CUDAToolkit_INCLUDE_DIRS}"
+)
+target_link_libraries(nvrtcc_device_compiler PRIVATE ${nvrtcc_cufilt_lib})
+set_target_properties(
+  nvrtcc_device_compiler
+  PROPERTIES RUNTIME_OUTPUT_DIRECTORY "${nvrtcc_exec_output_dir}"
+)
+
+# Add nvrtcc target that copies nvrtcc_cicc to the output directory and depends on nvrtcc_device_compiler.
+add_custom_target(
+  nvrtcc
+  ALL
+  COMMAND
+    ${CMAKE_COMMAND} -E copy_directory ${CMAKE_SOURCE_DIR}/bin
+    ${nvrtcc_exec_output_dir}
+)
+add_dependencies(nvrtcc nvrtcc_device_compiler)

nvrtcc/README.md

@@ -0,0 +1,55 @@
+> [!CAUTION]
+> This is an internal tool not intended for public use.
+
+> [!WARNING]
+> This tool is experimental.
+
+# NVRTCC
+
+## Overview
+
+`nvrtcc` is a tool to simplify [NVRTC](https://docs.nvidia.com/cuda/nvrtc/index.html) testing. It follows the `nvcc` compilation trajectory and replaces the `nvcc` generated PTX code with the NVRTC compiled one. The main advantage is that you can keep the source files almost the same as if compiled with `nvcc`. `nvrtcc` makes sure all of the necessary symbols are present in the generated PTX and let's `nvcc` do the host compilation and linking.
+
+The compilation with NVRTC is optional and can be controlled by the `-use-nvrtc` flag. This allows `nvrtcc` to be used as the `CMAKE_CUDA_COMPILER` where we can trick CMake to think it's using `nvcc`. `nvrtcc` supports almost all `nvcc` options except for those that are unsupported by NVRTC.
+
+## Example
+
+When compiling with `nvrtcc`, we need to make sure NVRTC wouldn't see the host includes and symbols, so they must be guarded by `#ifndef __CUDACC_RTC__` preprocessor directive.
+```cpp
+// hello_world.cu
+
+#ifndef __CUDACC_RTC__
+#  include <cstdio>
+#endif
+
+__global__ void kernel()
+{
+#ifdef __CUDACC_RTC__
+    printf("Hello world from NVRTC!\n");
+#else
+    printf("Hello world from NVCC!\n");
+#endif
+}
+
+#ifndef __CUDACC_RTC__
+int main()
+{
+    kernel<<<1, 1>>>();
+    cudaDeviceSynchronize();
+}
+#endif
+```
+
+If we compile with `nvcc` the output should look as:
+```sh
+$ nvcc hello_world.cu -o hello_world
+$ ./hello_world
+Hello world from NVCC!
+```
+
+On the other hand if compiled with `nvrtcc` with `-use-nvrtc` flag present, the output should be:
+```sh
+$ nvrtcc hello_world.cu -o hello_world -use-nvrtc
+$ ./hello_world
+Hello world from NVRTCC!
+```