feat: migrate HIP implementation to CUDA with HIPIFY support.

nsy/cuda/.gitignore

@@ -0,0 +1,4 @@
+nsy
+nsy.o
+nsy.cpp
+nsy.cu.hip

nsy/cuda/Makefile

@@ -0,0 +1,46 @@
+ifndef CXX
+CXX := nvcc
+endif
+CXXFLAGS := -O3 -Wall -Wno-pedantic
+ifeq ($(CXX), nvcc)
+CXXFLAGS := -Xcompiler "$(CXXFLAGS)"
+endif
+LDFLAGS :=
+
+TARGET := nsy
+
+SOURCES := $(wildcard ./*.cu)
+OBJECTS := $(patsubst %.cu, %.o, $(SOURCES))
+
+all: $(TARGET)
+
+$(TARGET): $(OBJECTS)
+	$(CXX) $(CXXFLAGS) $(LDFLAGS) -o $@ $^
+
+ifneq (,$(filter $(CXX),clang clang++))
+CXXFLAGS += -march=native -std=c++17
+LDFLAGS += -ltbb
+
+_SOURCES := $(SOURCES)
+SOURCES := $(patsubst %.cu, %.cpp, $(SOURCES))
+
+$(SOURCES): $(_SOURCES)
+	hipify-clang --hip-kernel-execution-syntax -o $@ $<
+endif
+
+ifeq ($(CXX), hipcc)
+_SOURCES := $(SOURCES)
+SOURCES := $(patsubst %.cu, %.cu.hip, $(SOURCES))
+
+$(SOURCES): $(_SOURCES)
+	hipify-clang -o $@ $<
+endif
+
+$(OBJECTS): $(SOURCES)
+	$(CXX) $(CXXFLAGS) -c -o $@ $<
+
+clean:
+	rm -f nsy
+	rm -f nsy.o
+	rm -f nsy.cpp
+	rm -f nsy.cu.hip

nsy/cuda/README.md

@@ -0,0 +1,26 @@
+# Build for CPU
+
+Requirements: clang compiler, [HIP-CPU](https://github.com/ROCm/HIP-CPU)
+
+```
+export CXX=clang++
+make
+```
+
+# Build for NVIDIA GPUs
+
+Requirements: NVIDIA Cuda compiler
+
+```
+export CXX=nvcc
+make
+```
+
+# Build for AMDGPUs
+
+Requirements: HIP C++ compiler, AMD clang compiler
+
+```
+export CXX=hipcc
+make
+```

nsy/hip/nsy_hip.cu → nsy/cuda/nsy.cu

@@ -1,27 +1,21 @@
-#ifdef _MSC_VER
-#include "msvc_defines.h"
-#endif
-#include <hip/hip_runtime.h>
+#include <cmath>
+#include <cuda_runtime.h>
 #include <iostream>
 #include <locale>
 #include <string>
 #include <vector>
 
-#define HIP_CHECK(status)                                                      \
-  if (status != hipSuccess) {                                                  \
-    fprintf(stderr, "error: '%s' at %s:%d\n", hipGetErrorString(status),       \
+#define CUDA_CHECK(status)                                                     \
+  if (status != cudaSuccess) {                                                 \
+    fprintf(stderr, "error: '%s' at %s:%d\n", cudaGetErrorString(status),      \
             __FILE__, __LINE__);                                               \
     exit(0);                                                                   \
   }
 
 __global__ void kernel(const wchar_t *in, wchar_t *out, const size_t dim,
                        const wchar_t mark) {
-  const int idx = hipThreadIdx_x * dim + hipThreadIdx_y;
-  if (hipBlockIdx_x) {
-    out[idx * 2 + 1] = mark;
-  } else {
-    out[idx * 2] = in[idx];
-  }
+  const int idx = threadIdx.x * dim + threadIdx.y;
+  out[idx * 2 + blockIdx.x] = blockIdx.x == 0 ? in[idx] : mark;
 }
 
 wchar_t get_mark(char *s) {
@@ -41,7 +35,7 @@ int main(int argc, char *argv[]) {
   std::vector<wchar_t> str;
 
   wchar_t c;
-  while ((c = std::wcin.get()) != WEOF) {
+  while ((c = std::wcin.get()) != (wchar_t)WEOF) {
     if (c == ' ') {
       continue;
     }
@@ -55,21 +49,22 @@ int main(int argc, char *argv[]) {
   const size_t input_size = sizeof(wchar_t) * dim * dim;
   const size_t output_length = dim * 2 * dim;
   const size_t output_size = sizeof(wchar_t) * output_length;
-  HIP_CHECK(hipMalloc((void **)&input, input_size));
-  HIP_CHECK(hipMalloc((void **)&output, output_size));
-  HIP_CHECK(hipMemcpy(input, str.data(), sizeof(wchar_t) * length,
-                      hipMemcpyHostToDevice));
+  CUDA_CHECK(cudaMalloc((void **)&input, input_size));
+  CUDA_CHECK(cudaMalloc((void **)&output, output_size));
+  CUDA_CHECK(cudaMemcpy(input, str.data(), sizeof(wchar_t) * length,
+                        cudaMemcpyHostToDevice));
 
   const wchar_t mark = argc < 2 ? L'\xFF01' : get_mark(argv[1]);
   kernel<<<2, dim3(dim, dim), 0, 0>>>(input, output, dim, mark);
 
   auto result = new wchar_t[output_length + 2]; // mark ... \x0000
   *result = mark;
-  HIP_CHECK(hipMemcpy(result + 1, output, output_size, hipMemcpyDeviceToHost));
-  result[length * 2 + 1] = 0;
+  CUDA_CHECK(
+      cudaMemcpy(result + 1, output, output_size, cudaMemcpyDeviceToHost));
+  result[length * 2 + 1] = L'\0';
 
-  HIP_CHECK(hipFree(input));
-  HIP_CHECK(hipFree(output));
+  CUDA_CHECK(cudaFree(input));
+  CUDA_CHECK(cudaFree(output));
 
   std::wcout << result << L'\n';