update examples + readme for nim 2.1.9

.gitignore

@@ -10,3 +10,5 @@ nimcache
 *.ilk
 *.out
 .*
+
+!.vscode

.vscode/extensions.json

@@ -0,0 +1,5 @@
+{
+  "recommendations": [
+    "nimlang.nimlang"
+  ]
+}

.vscode/launch.json

@@ -0,0 +1,14 @@
+{
+  "version": "0.2.0",
+  "configurations": [
+      {
+          "type": "lldb",
+          "request": "launch",
+          "name": "nim: debug current file",
+          "preLaunchTask": "nim: build current file (for debugging)",
+          "program": "${workspaceFolder}/bin/${fileBasenameNoExtension}",
+          "args": [],
+          "cwd": "${workspaceFolder}",
+      }
+  ]
+}

.vscode/settings.json

@@ -0,0 +1,5 @@
+{
+  "files.associations": {
+    "version": "cpp"
+  }
+}

.vscode/tasks.json

@@ -0,0 +1,21 @@
+{
+  "version": "2.0.0",
+  "tasks": [
+      {
+          "label": "nim: build current file (for debugging)",
+          "command": "nim",
+          "args": [
+              "cpp",
+              "-g",
+              "--define:HippoRuntime=HIP_CPU",
+              "--debugger:native",
+              "-o:${workspaceRoot}/bin/${fileBasenameNoExtension}",
+              "${relativeFile}"
+          ],
+          "options": {
+              "cwd": "${workspaceRoot}"
+          },
+          "type": "shell",
+      }
+  ]
+}

README.md

@@ -5,6 +5,9 @@ A Julia Set fractal generated with HIP
 
 ## Minimal HIP Example
 
+- IMPORTANT: requires at least nim 2.1.9 for gpu usage
+- `--cc:nvcc` and `--cc:hipcc`
+
 ```
 import hippo
 
@@ -74,15 +77,13 @@ handleError(hipFree(dev_c))
     - You also have to set the `HIP_PLATFORM=nvidia` environment variable to build for nvidia GPUs if you don't have an nvidia GPU in your system
     - hipcc will pick nvcc by default if you have nvcc but do not have the amd rocm stack installed
 
+- If you want to build for CPU with `-d:"HippoRuntime:HIP_CPU"`, you must have intel tbb libraries installed. `yum install -y tbb-devel`
 
 ### Required flags for HIP
 
 - You must compile with nim cpp for c++
 - Please refer to the [examples](examples/) for the nim compiler settings needed for each target
 
-- The HIP Nvidia and the CUDA targets currently require some changes to the nim compiler
-  - I have an experimental branch here [Nim](https://github.com/monofuel/Nim/tree/hipcc-nvcc)
-
 ### Optional flags
 
 - `-d:HippoRuntime=HIP` (default) (requires hipcc)

examples/vector_sum.nims

@@ -1,7 +1,2 @@
 --path:"../src"
-
---cc:clang
---clang.cpp.exe:hipcc
---clang.cpp.linkerexe:hipcc
---passC: "--offload-arch=gfx1100"
---passC: "--offload-arch=gfx1102"
+--cc:hipcc

examples/vector_sum_hippo.nims

@@ -1,7 +1,2 @@
 --path:"../src"
-
---cc:clang
---clang.cpp.exe:hipcc
---clang.cpp.linkerexe:hipcc
---passC: "--offload-arch=gfx1100"
---passC: "--offload-arch=gfx1102"
+--cc:hipcc

hippo.nimble

@@ -1,11 +1,11 @@
-version     = "0.0.1"
+version     = "0.0.2"
 author      = "Andrew Brower"
 description = "HIP library for Nim"
 license     = "MIT"
 
 srcDir = "src"
 
-requires "nim >= 2.0.0"
+requires "nim >= 2.1.9"
 
 
 task test, "Run all tests":

tests/cuda/hello_kernel.nim

@@ -1,32 +1,12 @@
 # nim cpp --cc:nvcc --define:"useMalloc" hello_kernel.nim
+# minimal cuda example for nim without using hippo
 {.emit: """
-__global__ void add(int a, int b, int *c) {
-    *c = a + b;
+__global__ void add(int a, int b) {
+  int c;
+  c = a + b;
 }
 """.}
 
-type
-  cudaMemcpyKind* {.size: sizeof(cint), importcpp: "cudaMemcpyKind".} = enum
-    cudaMemcpyHostToHost = 0,    ## < Host-to-Host Copy
-    cudaMemcpyHostToDevice = 1,  ## < Host-to-Device Copy
-    cudaMemcpyDeviceToHost = 2,  ## < Device-to-Host Copy
-    cudaMemcpyDeviceToDevice = 3, ## < Device-to-Device Copy
-    cudaMemcpyDefault = 4        ## < Runtime will automatically determine copy-kind based on virtual addresses.
-
-proc cudaMalloc*(`ptr`: ptr pointer; size: cint): cint {.importcpp: "cudaMalloc(@)".}
-proc cudaMemcpy*(dst: pointer; src: pointer; size: cint; kind: cudaMemcpyKind): cint {.importcpp: "cudaMemcpy(@)".}
-proc cudaFree*(`ptr`: pointer): cint {.importcpp: "cudaFree(@)".}
-
-proc main() =
-  var c: int32
-  var dev_c: ptr[int32]
-  discard cudaMalloc(cast[ptr pointer](addr dev_c), sizeof(int32).cint)
-  {.emit: """
-  add<<<1,1>>>(2,7,dev_c);
-  """.}
-  discard cudaMemcpy(addr c, dev_c, sizeof(int32).cint, cudaMemcpyDeviceToHost)
-  echo "2 + 7 = ", c
-  discard cudaFree(dev_c)
-
-when isMainModule:
-  main()
+{.emit: """
+add<<<1,1>>>(2,7);
+""".}