From 209f31525fee2c909c6022b5fddbdcc01d09699a Mon Sep 17 00:00:00 2001 From: mhbliao <47895780+mhbliao@users.noreply.github.com> Date: Fri, 6 Dec 2019 12:49:17 -0500 Subject: [PATCH 1/5] Fix `hipExtLaunchMultiKernelMultiDevice` refactoring. (#1714) - Use the correct condition for HIP VDI runtime. --- hipamd/include/hip/hcc_detail/hip_runtime_api.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 4397e7fe50..e38c3a2dd7 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -2956,7 +2956,7 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor( hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk, unsigned int flags); -#if defined(__clang__) && defined(__HIP__) +#if __HIP_VDI__ /** * @brief Launches kernels on multiple devices and guarantees all specified kernels are dispatched * on respective streams before enqueuing any other work on the specified streams from any other threads From ef405daaeed9fa0931bc9ad4cd9ace5fcef49b20 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Sat, 7 Dec 2019 11:20:17 +0300 Subject: [PATCH 2/5] [HIPIFY][fix][#1246][#1655] Sort input files based on their dependency graph + Implemented by using clang Driver infrastructure [ToDo] Add tests for in-place hipification of multiple files --- hipamd/hipify-clang/src/main.cpp | 32 ++++++++++++++++++++++++++++++++ 1 file changed, 32 insertions(+) diff --git a/hipamd/hipify-clang/src/main.cpp b/hipamd/hipify-clang/src/main.cpp index b4200bd109..62c4e1d5eb 100644 --- a/hipamd/hipify-clang/src/main.cpp +++ b/hipamd/hipify-clang/src/main.cpp @@ -34,6 +34,14 @@ THE SOFTWARE. #include "ArgParse.h" #include "StringUtils.h" #include "llvm/Support/Debug.h" +#include "clang/Basic/Diagnostic.h" +#include "clang/Basic/DiagnosticIDs.h" +#include "clang/Basic/DiagnosticOptions.h" +#include "clang/Driver/Driver.h" +#include "clang/Driver/Compilation.h" +#include "clang/Driver/Tool.h" +#include "clang/Frontend/TextDiagnosticPrinter.h" + #if LLVM_VERSION_MAJOR < 8 #include "llvm/Support/Path.h" #endif @@ -139,6 +147,30 @@ int main(int argc, const char **argv) { if (PrintStats) { statPrint = &llvm::errs(); } + if (fileSources.size() > 1) { + IntrusiveRefCntPtr diagOpts(new clang::DiagnosticOptions()); + clang::TextDiagnosticPrinter diagClient(llvm::errs(), &*diagOpts); + clang::DiagnosticsEngine Diagnostics(IntrusiveRefCntPtr(new clang::DiagnosticIDs()), &*diagOpts, &diagClient, false); + std::unique_ptr driver(new clang::driver::Driver("", "nvptx64-nvidia-cuda", Diagnostics)); + SmallVector Args(argv, argv + argc); + std::unique_ptr C(driver->BuildCompilation(Args)); + std::vector fileSourcesOrdered; + for (const auto &J : C->getJobs()) { + if (std::string(J.getCreator().getName()) != "clang") continue; + const auto &JA = J.getArguments(); + for (size_t i = 0; i < JA.size(); ++i) { + const auto &A = std::string(JA[i]); + if (std::find(fileSources.begin(), fileSources.end(), A) != fileSources.end() && + i > 0 && std::string(JA[i-1]) == "-main-file-name") { + fileSourcesOrdered.push_back(A); + } + } + } + if (!fileSourcesOrdered.empty()) { + std::reverse(fileSourcesOrdered.begin(), fileSourcesOrdered.end()); + fileSources = fileSourcesOrdered; + } + } for (const auto & src : fileSources) { // Create a copy of the file to work on. When we're done, we'll move this onto the // output (which may mean overwriting the input, if we're in-place). From 72b1d16115d93b1a9ada0fab9eeb8193ca60e725 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Sat, 7 Dec 2019 13:50:13 +0300 Subject: [PATCH 3/5] [HIPIFY] main source file refactoring --- hipamd/hipify-clang/src/main.cpp | 156 ++++++++++++++++--------------- 1 file changed, 81 insertions(+), 75 deletions(-) diff --git a/hipamd/hipify-clang/src/main.cpp b/hipamd/hipify-clang/src/main.cpp index 62c4e1d5eb..91ff9c36d4 100644 --- a/hipamd/hipify-clang/src/main.cpp +++ b/hipamd/hipify-clang/src/main.cpp @@ -50,6 +50,81 @@ constexpr auto DEBUG_TYPE = "cuda2hip"; namespace ct = clang::tooling; +void sortInputFiles(int argc, const char **argv, std::vector &files) { + if (files.size() < 2) return; + IntrusiveRefCntPtr diagOpts(new clang::DiagnosticOptions()); + clang::TextDiagnosticPrinter diagClient(llvm::errs(), &*diagOpts); + clang::DiagnosticsEngine Diagnostics(IntrusiveRefCntPtr(new clang::DiagnosticIDs()), &*diagOpts, &diagClient, false); + std::unique_ptr driver(new clang::driver::Driver("", "nvptx64-nvidia-cuda", Diagnostics)); + SmallVector Args(argv, argv + argc); + std::unique_ptr C(driver->BuildCompilation(Args)); + std::vector sortedFiles; + for (const auto &J : C->getJobs()) { + if (std::string(J.getCreator().getName()) != "clang") continue; + const auto &JA = J.getArguments(); + for (size_t i = 0; i < JA.size(); ++i) { + const auto &A = std::string(JA[i]); + if (std::find(files.begin(), files.end(), A) != files.end() && + i > 0 && std::string(JA[i - 1]) == "-main-file-name") { + sortedFiles.push_back(A); + } + } + } + if (sortedFiles.empty()) return; + std::reverse(sortedFiles.begin(), sortedFiles.end()); + files.assign(sortedFiles.begin(), sortedFiles.end()); +} + +void appendArgumentsAdjusters(ct::RefactoringTool &Tool, const std::string &sSourceAbsPath) { + if (!IncludeDirs.empty()) { + for (std::string s : IncludeDirs) { + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(s.c_str(), ct::ArgumentInsertPosition::BEGIN)); + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-I", ct::ArgumentInsertPosition::BEGIN)); + } + } + if (!MacroNames.empty()) { + for (std::string s : MacroNames) { + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(s.c_str(), ct::ArgumentInsertPosition::BEGIN)); + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-D", ct::ArgumentInsertPosition::BEGIN)); + } + } + // Includes for clang's CUDA wrappers for using by packaged hipify-clang + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("./include", ct::ArgumentInsertPosition::BEGIN)); + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-isystem", ct::ArgumentInsertPosition::BEGIN)); + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("./include/cuda_wrappers", ct::ArgumentInsertPosition::BEGIN)); + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-isystem", ct::ArgumentInsertPosition::BEGIN)); + // Ensure at least c++11 is used. + std::string stdCpp = "-std=c++11"; +#if defined(_MSC_VER) + stdCpp = "-std=c++14"; +#endif + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(stdCpp.c_str(), ct::ArgumentInsertPosition::BEGIN)); + std::string sInclude = "-I" + sys::path::parent_path(sSourceAbsPath).str(); +#if defined(HIPIFY_CLANG_RES) + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES, ct::ArgumentInsertPosition::BEGIN)); +#endif + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sInclude.c_str(), ct::ArgumentInsertPosition::BEGIN)); + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-fno-delayed-template-parsing", ct::ArgumentInsertPosition::BEGIN)); + if (llcompat::pragma_once_outside_header()) { + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-Wno-pragma-once-outside-header", ct::ArgumentInsertPosition::BEGIN)); + } + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("--cuda-host-only", ct::ArgumentInsertPosition::BEGIN)); + if (!CudaGpuArch.empty()) { + std::string sCudaGpuArch = "--cuda-gpu-arch=" + CudaGpuArch; + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sCudaGpuArch.c_str(), ct::ArgumentInsertPosition::BEGIN)); + } + if (!CudaPath.empty()) { + std::string sCudaPath = "--cuda-path=" + CudaPath; + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sCudaPath.c_str(), ct::ArgumentInsertPosition::BEGIN)); + } + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("cuda", ct::ArgumentInsertPosition::BEGIN)); + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-x", ct::ArgumentInsertPosition::BEGIN)); + if (Verbose) { + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-v", ct::ArgumentInsertPosition::END)); + } + Tool.appendArgumentsAdjuster(ct::getClangSyntaxOnlyAdjuster()); +} + int main(int argc, const char **argv) { std::vector new_argv(argv, argv + argc); if (std::find(new_argv.begin(), new_argv.end(), std::string("--")) == new_argv.end()) { @@ -127,7 +202,7 @@ int main(int argc, const char **argv) { } // Arguments for the Statistics print routines. std::unique_ptr csv = nullptr; - llvm::raw_ostream* statPrint = nullptr; + llvm::raw_ostream *statPrint = nullptr; bool create_csv = false; if (!OutputStatsFilename.empty()) { PrintStatsCSV = true; @@ -147,31 +222,8 @@ int main(int argc, const char **argv) { if (PrintStats) { statPrint = &llvm::errs(); } - if (fileSources.size() > 1) { - IntrusiveRefCntPtr diagOpts(new clang::DiagnosticOptions()); - clang::TextDiagnosticPrinter diagClient(llvm::errs(), &*diagOpts); - clang::DiagnosticsEngine Diagnostics(IntrusiveRefCntPtr(new clang::DiagnosticIDs()), &*diagOpts, &diagClient, false); - std::unique_ptr driver(new clang::driver::Driver("", "nvptx64-nvidia-cuda", Diagnostics)); - SmallVector Args(argv, argv + argc); - std::unique_ptr C(driver->BuildCompilation(Args)); - std::vector fileSourcesOrdered; - for (const auto &J : C->getJobs()) { - if (std::string(J.getCreator().getName()) != "clang") continue; - const auto &JA = J.getArguments(); - for (size_t i = 0; i < JA.size(); ++i) { - const auto &A = std::string(JA[i]); - if (std::find(fileSources.begin(), fileSources.end(), A) != fileSources.end() && - i > 0 && std::string(JA[i-1]) == "-main-file-name") { - fileSourcesOrdered.push_back(A); - } - } - } - if (!fileSourcesOrdered.empty()) { - std::reverse(fileSourcesOrdered.begin(), fileSourcesOrdered.end()); - fileSources = fileSourcesOrdered; - } - } - for (const auto & src : fileSources) { + sortInputFiles(argc, argv, fileSources); + for (const auto &src : fileSources) { // Create a copy of the file to work on. When we're done, we'll move this onto the // output (which may mean overwriting the input, if we're in-place). // Should we fail for some reason, we'll just leak this file and not corrupt the input. @@ -224,56 +276,10 @@ int main(int argc, const char **argv) { // because that'll break relative includes, and we don't want to overwrite the input file. // So what we do is operate on a copy, which we then move to the output. ct::RefactoringTool Tool(OptionsParser.getCompilations(), std::string(tmpFile.c_str())); - ct::Replacements& replacementsToUse = llcompat::getReplacements(Tool, tmpFile.c_str()); + ct::Replacements &replacementsToUse = llcompat::getReplacements(Tool, tmpFile.c_str()); ReplacementsFrontendActionFactory actionFactory(&replacementsToUse); - if (!IncludeDirs.empty()) { - for (std::string s : IncludeDirs) { - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(s.c_str(), ct::ArgumentInsertPosition::BEGIN)); - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-I", ct::ArgumentInsertPosition::BEGIN)); - } - } - if (!MacroNames.empty()) { - for (std::string s : MacroNames) { - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(s.c_str(), ct::ArgumentInsertPosition::BEGIN)); - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-D", ct::ArgumentInsertPosition::BEGIN)); - } - } - // Includes for clang's CUDA wrappers for using by packaged hipify-clang - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("./include", ct::ArgumentInsertPosition::BEGIN)); - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-isystem", ct::ArgumentInsertPosition::BEGIN)); - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("./include/cuda_wrappers", ct::ArgumentInsertPosition::BEGIN)); - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-isystem", ct::ArgumentInsertPosition::BEGIN)); - // Ensure at least c++11 is used. - std::string stdCpp = "-std=c++11"; -#if defined(_MSC_VER) - stdCpp = "-std=c++14"; -#endif - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(stdCpp.c_str(), ct::ArgumentInsertPosition::BEGIN)); - std::string sInclude = "-I" + sys::path::parent_path(sSourceAbsPath).str(); -#if defined(HIPIFY_CLANG_RES) - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES, ct::ArgumentInsertPosition::BEGIN)); -#endif - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sInclude.c_str(), ct::ArgumentInsertPosition::BEGIN)); - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-fno-delayed-template-parsing", ct::ArgumentInsertPosition::BEGIN)); - if (llcompat::pragma_once_outside_header()) { - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-Wno-pragma-once-outside-header", ct::ArgumentInsertPosition::BEGIN)); - } - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("--cuda-host-only", ct::ArgumentInsertPosition::BEGIN)); - if (!CudaGpuArch.empty()) { - std::string sCudaGpuArch = "--cuda-gpu-arch=" + CudaGpuArch; - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sCudaGpuArch.c_str(), ct::ArgumentInsertPosition::BEGIN)); - } - if (!CudaPath.empty()) { - std::string sCudaPath = "--cuda-path=" + CudaPath; - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sCudaPath.c_str(), ct::ArgumentInsertPosition::BEGIN)); - } - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("cuda", ct::ArgumentInsertPosition::BEGIN)); - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-x", ct::ArgumentInsertPosition::BEGIN)); - if (Verbose) { - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-v", ct::ArgumentInsertPosition::END)); - } - Tool.appendArgumentsAdjuster(ct::getClangSyntaxOnlyAdjuster()); - Statistics& currentStat = Statistics::current(); + appendArgumentsAdjusters(Tool, sSourceAbsPath); + Statistics ¤tStat = Statistics::current(); // Hipify _all_ the things! if (Tool.runAndSave(&actionFactory)) { currentStat.hasErrors = true; From 486ec185e76add51e3cd15fc8f370a9e38b5f718 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Sun, 8 Dec 2019 15:13:28 +0300 Subject: [PATCH 4/5] [HIPIFY] Add cleanupHipifyOptions functionality Needed for sorting input files to avoid clang errors on unsupported options --- hipamd/hipify-clang/src/main.cpp | 35 +++++++++++++++++++++++++++++++- 1 file changed, 34 insertions(+), 1 deletion(-) diff --git a/hipamd/hipify-clang/src/main.cpp b/hipamd/hipify-clang/src/main.cpp index 91ff9c36d4..65376df22c 100644 --- a/hipamd/hipify-clang/src/main.cpp +++ b/hipamd/hipify-clang/src/main.cpp @@ -50,13 +50,46 @@ constexpr auto DEBUG_TYPE = "cuda2hip"; namespace ct = clang::tooling; +void cleanupHipifyOptions(std::vector &args) { + std::vector hipifyOptions = {"-perl", "-python", "-roc", "-inplace", + "-no-backup", "-no-output", "-print-stats", + "-print-stats-csv", "-examine", "-save-temps", + "-skip-excluded-preprocessor-conditional-blocks"}; + for (const auto &a : hipifyOptions) { + args.erase(std::remove(args.begin(), args.end(), a), args.end()); + args.erase(std::remove(args.begin(), args.end(), "-" + a), args.end()); + } + std::vector hipifyDirOptions = {"-o-dir", "-o-hipify-perl-dir", "-o-stats", + "-o-python-map-dir", "-temp-dir"}; + for (const auto &a : hipifyDirOptions) { + // remove all pairs of arguments "-option value" + auto it = args.erase(std::remove(args.begin(), args.end(), a), args.end()); + if (it != args.end()) { + args.erase(it); + } + // remove all pairs of arguments "--option value" + it = args.erase(std::remove(args.begin(), args.end(), "-" + a), args.end()); + if (it != args.end()) { + args.erase(it); + } + // remove all "-option=value" and "--option=value" + args.erase( + std::remove_if(args.begin(), args.end(), + [a](const std::string &s) { return s.find(a + "=") == 0 || s.find("-" + a + "=") == 0; } + ), + args.end() + ); + } +} + void sortInputFiles(int argc, const char **argv, std::vector &files) { if (files.size() < 2) return; IntrusiveRefCntPtr diagOpts(new clang::DiagnosticOptions()); clang::TextDiagnosticPrinter diagClient(llvm::errs(), &*diagOpts); clang::DiagnosticsEngine Diagnostics(IntrusiveRefCntPtr(new clang::DiagnosticIDs()), &*diagOpts, &diagClient, false); std::unique_ptr driver(new clang::driver::Driver("", "nvptx64-nvidia-cuda", Diagnostics)); - SmallVector Args(argv, argv + argc); + std::vector Args(argv, argv + argc); + cleanupHipifyOptions(Args); std::unique_ptr C(driver->BuildCompilation(Args)); std::vector sortedFiles; for (const auto &J : C->getJobs()) { From 280a8ebadf2d690533c5e5268780395d7d12fd25 Mon Sep 17 00:00:00 2001 From: jglaser Date: Mon, 9 Dec 2019 23:10:15 -0500 Subject: [PATCH 5/5] fix linking of vector types with gcc (#1690) * fix linking of vector types when linking hipcc objects with gcc * use __atribute__((vector_size)) with both clang and gcc and reinstate nonaligned n=3 vector type * use implicit conversion to value and ext_vector_type when available * Alternate formulation for GCC compatibility * Built-in arrays don't mix well with placement new * Fix typo * Add conversions to enum * Fix Scalar_accessor assignment. * Update hip_vector_types.h * stir up the underlying_type hideous mess This fixes the HIP build issue "error: only enumeration types have underlying types". --- .../include/hip/hcc_detail/hip_vector_types.h | 114 +++++++++++++++--- 1 file changed, 98 insertions(+), 16 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/hip_vector_types.h b/hipamd/include/hip/hcc_detail/hip_vector_types.h index 582a359fbd..0e8dad595a 100644 --- a/hipamd/include/hip/hcc_detail/hip_vector_types.h +++ b/hipamd/include/hip/hcc_detail/hip_vector_types.h @@ -34,15 +34,31 @@ THE SOFTWARE. #include "hip/hcc_detail/host_defines.h" -#if !defined(_MSC_VER) || __clang__ -#if defined(__clang__) - #define __NATIVE_VECTOR__(n, ...) __attribute__((ext_vector_type(n))) +#if !defined(_MSC_VER) + #if __has_attribute(ext_vector_type) + #define __NATIVE_VECTOR__(n, ...) __attribute__((ext_vector_type(n))) + #else + #define __NATIVE_VECTOR__(n, ...) [n] #endif -#if defined(__cplusplus) && defined(__clang__) +#if defined(__cplusplus) + #include #include #include + namespace hip_impl { + template struct Scalar_accessor; + } // Namespace hip_impl. + + namespace std { + template + struct is_integral> + : is_integral {}; + template + struct is_floating_point> + : is_floating_point {}; + } // Namespace std. + namespace hip_impl { template struct Scalar_accessor { @@ -93,6 +109,27 @@ THE SOFTWARE. operator T() const noexcept { return data[idx]; } __host__ __device__ operator T() const volatile noexcept { return data[idx]; } + + // The conversions to enum are fairly ghastly, but unfortunately used in + // some pre-existing, difficult to modify, code. + template< + typename U, + typename std::enable_if< + !std::is_same{} && + std::is_enum{} && + std::is_convertible< + T, typename std::enable_if::value, std::underlying_type>::type::type>{}>::type* = nullptr> + __host__ __device__ + operator U() const noexcept { return static_cast(data[idx]); } + template< + typename U, + typename std::enable_if< + !std::is_same{} && + std::is_enum{} && + std::is_convertible< + T, typename std::enable_if::value, std::underlying_type>::type::type>{}>::type* = nullptr> + __host__ __device__ + operator U() const volatile noexcept { return static_cast(data[idx]); } __host__ __device__ operator T&() noexcept { @@ -107,7 +144,13 @@ THE SOFTWARE. __host__ __device__ Address operator&() const noexcept { return Address{this}; } + + __host__ __device__ + Scalar_accessor& operator=(const Scalar_accessor& x) noexcept { + data[idx] = x.data[idx]; + return *this; + } __host__ __device__ Scalar_accessor& operator=(T x) noexcept { data[idx] = x; @@ -179,7 +222,7 @@ THE SOFTWARE. typename std::enable_if< std::is_convertible{}>::type* = nullptr> __host__ __device__ - Scalar_accessor& operator/=(T x) noexcept { + Scalar_accessor& operator/=(U x) noexcept { data[idx] /= x; return *this; } @@ -245,7 +288,7 @@ THE SOFTWARE. template struct HIP_vector_base { - typedef T Native_vec_ __NATIVE_VECTOR__(1, T); + using Native_vec_ = T __NATIVE_VECTOR__(1, T); union { Native_vec_ data; @@ -253,11 +296,22 @@ THE SOFTWARE. }; using value_type = T; + + __host__ __device__ + HIP_vector_base& operator=(const HIP_vector_base& x) noexcept { + #if __has_attribute(ext_vector_type) + data = x.data; + #else + data[0] = x.data[0]; + #endif + + return *this; + } }; template struct HIP_vector_base { - typedef T Native_vec_ __NATIVE_VECTOR__(2, T); + using Native_vec_ = T __NATIVE_VECTOR__(2, T); union { Native_vec_ data; @@ -266,6 +320,18 @@ THE SOFTWARE. }; using value_type = T; + + __host__ __device__ + HIP_vector_base& operator=(const HIP_vector_base& x) noexcept { + #if __has_attribute(ext_vector_type) + data = x.data; + #else + data[0] = x.data[0]; + data[1] = x.data[1]; + #endif + + return *this; + } }; template @@ -404,15 +470,11 @@ THE SOFTWARE. return *this; } - using Vec3_cmp = int __NATIVE_VECTOR__(3, int); + using Vec3_cmp = int __attribute__((vector_size(4 * sizeof(int)))); __host__ __device__ Vec3_cmp operator==(const Native_vec_& x) const noexcept { - Vec3_cmp r; - r[0] = d[0] == x.d[0]; - r[1] = d[1] == x.d[1]; - r[2] = d[2] == x.d[2]; - return r; + return Vec3_cmp{d[0] == x.d[0], d[1] == x.d[1], d[2] == x.d[2]}; } }; @@ -430,7 +492,7 @@ THE SOFTWARE. template struct HIP_vector_base { - typedef T Native_vec_ __NATIVE_VECTOR__(4, T); + using Native_vec_ = T __NATIVE_VECTOR__(4, T); union { Native_vec_ data; @@ -441,6 +503,20 @@ THE SOFTWARE. }; using value_type = T; + + __host__ __device__ + HIP_vector_base& operator=(const HIP_vector_base& x) noexcept { + #if __has_attribute(ext_vector_type) + data = x.data; + #else + data[0] = x.data[0]; + data[1] = x.data[1]; + data[2] = x.data[2]; + data[3] = x.data[3]; + #endif + + return *this; + } }; template @@ -464,7 +540,14 @@ THE SOFTWARE. typename std::enable_if< (rank > 1) && sizeof...(Us) == rank>::type* = nullptr> inline __host__ __device__ - HIP_vector_type(Us... xs) noexcept { data = Native_vec_{static_cast(xs)...}; } + HIP_vector_type(Us... xs) noexcept + { + #if __has_attribute(ext_vector_type) + new (&data) Native_vec_{static_cast(xs)...}; + #else + new (&data) std::array{static_cast(xs)...}; + #endif + } inline __host__ __device__ HIP_vector_type(const HIP_vector_type&) = default; inline __host__ __device__ @@ -651,7 +734,6 @@ THE SOFTWARE. } }; - template inline __host__ __device__ HIP_vector_type operator+(