Add support for nvcc & hipcc (cuda/rocm) (#23805)

I've been working on making some basic cuda examples work, both with
cuda (nvcc) and with AMD HIP (hipcc) https://github.com/monofuel/hippo

- hipcc is just a drop-in replacement for clang and works out of the box
with clang settings in Nim. hipcc is capable of compiling for AMD ROCm
or to CUDA, depending on how HIP_PLATFORM is set.
- nvcc is a little quirky. we can use `-x cu` to tell it to handle nim's
`.cpp` files as if they were `.cu` files. nvcc expects all backend
compiler flags to be wrapped with a special `-Xcompiler=""` flag when
compiling and also when linking.

I manually tested on a linux desktop with amd and a laptop with nvidia.

3 files changed, 42 insertions, 2 deletions

diff --git a/compiler/extccomp.nim b/compiler/extccomp.nim
index 0ea458e3b..dc9b43514 100644
--- a/compiler/extccomp.nim
+++ b/compiler/extccomp.nim
@@ -172,6 +172,22 @@ compiler vcc:
     cppXsupport: "",
     props: {hasCpp, hasAssume, hasDeclspec})
 
+# Nvidia CUDA NVCC Compiler
+compiler nvcc:
+  result = gcc()
+  result.name = "nvcc"
+  result.compilerExe = "nvcc"
+  result.cppCompiler = "nvcc"
+  result.compileTmpl = "-c -x cu -Xcompiler=\"$options\" $include -o $objfile $file"
+  result.linkTmpl = "$buildgui $builddll -o $exefile $objfiles -Xcompiler=\"$options\""
+
+# AMD HIPCC Compiler (rocm/cuda)
+compiler hipcc:
+  result = clang()
+  result.name = "hipcc"
+  result.compilerExe = "hipcc"
+  result.cppCompiler = "hipcc"
+
 compiler clangcl:
   result = vcc()
   result.name = "clang_cl"
@@ -285,7 +301,9 @@ const
     envcc(),
     icl(),
     icc(),
-    clangcl()]
+    clangcl(),
+    hipcc(),
+    nvcc()]
 
   hExt* = ".h"
 
diff --git a/compiler/options.nim b/compiler/options.nim
index a6a2b084e..bdea506d2 100644
--- a/compiler/options.nim
+++ b/compiler/options.nim
@@ -254,7 +254,7 @@ type
 
   TSystemCC* = enum
     ccNone, ccGcc, ccNintendoSwitch, ccLLVM_Gcc, ccCLang, ccBcc, ccVcc,
-    ccTcc, ccEnv, ccIcl, ccIcc, ccClangCl
+    ccTcc, ccEnv, ccIcl, ccIcc, ccClangCl, ccHipcc, ccNvcc
 
   ExceptionSystem* = enum
     excNone,   # no exception system selected yet
diff --git a/doc/nimc.md b/doc/nimc.md
index 25acf31e8..38558454b 100644
--- a/doc/nimc.md
+++ b/doc/nimc.md
@@ -481,6 +481,28 @@ They are:
 5. nl_types. No headers for this.
 6. As mmap is not supported, the nimAllocPagesViaMalloc option has to be used.
 
+GPU Compilation
+===============
+
+Compiling for GPU computation can be achieved with `--cc:nvcc` for CUDA with nvcc, or with `--cc:hipcc` for AMD GPUs with HIP. Both compilers require building for C++ with `nim cpp`.
+
+Here's a very simple CUDA kernel example using emit, which can be compiled with `nim cpp --cc:nvcc --define:"useMalloc" hello_kernel.nim` assuming you have the CUDA toolkit installed.
+
+```nim
+{.emit: """
+__global__ void add(int a, int b) {
+  int c;
+  c = a + b;
+}
+""".}
+
+proc main() =
+  {.emit: """
+  add<<<1,1>>>(2,7);
+  """.}
+
+main()
+```
 
 DLL generation
 ==============