diff --git a/nsy/hip/.clang-format b/nsy/cuda/.clang-format similarity index 100% rename from nsy/hip/.clang-format rename to nsy/cuda/.clang-format diff --git a/nsy/cuda/.gitignore b/nsy/cuda/.gitignore new file mode 100644 index 0000000..ec0b70b --- /dev/null +++ b/nsy/cuda/.gitignore @@ -0,0 +1,4 @@ +nsy +nsy.o +nsy.cpp +nsy.cu.hip \ No newline at end of file diff --git a/nsy/cuda/Makefile b/nsy/cuda/Makefile new file mode 100644 index 0000000..f56dece --- /dev/null +++ b/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 \ No newline at end of file diff --git a/nsy/cuda/README.md b/nsy/cuda/README.md new file mode 100644 index 0000000..3368b19 --- /dev/null +++ b/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 +``` diff --git a/nsy/hip/nsy_hip.cu b/nsy/cuda/nsy.cu similarity index 61% rename from nsy/hip/nsy_hip.cu rename to nsy/cuda/nsy.cu index e79cc9c..da4c61c 100644 --- a/nsy/hip/nsy_hip.cu +++ b/nsy/cuda/nsy.cu @@ -1,27 +1,21 @@ -#ifdef _MSC_VER -#include "msvc_defines.h" -#endif -#include +#include +#include #include #include #include #include -#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 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'; diff --git a/nsy/hip/.gitignore b/nsy/hip/.gitignore deleted file mode 100644 index 37c25a7..0000000 --- a/nsy/hip/.gitignore +++ /dev/null @@ -1,3 +0,0 @@ -nsy_hip -nsy_hip.exe -nsy_hip.o \ No newline at end of file diff --git a/nsy/hip/Makefile b/nsy/hip/Makefile deleted file mode 100644 index 45c1cd2..0000000 --- a/nsy/hip/Makefile +++ /dev/null @@ -1,20 +0,0 @@ -HIP_PATH ?= $(wildcard /opt/rocm) -CXX = $(HIP_PATH)/bin/hipcc -CXXFLAGS = -O -Wall -W -pedantic - -TARGET = nsy_hip - -ifeq ($(OS), Windows_NT) - TARGET = nsy_hip.exe -endif - -SOURCES = $(wildcard ./*.cu) -OBJECTS = $(patsubst %.cu, %.o, $(SOURCES)) - -all: $(TARGET) - -$(TARGET): $(OBJECTS) - $(CXX) $(CXXFLAGS) -o $@ $^ - -%.o: %.cu - $(CXX) $(CXXFLAGS) -c -o $@ $< diff --git a/nsy/hip/msvc_defines.h b/nsy/hip/msvc_defines.h deleted file mode 100644 index 924bd16..0000000 --- a/nsy/hip/msvc_defines.h +++ /dev/null @@ -1,62 +0,0 @@ -/* -Copyright (c) 2022-present Advanced Micro Devices, Inc. All rights reserved. - -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: - -The above copyright notice and this permission notice shall be included in -all copies or substantial portions of the Software. - -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -THE SOFTWARE. -*/ - -/** - * @file msvc_defines.h - * @brief TODO-doc - */ - -#ifndef HIP_INCLUDE_HIP_AMD_DETAIL_MSVC_DEFINES_H -#define HIP_INCLUDE_HIP_AMD_DETAIL_MSVC_DEFINES_H - -#if defined(_MSC_VER) && defined(__INTELLISENSE__) -#define __CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ -#define __restrict__ -#define __inline__ \ - __inline -#define __no_return__ \ - __declspec(noreturn) -#define __noinline__ \ - __declspec(noinline) -#define __forceinline__ \ - __forceinline -#define __align__(n) \ - __declspec(align(n)) -#define __thread__ \ - __declspec(thread) -#define __import__ \ - __declspec(dllimport) -#define __export__ \ - __declspec(dllexport) -#define __annotate__(a) \ - __declspec(a) -#define __location__(a) \ - __annotate__(__##a##__) -#define __host__ __location__(host) -#define __device__ __location__(device) -#define __global__ __location__(global) -#define __shared__ __location__(shared) -#define __constant__ __location__(constant) -#define __attribute__(x) -#endif - -#endif diff --git a/nsy/hip/nsy_hip.sln b/nsy/hip/nsy_hip.sln deleted file mode 100644 index 574b0b0..0000000 --- a/nsy/hip/nsy_hip.sln +++ /dev/null @@ -1,25 +0,0 @@ - -Microsoft Visual Studio Solution File, Format Version 12.00 -# Visual Studio Version 17 -VisualStudioVersion = 17.10.34928.147 -MinimumVisualStudioVersion = 10.0.40219.1 -Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "nsy_hip", "nsy_hip.vcxproj", "{7DE6465A-D391-462B-9EB9-439DBDBBAAE9}" -EndProject -Global - GlobalSection(SolutionConfigurationPlatforms) = preSolution - Debug|x64 = Debug|x64 - Release|x64 = Release|x64 - EndGlobalSection - GlobalSection(ProjectConfigurationPlatforms) = postSolution - {7DE6465A-D391-462B-9EB9-439DBDBBAAE9}.Debug|x64.ActiveCfg = Debug|x64 - {7DE6465A-D391-462B-9EB9-439DBDBBAAE9}.Debug|x64.Build.0 = Debug|x64 - {7DE6465A-D391-462B-9EB9-439DBDBBAAE9}.Release|x64.ActiveCfg = Release|x64 - {7DE6465A-D391-462B-9EB9-439DBDBBAAE9}.Release|x64.Build.0 = Release|x64 - EndGlobalSection - GlobalSection(SolutionProperties) = preSolution - HideSolutionNode = FALSE - EndGlobalSection - GlobalSection(ExtensibilityGlobals) = postSolution - SolutionGuid = {24E35829-4F1C-48DE-8F96-B53242339FE4} - EndGlobalSection -EndGlobal diff --git a/nsy/hip/nsy_hip.vcxproj b/nsy/hip/nsy_hip.vcxproj deleted file mode 100644 index ac34574..0000000 --- a/nsy/hip/nsy_hip.vcxproj +++ /dev/null @@ -1,106 +0,0 @@ - - - - - Debug - x64 - - - Release - x64 - - - - 17.0 - 5.7 - {7de6465a-d391-462b-9eb9-439dbdbbaae9} - Win32Proj - nsy_hip - $(LatestTargetPlatformVersion) - - - - Application - true - HIP clang 5.7 - Unicode - - - Application - false - HIP clang 5.7 - Unicode - - - - - - - - - - - - - - - - - true - - - false - - - - Level1 - __clang__;__HIP__;_DEBUG;_CONSOLE;%(PreprocessorDefinitions) - - - Console - true - - - - - __CUDACC__;_DEBUG;_CONSOLE;%(PreprocessorDefinitions) - - - - - true - - - - - Level2 - true - __clang__;__HIP__;NDEBUG;_CONSOLE;%(PreprocessorDefinitions) - - - Console - true - true - - - - - true - __CUDACC__;NDEBUG;_CONSOLE;%(PreprocessorDefinitions) - - - UseLinkTimeCodeGeneration - - - - - - - - - - - - - - \ No newline at end of file diff --git a/nsy/hip/nsy_hip.vcxproj.filters b/nsy/hip/nsy_hip.vcxproj.filters deleted file mode 100644 index e7ed3d0..0000000 --- a/nsy/hip/nsy_hip.vcxproj.filters +++ /dev/null @@ -1,27 +0,0 @@ - - - - - {610c35ef-46f4-4478-aecc-5e1b74587212} - cpp;c;cc;cxx;c++;def;odl;idl;hpj;bat;asm;asmx;hip;cu - - - {812edc2e-f1be-4c49-8507-a0a86d776e17} - h;hh;hpp;hxx;h++;hm;inl;inc;ipp;xsd;cuh - - - {29b5708c-0216-4869-a869-cfb23b853454} - rc;ico;cur;bmp;dlg;rc2;rct;bin;rgs;gif;jpg;jpeg;jpe;resx;tiff;tif;png;wav;mfcribbon-ms - - - - - Source Files - - - - - Header Files - - - \ No newline at end of file diff --git a/nsy/hip/nsy_hip.vcxproj.user b/nsy/hip/nsy_hip.vcxproj.user deleted file mode 100644 index 88a5509..0000000 --- a/nsy/hip/nsy_hip.vcxproj.user +++ /dev/null @@ -1,4 +0,0 @@ - - - - \ No newline at end of file