From 954535e0b5b0a8d87b9b896479a41d72fb61ffb3 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Thu, 29 Nov 2018 11:35:00 -0500 Subject: [PATCH 01/10] Revert "Fix issue in kernarg metadata parsing due to early finalization" This reverts commit 48c2d69d323b7866b8008d102675393cd4f1bd05. [ROCm/hip commit: 909609773ca1587d32a816cafe5e62bc4e24d2d3] --- projects/hip/src/program_state.cpp | 14 ++++++-------- 1 file changed, 6 insertions(+), 8 deletions(-) diff --git a/projects/hip/src/program_state.cpp b/projects/hip/src/program_state.cpp index 2fda546b54..7e42a44245 100644 --- a/projects/hip/src/program_state.cpp +++ b/projects/hip/src/program_state.cpp @@ -590,16 +590,14 @@ unordered_map>>& kernargs() { static once_flag f; call_once(f, []() { - for (auto&& blobs_for_one_arch : code_object_blobs()) { - for (auto && blob : blobs_for_one_arch.second) { - stringstream tmp{std::string{ - blob.cbegin(), blob.cend()}}; + for (auto&& blob : code_object_blobs()) { + stringstream tmp{std::string{ + blob.second.front().cbegin(), blob.second.front().cend()}}; - elfio reader; - if (!reader.load(tmp)) continue; + elfio reader; + if (!reader.load(tmp)) continue; - read_kernarg_metadata(reader, r); - } + read_kernarg_metadata(reader, r); } }); From a715dea79cd8ae8992d60256c0d036c997c6c9af Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Thu, 29 Nov 2018 11:38:31 -0500 Subject: [PATCH 02/10] Revert "Missing handling nullary `__global__` functions for mixed arity cases." This reverts commit 51c47fcc2e781ee0c6ce2bc58081f93fb7cc09db. [ROCm/hip commit: 8eb9b38e76b265aeaad3eca29aa77a1431e01e14] --- projects/hip/src/program_state.cpp | 9 ++------- 1 file changed, 2 insertions(+), 7 deletions(-) diff --git a/projects/hip/src/program_state.cpp b/projects/hip/src/program_state.cpp index 7e42a44245..97e9035e0d 100644 --- a/projects/hip/src/program_state.cpp +++ b/projects/hip/src/program_state.cpp @@ -409,18 +409,13 @@ void read_kernarg_metadata( auto fn = tmp.substr(dx, tmp.find_first_of("'\n", dx) - dx); dx += fn.size(); - - auto dx1 = tmp.find("CodeProps", dx); dx = tmp.find("Args:", dx); - if (dx1 < dx) { - dx = dx1; - continue; - } if (dx == string::npos) break; static constexpr decltype(tmp.size()) args_sz{5}; - dx = parse_args(tmp, dx + args_sz, dx1, kernargs[fn]); + dx = parse_args( + tmp, dx + args_sz, tmp.find("CodeProps", dx), kernargs[fn]); } while (true); } } From f04655ff6e5da0d211f50b26121398fa24fbc7b7 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Thu, 29 Nov 2018 11:38:34 -0500 Subject: [PATCH 03/10] Revert "Handle (odd) corner case of argumentless __global__ function." This reverts commit 591d4d1b5ff6d147ac293a3527dc252b868dd5b8. [ROCm/hip commit: aeca2c8cdcef36ebd5dab47dc23aa3a8c6f65e12] --- projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp b/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp index 5edddad6c5..ba9929c0a6 100644 --- a/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp +++ b/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp @@ -102,8 +102,6 @@ inline std::vector make_kernarg( static_assert(sizeof...(Formals) == sizeof...(Actuals), "The count of formal arguments must match the count of actuals."); - if (sizeof...(Formals) == 0) return {}; - const auto it = function_names().find( reinterpret_cast(kernel)); From 8bf51361e45147c935a1165d9596d481bc87a3ac Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Thu, 29 Nov 2018 11:38:35 -0500 Subject: [PATCH 04/10] Revert "If we've already seen a `__global__` function we do not need to re-parse" This reverts commit d3beba42d37d535eccd664b603f13a11f7abd2e5. [ROCm/hip commit: 09f87e41d948883878ea0602c94aabf593dafbf6] --- projects/hip/src/program_state.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/src/program_state.cpp b/projects/hip/src/program_state.cpp index 97e9035e0d..6b88da288e 100644 --- a/projects/hip/src/program_state.cpp +++ b/projects/hip/src/program_state.cpp @@ -347,7 +347,6 @@ size_t parse_args( size_t l, vector>& size_align) { if (f == l) return f; - if (!size_align.empty()) return l; do { static constexpr size_t size_sz{5}; @@ -596,6 +595,7 @@ unordered_map>>& kernargs() { } }); + for (auto&& x : r) std::cerr << x.first << std::endl; return r; } From a72a54996e9dad5d925a84638b1c7ccad4cc64cc Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Thu, 29 Nov 2018 11:38:36 -0500 Subject: [PATCH 05/10] Revert "Handle the very confusing dual encoding of the symbol name." This reverts commit 2a69dd1fb6f2022b0c05ef123501887f252959dd. [ROCm/hip commit: 71189c10c19ff840dd4b08c4c971fa372d2f4a95] --- projects/hip/src/program_state.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/hip/src/program_state.cpp b/projects/hip/src/program_state.cpp index 6b88da288e..b490c0ee25 100644 --- a/projects/hip/src/program_state.cpp +++ b/projects/hip/src/program_state.cpp @@ -404,10 +404,11 @@ void read_kernarg_metadata( if (dx == string::npos) break; static constexpr decltype(tmp.size()) name_sz{5}; - dx = tmp.find_first_not_of(" '", dx + name_sz); + dx = tmp.find_first_not_of(' ', dx + name_sz); - auto fn = tmp.substr(dx, tmp.find_first_of("'\n", dx) - dx); + auto fn = tmp.substr(dx, tmp.find('\n', dx) - dx); dx += fn.size(); + dx = tmp.find("Args:", dx); if (dx == string::npos) break; @@ -595,7 +596,6 @@ unordered_map>>& kernargs() { } }); - for (auto&& x : r) std::cerr << x.first << std::endl; return r; } From bca8022aa0e65b5da64e0d6bc095148285d8e3c0 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Thu, 29 Nov 2018 11:38:37 -0500 Subject: [PATCH 06/10] Revert "Rely on code object metadat for kernarg arguments alignof and sizeof." This reverts commit 6d64f5e112c17e2490d9d78d99ef09e4c5b3eeb4. [ROCm/hip commit: 1fbf6399620aededcd558e9ec630af50ecfd03bc] --- .../hip/hcc_detail/functional_grid_launch.hpp | 37 ++----- .../include/hip/hcc_detail/program_state.hpp | 2 - projects/hip/src/program_state.cpp | 102 +----------------- 3 files changed, 9 insertions(+), 132 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp b/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp index ba9929c0a6..e678f25aa2 100644 --- a/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp +++ b/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp @@ -33,7 +33,6 @@ THE SOFTWARE. #include #include -#include #include #include #include @@ -57,9 +56,7 @@ template < typename... Ts, typename std::enable_if::type* = nullptr> inline std::vector make_kernarg( - const std::tuple&, - const std::vector>&, - std::vector kernarg) { + std::vector kernarg, const std::tuple&) { return kernarg; } @@ -68,9 +65,7 @@ template < typename... Ts, typename std::enable_if::type* = nullptr> inline std::vector make_kernarg( - const std::tuple& formals, - const std::vector>& size_align, - std::vector kernarg) { + std::vector kernarg, const std::tuple& formals) { using T = typename std::tuple_element>::type; static_assert( @@ -85,42 +80,24 @@ inline std::vector make_kernarg( #endif kernarg.resize(round_up_to_next_multiple_nonnegative( - kernarg.size(), size_align[n].second) + - size_align[n].first); + kernarg.size(), alignof(T)) + sizeof(T)); - std::memcpy( - kernarg.data() + kernarg.size() - size_align[n].first, - &std::get(formals), - size_align[n].first); + new (kernarg.data() + kernarg.size() - sizeof(T)) T{std::get(formals)}; - return make_kernarg(formals, size_align, std::move(kernarg)); + return make_kernarg(std::move(kernarg), formals); } template inline std::vector make_kernarg( - void (*kernel)(Formals...), std::tuple actuals) { + void (*)(Formals...), std::tuple actuals) { static_assert(sizeof...(Formals) == sizeof...(Actuals), "The count of formal arguments must match the count of actuals."); - const auto it = function_names().find( - reinterpret_cast(kernel)); - - if (it == function_names().cend()) { - throw std::runtime_error{"Undefined __global__ function."}; - } - - const auto it1 = kernargs().find(it->second); - - if (it1 == kernargs().end()) { - throw std::runtime_error{ - "Missing metadata for __global__ function: " + it->second}; - } - std::tuple to_formals{std::move(actuals)}; std::vector kernarg; kernarg.reserve(sizeof(to_formals)); - return make_kernarg<0>(to_formals, it1->second, std::move(kernarg)); + return make_kernarg<0>(std::move(kernarg), to_formals); } void hipLaunchKernelGGLImpl(std::uintptr_t function_address, const dim3& numBlocks, diff --git a/projects/hip/include/hip/hcc_detail/program_state.hpp b/projects/hip/include/hip/hcc_detail/program_state.hpp index 92bef22172..bdb87b3509 100644 --- a/projects/hip/include/hip/hcc_detail/program_state.hpp +++ b/projects/hip/include/hip/hcc_detail/program_state.hpp @@ -99,8 +99,6 @@ const std::unordered_map& function_names(bool rebuild = false); std::unordered_map& globals(bool rebuild = false); -std::unordered_map< - std::string, std::vector>>& kernargs(); hsa_executable_t load_executable(const std::string& file, hsa_executable_t executable, hsa_agent_t agent); diff --git a/projects/hip/src/program_state.cpp b/projects/hip/src/program_state.cpp index b490c0ee25..88cdeeb404 100644 --- a/projects/hip/src/program_state.cpp +++ b/projects/hip/src/program_state.cpp @@ -312,8 +312,8 @@ const unordered_map>& kernels(bool rebui void load_code_object_and_freeze_executable( const string& file, hsa_agent_t agent, - hsa_executable_t executable) { - // TODO: the following sequence is inefficient, should be refactored + hsa_executable_t + executable) { // TODO: the following sequence is inefficient, should be refactored // into a single load of the file and subsequent ELFIO // processing. static const auto cor_deleter = [](hsa_code_object_reader_t* p) { @@ -340,85 +340,6 @@ void load_code_object_and_freeze_executable( code_readers.push_back(move(tmp)); } } - -size_t parse_args( - const string& metadata, - size_t f, - size_t l, - vector>& size_align) { - if (f == l) return f; - - do { - static constexpr size_t size_sz{5}; - f = metadata.find("Size:", f) + size_sz; - - if (l <= f) return f; - - auto size = strtoul(&metadata[f], nullptr, 10); - - static constexpr size_t align_sz{6}; - f = metadata.find("Align:", f) + align_sz; - - char* l{}; - auto align = strtoul(&metadata[f], &l, 10); - - f += (l - &metadata[f]) + 1; - - size_align.emplace_back(size, align); - } while (true); -} - -void read_kernarg_metadata( - elfio& reader, - unordered_map>>& kernargs) -{ // TODO: this is inefficient. - auto it = find_section_if( - reader, [](const section* x) { return x->get_type() == SHT_NOTE; }); - - if (!it) return; - - const note_section_accessor acc{reader, it}; - for (decltype(acc.get_notes_num()) i = 0; i != acc.get_notes_num(); ++i) { - ELFIO::Elf_Word type{}; - string name{}; - void* desc{}; - Elf_Word desc_size{}; - - acc.get_note(i, type, name, desc, desc_size); - - if (name != "AMD") continue; // TODO: switch to using NT_AMD_AMDGPU_HSA_METADATA. - - string tmp{ - static_cast(desc), static_cast(desc) + desc_size}; - - auto dx = tmp.find("Kernels:"); - - if (dx == string::npos) continue; - - static constexpr decltype(tmp.size()) kernels_sz{8}; - dx += kernels_sz; - - do { - dx = tmp.find("Name:", dx); - - if (dx == string::npos) break; - - static constexpr decltype(tmp.size()) name_sz{5}; - dx = tmp.find_first_not_of(' ', dx + name_sz); - - auto fn = tmp.substr(dx, tmp.find('\n', dx) - dx); - dx += fn.size(); - - dx = tmp.find("Args:", dx); - - if (dx == string::npos) break; - - static constexpr decltype(tmp.size()) args_sz{5}; - dx = parse_args( - tmp, dx + args_sz, tmp.find("CodeProps", dx), kernargs[fn]); - } while (true); - } -} } // namespace namespace hip_impl { @@ -580,25 +501,6 @@ unordered_map& globals(bool rebuild) { return r; } -unordered_map>>& kernargs() { - static unordered_map>> r; - static once_flag f; - - call_once(f, []() { - for (auto&& blob : code_object_blobs()) { - stringstream tmp{std::string{ - blob.second.front().cbegin(), blob.second.front().cend()}}; - - elfio reader; - if (!reader.load(tmp)) continue; - - read_kernarg_metadata(reader, r); - } - }); - - return r; -} - hsa_executable_t load_executable(const string& file, hsa_executable_t executable, hsa_agent_t agent) { elfio reader; From 71005f60da510ab599461a36a8aaab61a6f08087 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Fri, 30 Nov 2018 15:33:57 +0300 Subject: [PATCH 07/10] [HIPIFY][SPARSE] Matrix Reorderings and Format Conversion Reference + cuSPARSE is supported up to CUDA 10.0 + cuSPARSE_08 test (CUDA 10.0) + update CUSPARSE_API_supported_by_HIP.md + lit: add a rule for CUDA 10.0 tests excluding [ROCm/hip commit: 116b9191f78b6a4d2d3b8c57238f9b1d255adb74] --- .../markdown/CUSPARSE_API_supported_by_HIP.md | 156 +++++++ .../src/CUDA2HIP_SPARSE_API_functions.cpp | 190 ++++++++ .../hipify-clang/cuSPARSE/cuSPARSE_08.cu | 413 ++++++++++++++++++ projects/hip/tests/hipify-clang/lit.cfg | 4 + 4 files changed, 763 insertions(+) create mode 100644 projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_08.cu diff --git a/projects/hip/docs/markdown/CUSPARSE_API_supported_by_HIP.md b/projects/hip/docs/markdown/CUSPARSE_API_supported_by_HIP.md index 0d0ded1406..f552246f11 100644 --- a/projects/hip/docs/markdown/CUSPARSE_API_supported_by_HIP.md +++ b/projects/hip/docs/markdown/CUSPARSE_API_supported_by_HIP.md @@ -300,6 +300,7 @@ |`cusparseScsrgeam` | | |`cusparseDcsrgeam` | | |`cusparseCcsrgeam` | | +|`cusparseZcsrgeam` | | |`cusparseScsrgeam2_bufferSizeExt` | | |`cusparseDcsrgeam2_bufferSizeExt` | | |`cusparseCcsrgeam2_bufferSizeExt` | | @@ -308,6 +309,7 @@ |`cusparseScsrgemm` | | |`cusparseDcsrgemm` | | |`cusparseCcsrgemm` | | +|`cusparseZcsrgemm` | | |`cusparseScsrgemm2_bufferSizeExt` | | |`cusparseDcsrgemm2_bufferSizeExt` | | |`cusparseCcsrgemm2_bufferSizeExt` | | @@ -460,3 +462,157 @@ |`cusparseDgpsvInterleavedBatch` | | |`cusparseCgpsvInterleavedBatch` | | |`cusparseZgpsvInterleavedBatch` | | + +## **8. cuSPARSE Matrix Reorderings Reference** + +| **CUDA** | **HIP** | +|-----------------------------------------------------------|-------------------------------------------------| +|`cusparseScsrcolor` | | +|`cusparseDcsrcolor` | | +|`cusparseCcsrcolor` | | +|`cusparseZcsrcolor` | | + +## **9. cuSPARSE Format Conversion Reference** + +| **CUDA** | **HIP** | +|-----------------------------------------------------------|-------------------------------------------------| +|`cusparseSbsr2csr` | | +|`cusparseDbsr2csr` | | +|`cusparseCbsr2csr` | | +|`cusparseZbsr2csr` | | +|`cusparseSgebsr2gebsc_bufferSize` | | +|`cusparseDgebsr2gebsc_bufferSize` | | +|`cusparseCgebsr2gebsc_bufferSize` | | +|`cusparseZgebsr2gebsc_bufferSize` | | +|`cusparseSgebsr2gebsc` | | +|`cusparseDgebsr2gebsc` | | +|`cusparseCgebsr2gebsc` | | +|`cusparseZgebsr2gebsc` | | +|`cusparseSgebsr2gebsr_bufferSize` | | +|`cusparseDgebsr2gebsr_bufferSize` | | +|`cusparseCgebsr2gebsr_bufferSize` | | +|`cusparseZgebsr2gebsr_bufferSize` | | +|`cusparseXgebsr2gebsrNnz` | | +|`cusparseSgebsr2gebsr` | | +|`cusparseDgebsr2gebsr` | | +|`cusparseCgebsr2gebsr` | | +|`cusparseZgebsr2gebsr` | | +|`cusparseSgebsr2csr` | | +|`cusparseDgebsr2csr` | | +|`cusparseCgebsr2csr` | | +|`cusparseZgebsr2csr` | | +|`cusparseScsr2gebsr_bufferSize` | | +|`cusparseDcsr2gebsr_bufferSize` | | +|`cusparseCcsr2gebsr_bufferSize` | | +|`cusparseZcsr2gebsr_bufferSize` | | +|`cusparseXcsr2gebsrNnz` | | +|`cusparseScsr2gebsr` | | +|`cusparseDcsr2gebsr` | | +|`cusparseCcsr2gebsr` | | +|`cusparseZcsr2gebsr` | | +|`cusparseXcoo2csr` |`hipsparseXcoo2csr` | +|`cusparseScsc2dense` | | +|`cusparseDcsc2dense` | | +|`cusparseCcsc2dense` | | +|`cusparseZcsc2dense` | | +|`cusparseScsc2hyb` | | +|`cusparseDcsc2hyb` | | +|`cusparseCcsc2hyb` | | +|`cusparseZcsc2hyb` | | +|`cusparseXcsr2bsrNnz` | | +|`cusparseScsr2bsr` | | +|`cusparseDcsr2bsr` | | +|`cusparseCcsr2bsr` | | +|`cusparseZcsr2bsr` | | +|`cusparseXcsr2coo` |`hipsparseXcsr2coo` | +|`cusparseScsr2csc` |`hipsparseScsr2csc` | +|`cusparseDcsr2csc` |`hipsparseDcsr2csc` | +|`cusparseCcsr2csc` | | +|`cusparseZcsr2csc` | | +|`cusparseCsr2cscEx` | | +|`cusparseScsr2dense` | | +|`cusparseDcsr2dense` | | +|`cusparseCcsr2dense` | | +|`cusparseZcsr2dense` | | +|`cusparseScsr2csr_compress` | | +|`cusparseDcsr2csr_compress` | | +|`cusparseCcsr2csr_compress` | | +|`cusparseZcsr2csr_compress` | | +|`cusparseScsr2hyb` |`hipsparseScsr2hyb` | +|`cusparseDcsr2hyb` |`hipsparseDcsr2hyb` | +|`cusparseCcsr2hyb` | | +|`cusparseZcsr2hyb` | | +|`cusparseSdense2csc` | | +|`cusparseDdense2csc` | | +|`cusparseCdense2csc` | | +|`cusparseZdense2csc` | | +|`cusparseSdense2csr` | | +|`cusparseDdense2csr` | | +|`cusparseCdense2csr` | | +|`cusparseZdense2csr` | | +|`cusparseSdense2hyb` | | +|`cusparseDdense2hyb` | | +|`cusparseCdense2hyb` | | +|`cusparseZdense2hyb` | | +|`cusparseShyb2csc` | | +|`cusparseDhyb2csc` | | +|`cusparseChyb2csc` | | +|`cusparseZhyb2csc` | | +|`cusparseShyb2csr` | | +|`cusparseDhyb2csr` | | +|`cusparseChyb2csr` | | +|`cusparseZhyb2csr` | | +|`cusparseShyb2dense` | | +|`cusparseDhyb2dense` | | +|`cusparseChyb2dense` | | +|`cusparseZhyb2dense` | | +|`cusparseSnnz` | | +|`cusparseDnnz` | | +|`cusparseCnnz` | | +|`cusparseZnnz` | | +|`cusparseCreateIdentityPermutation` |`hipsparseCreateIdentityPermutation` | +|`cusparseXcoosort_bufferSizeExt` |`hipsparseXcoosort_bufferSizeExt` | +|`cusparseXcoosortByRow` |`hipsparseXcoosortByRow` | +|`cusparseXcoosortByColumn` |`hipsparseXcoosortByColumn` | +|`cusparseXcsrsort_bufferSizeExt` |`hipsparseXcsrsort_bufferSizeExt` | +|`cusparseXcsrsort` |`hipsparseXcsrsort` | +|`cusparseScusparseXcscsort_bufferSizeExtnnz` | | +|`cusparseXcscsort` | | +|`cusparseCreateCsru2csrInfo` | | +|`cusparseDestroyCsru2csrInfo` | | +|`cusparseScsru2csr_bufferSizeExt` | | +|`cusparseDcsru2csr_bufferSizeExt` | | +|`cusparseCcsru2csr_bufferSizeExt` | | +|`cusparseZcsru2csr_bufferSizeExt` | | +|`cusparseScsru2csr` | | +|`cusparseDcsru2csr` | | +|`cusparseCcsru2csr` | | +|`cusparseZcsru2csr` | | +|`cusparseHpruneDense2csr_bufferSizeExt` | | +|`cusparseSpruneDense2csr_bufferSizeExt` | | +|`cusparseDpruneDense2csr_bufferSizeExt` | | +|`cusparseHpruneDense2csrNnz` | | +|`cusparseSpruneDense2csrNnz` | | +|`cusparseDpruneDense2csrNnz` | | +|`cusparseHpruneCsr2csr_bufferSizeExt` | | +|`cusparseSpruneCsr2csr_bufferSizeExt` | | +|`cusparseDpruneCsr2csr_bufferSizeExt` | | +|`cusparseHpruneCsr2csrNnz` | | +|`cusparseSpruneCsr2csrNnz` | | +|`cusparseDpruneCsr2csrNnz` | | +|`cusparseHpruneDense2csrByPercentage_bufferSizeExt` | | +|`cusparseSpruneDense2csrByPercentage_bufferSizeExt` | | +|`cusparseDpruneDense2csrByPercentage_bufferSizeExt` | | +|`cusparseHpruneDense2csrNnzByPercentage` | | +|`cusparseSpruneDense2csrNnzByPercentage` | | +|`cusparseDpruneDense2csrNnzByPercentage` | | +|`cusparseHpruneCsr2csrByPercentage_bufferSizeExt` | | +|`cusparseSpruneCsr2csrByPercentage_bufferSizeExt` | | +|`cusparseDpruneCsr2csrByPercentage_bufferSizeExt` | | +|`cusparseHpruneCsr2csrNnzByPercentage` | | +|`cusparseSpruneCsr2csrNnzByPercentage` | | +|`cusparseDpruneCsr2csrNnzByPercentage` | | +|`cusparseSnnz_compress` | | +|`cusparseDnnz_compress` | | +|`cusparseCnnz_compress` | | +|`cusparseZnnz_compress` | | diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp index fdb1a8381c..f3b0f0eb99 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp @@ -245,6 +245,7 @@ const std::map CUDA_SPARSE_FUNCTION_MAP{ {"cusparseScsrgeam", {"hipsparseScsrgeam", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseDcsrgeam", {"hipsparseDcsrgeam", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseCcsrgeam", {"hipsparseCcsrgeam", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZcsrgeam", {"hipsparseZcsrgeam", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseScsrgeam2_bufferSizeExt", {"hipsparseScsrgeam2_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseDcsrgeam2_bufferSizeExt", {"hipsparseDcsrgeam2_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, @@ -255,6 +256,7 @@ const std::map CUDA_SPARSE_FUNCTION_MAP{ {"cusparseScsrgemm", {"hipsparseScsrgemm", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseDcsrgemm", {"hipsparseDcsrgemm", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseCcsrgemm", {"hipsparseCcsrgemm", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZcsrgemm", {"hipsparseZcsrgemm", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseScsrgemm2_bufferSizeExt", {"hipsparseScsrgemm2_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseDcsrgemm2_bufferSizeExt", {"hipsparseDcsrgemm2_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, @@ -416,4 +418,192 @@ const std::map CUDA_SPARSE_FUNCTION_MAP{ {"cusparseDgpsvInterleavedBatch", {"hipsparseDgpsvInterleavedBatch", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseCgpsvInterleavedBatch", {"hipsparseCgpsvInterleavedBatch", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZgpsvInterleavedBatch", {"hipsparseZgpsvInterleavedBatch", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + // 11. cuSPARSE Matrix Reorderings Reference + {"cusparseScsrcolor", {"hipsparseScsrcolor", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDcsrcolor", {"hipsparseDcsrcolor", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsrcolor", {"hipsparseCcsrcolor", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZcsrcolor", {"hipsparseZcsrcolor", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + // 12. cuSPARSE Format Conversion Reference + {"cusparseSbsr2csr", {"hipsparseSbsr2csr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDbsr2csr", {"hipsparseDbsr2csr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCbsr2csr", {"hipsparseCbsr2csr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZbsr2csr", {"hipsparseZbsr2csr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseSgebsr2gebsc_bufferSize", {"hipsparseSgebsr2gebsc_bufferSize", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDgebsr2gebsc_bufferSize", {"hipsparseDgebsr2gebsc_bufferSize", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCgebsr2gebsc_bufferSize", {"hipsparseCgebsr2gebsc_bufferSize", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZgebsr2gebsc_bufferSize", {"hipsparseZgebsr2gebsc_bufferSize", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseSgebsr2gebsc", {"hipsparseSgebsr2gebsc", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDgebsr2gebsc", {"hipsparseDgebsr2gebsc", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCgebsr2gebsc", {"hipsparseCgebsr2gebsc", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZgebsr2gebsc", {"hipsparseZgebsr2gebsc", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseSgebsr2gebsr_bufferSize", {"hipsparseSgebsr2gebsr_bufferSize", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDgebsr2gebsr_bufferSize", {"hipsparseDgebsr2gebsr_bufferSize", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCgebsr2gebsr_bufferSize", {"hipsparseCgebsr2gebsr_bufferSize", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZgebsr2gebsr_bufferSize", {"hipsparseZgebsr2gebsr_bufferSize", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseSgebsr2csr", {"hipsparseSgebsr2csr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDgebsr2csr", {"hipsparseDgebsr2csr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCgebsr2csr", {"hipsparseCgebsr2csr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZgebsr2csr", {"hipsparseZgebsr2csr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseXgebsr2gebsrNnz", {"hipsparseXgebsr2gebsrNnz", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseSgebsr2gebsr", {"hipsparseSgebsr2gebsr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDgebsr2gebsr", {"hipsparseDgebsr2gebsr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCgebsr2gebsr", {"hipsparseCgebsr2gebsr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZgebsr2gebsr", {"hipsparseZgebsr2gebsr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseScsr2gebsr_bufferSize", {"hipsparseScsr2gebsr_bufferSize", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDcsr2gebsr_bufferSize", {"hipsparseDcsr2gebsr_bufferSize", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsr2gebsr_bufferSize", {"hipsparseCcsr2gebsr_bufferSize", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZcsr2gebsr_bufferSize", {"hipsparseZcsr2gebsr_bufferSize", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseXcsr2gebsrNnz", {"hipsparseXcsr2gebsrNnz", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseScsr2gebsr", {"hipsparseScsr2gebsr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDcsr2gebsr", {"hipsparseDcsr2gebsr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsr2gebsr", {"hipsparseCcsr2gebsr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZcsr2gebsr", {"hipsparseZcsr2gebsr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseXcoo2csr", {"hipsparseXcoo2csr", CONV_LIB_FUNC, API_SPARSE}}, + + {"cusparseScsc2dense", {"hipsparseScsc2dense", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDcsc2dense", {"hipsparseDcsc2dense", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsc2dense", {"hipsparseCcsc2dense", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZcsc2dense", {"hipsparseZcsc2dense", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseScsc2hyb", {"hipsparseScsc2hyb", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDcsc2hyb", {"hipsparseDcsc2hyb", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsc2hyb", {"hipsparseCcsc2hyb", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZcsc2hyb", {"hipsparseZcsc2hyb", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseXcsr2bsrNnz", {"hipsparseXcsr2bsrNnz", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseScsr2bsr", {"hipsparseScsr2bsr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDcsr2bsr", {"hipsparseDcsr2bsr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsr2bsr", {"hipsparseCcsr2bsr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZcsr2bsr", {"hipsparseZcsr2bsr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseXcsr2coo", {"hipsparseXcsr2coo", CONV_LIB_FUNC, API_SPARSE}}, + + {"cusparseScsr2csc", {"hipsparseScsr2csc", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseDcsr2csc", {"hipsparseDcsr2csc", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseCcsr2csc", {"hipsparseCcsr2csc", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZcsr2csc", {"hipsparseZcsr2csc", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseCsr2cscEx", {"hipsparseCsr2cscEx", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseScsr2dense", {"hipsparseScsr2dense", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDcsr2dense", {"hipsparseDcsr2dense", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsr2dense", {"hipsparseCcsr2dense", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZcsr2dense", {"hipsparseZcsr2dense", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseScsr2csr_compress", {"hipsparseScsr2csr_compress", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDcsr2csr_compress", {"hipsparseDcsr2csr_compress", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDcsr2csr_compress", {"hipsparseDcsr2csr_compress", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZcsr2csr_compress", {"hipsparseZcsr2csr_compress", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseScsr2hyb", {"hipsparseScsr2hyb", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseDcsr2hyb", {"hipsparseDcsr2hyb", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseCcsr2hyb", {"hipsparseCcsr2hyb", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZcsr2hyb", {"hipsparseZcsr2hyb", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseSdense2csc", {"hipsparseSdense2csc", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDdense2csc", {"hipsparseDdense2csc", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCdense2csc", {"hipsparseCdense2csc", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZdense2csc", {"hipsparseZdense2csc", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseSdense2csr", {"hipsparseSdense2csr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDdense2csr", {"hipsparseDdense2csr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCdense2csr", {"hipsparseCdense2csr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZdense2csr", {"hipsparseZdense2csr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseSdense2hyb", {"hipsparseSdense2hyb", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDdense2hyb", {"hipsparseDdense2hyb", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCdense2hyb", {"hipsparseCdense2hyb", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZdense2hyb", {"hipsparseZdense2hyb", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseShyb2csc", {"hipsparseShyb2csc", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDhyb2csc", {"hipsparseDhyb2csc", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseChyb2csc", {"hipsparseChyb2csc", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZhyb2csc", {"hipsparseZhyb2csc", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseShyb2csr", {"hipsparseShyb2csr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDhyb2csr", {"hipsparseDhyb2csr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseChyb2csr", {"hipsparseChyb2csr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZhyb2csr", {"hipsparseZhyb2csr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseShyb2dense", {"hipsparseShyb2dense", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDhyb2dense", {"hipsparseDhyb2dense", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseChyb2dense", {"hipsparseChyb2dense", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZhyb2dense", {"hipsparseZhyb2dense", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseSnnz", {"hipsparseSnnz", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDnnz", {"hipsparseDnnz", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCnnz", {"hipsparseCnnz", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZnnz", {"hipsparseZnnz", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseCreateIdentityPermutation", {"hipsparseCreateIdentityPermutation", CONV_LIB_FUNC, API_SPARSE}}, + + {"cusparseXcoosort_bufferSizeExt", {"hipsparseXcoosort_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseXcoosortByRow", {"hipsparseXcoosortByRow", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseXcoosortByColumn", {"hipsparseXcoosortByColumn", CONV_LIB_FUNC, API_SPARSE}}, + + {"cusparseXcsrsort_bufferSizeExt", {"hipsparseXcsrsort_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseXcsrsort", {"hipsparseXcsrsort", CONV_LIB_FUNC, API_SPARSE}}, + + {"cusparseScusparseXcscsort_bufferSizeExtnnz", {"hipsparseXcscsort_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseXcscsort", {"hipsparseXcscsort", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseCreateCsru2csrInfo", {"hipsparseCreateCsru2csrInfo", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDestroyCsru2csrInfo", {"hipsparseDestroyCsru2csrInfo", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseScsru2csr_bufferSizeExt", {"hipsparseScsru2csr_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDcsru2csr_bufferSizeExt", {"hipsparseDcsru2csr_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsru2csr_bufferSizeExt", {"hipsparseCcsru2csr_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZcsru2csr_bufferSizeExt", {"hipsparseZcsru2csr_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseScsru2csr", {"hipsparseScsru2csr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDcsru2csr", {"hipsparseDcsru2csr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsru2csr", {"hipsparseCcsru2csr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZcsru2csr", {"hipsparseZcsru2csr", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseHpruneDense2csr_bufferSizeExt", {"hipsparseHpruneDense2csr_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseSpruneDense2csr_bufferSizeExt", {"hipsparseSpruneDense2csr_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDpruneDense2csr_bufferSizeExt", {"hipsparseDpruneDense2csr_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseHpruneDense2csrNnz", {"hipsparseHpruneDense2csrNnz", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseSpruneDense2csrNnz", {"hipsparseSpruneDense2csrNnz", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDpruneDense2csrNnz", {"hipsparseDpruneDense2csrNnz", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseHpruneCsr2csr_bufferSizeExt", {"hipsparseHpruneCsr2csr_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseSpruneCsr2csr_bufferSizeExt", {"hipsparseSpruneCsr2csr_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDpruneCsr2csr_bufferSizeExt", {"hipsparseDpruneCsr2csr_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseHpruneCsr2csrNnz", {"hipsparseHpruneCsr2csrNnz", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseSpruneCsr2csrNnz", {"hipsparseSpruneCsr2csrNnz", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDpruneCsr2csrNnz", {"hipsparseDpruneCsr2csrNnz", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseHpruneDense2csrByPercentage_bufferSizeExt", {"hipsparseHpruneDense2csrByPercentage_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseSpruneDense2csrByPercentage_bufferSizeExt", {"hipsparseSpruneDense2csrByPercentage_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDpruneDense2csrByPercentage_bufferSizeExt", {"hipsparseDpruneDense2csrByPercentage_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseHpruneDense2csrNnzByPercentage", {"hipsparseHpruneDense2csrNnzByPercentage", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseSpruneDense2csrNnzByPercentage", {"hipsparseSpruneDense2csrNnzByPercentage", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDpruneDense2csrNnzByPercentage", {"hipsparseDpruneDense2csrNnzByPercentage", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseHpruneCsr2csrByPercentage_bufferSizeExt", {"hipsparseHpruneCsr2csrByPercentage_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseSpruneCsr2csrByPercentage_bufferSizeExt", {"hipsparseSpruneCsr2csrByPercentage_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDpruneCsr2csrByPercentage_bufferSizeExt", {"hipsparseDpruneCsr2csrByPercentage_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseHpruneCsr2csrNnzByPercentage", {"hipsparseHpruneCsr2csrNnzByPercentage", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseSpruneCsr2csrNnzByPercentage", {"hipsparseSpruneCsr2csrNnzByPercentage", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDpruneCsr2csrNnzByPercentage", {"hipsparseDpruneCsr2csrNnzByPercentage", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseSnnz_compress", {"hipsparseSnnz_compress", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDnnz_compress", {"hipsparseDnnz_compress", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCnnz_compress", {"hipsparseCnnz_compress", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZnnz_compress", {"hipsparseZnnz_compress", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, }; diff --git a/projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_08.cu b/projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_08.cu new file mode 100644 index 0000000000..fcfde8d3b2 --- /dev/null +++ b/projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_08.cu @@ -0,0 +1,413 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args +#include +#include +#include +// CHECK: #include +#include +// CHECK: #include +#include +// CHECK: #include +#include + +// NOTE: CUDA 10.0 + +/* + * compute | b - A*x|_inf + */ +void residaul_eval( + int n, + const float *dl, + const float *d, + const float *du, + const float *b, + const float *x, + float *r_nrminf_ptr) +{ + float r_nrminf = 0; + for (int i = 0; i < n; i++) { + float dot = 0; + if (i > 0) { + dot += dl[i] * x[i - 1]; + } + dot += d[i] * x[i]; + if (i < (n - 1)) { + dot += du[i] * x[i + 1]; + } + float ri = b[i] - dot; + r_nrminf = (r_nrminf > fabs(ri)) ? r_nrminf : fabs(ri); + } + + *r_nrminf_ptr = r_nrminf; +} + +int main(int argc, char*argv[]) +{ + // CHECK: hipsparseHandle_t cusparseH = NULL; + cusparseHandle_t cusparseH = NULL; + // CHECK: hipblasHandle_t cublasH = NULL; + cublasHandle_t cublasH = NULL; + // CHECK: hipStream_t stream = NULL; + cudaStream_t stream = NULL; + // CHECK: hipsparseStatus_t status = HIPSPARSE_STATUS_SUCCESS; + cusparseStatus_t status = CUSPARSE_STATUS_SUCCESS; + // CHECK: hipblasStatus_t cublasStat = HIPBLAS_STATUS_SUCCESS; + cublasStatus_t cublasStat = CUBLAS_STATUS_SUCCESS; + // CHECK: hipError_t cudaStat1 = hipSuccess; + cudaError_t cudaStat1 = cudaSuccess; + + const int n = 3; + const int batchSize = 2; + /* + * | 1 6 0 | | 1 | | -0.603960 | + * A1 =| 4 2 7 |, b1 = | 2 |, x1 = | 0.267327 | + * | 0 5 3 | | 3 | | 0.554455 | + * + * | 8 13 0 | | 4 | | -0.063291 | + * A2 =| 11 9 14 |, b2 = | 5 |, x2 = | 0.346641 | + * | 0 12 10 | | 6 | | 0.184031 | + */ + + /* + * A = (dl, d, du), B and X are in aggregate format + */ + const float dl[n * batchSize] = { 0, 4, 5, 0, 11, 12 }; + const float d[n * batchSize] = { 1, 2, 3, 8, 9, 10 }; + const float du[n * batchSize] = { 6, 7, 0, 13, 14, 0 }; + const float B[n * batchSize] = { 1, 2, 3, 4, 5, 6 }; + float X[n * batchSize]; /* Xj = Aj \ Bj */ + +/* device memory + * (d_dl0, d_d0, d_du0) is aggregate format + * (d_dl, d_d, d_du) is interleaved format + */ + float *d_dl0 = NULL; + float *d_d0 = NULL; + float *d_du0 = NULL; + float *d_dl = NULL; + float *d_d = NULL; + float *d_du = NULL; + float *d_B = NULL; + float *d_X = NULL; + + size_t lworkInBytes = 0; + char *d_work = NULL; + + /* + * algo = 0: cuThomas (unstable) + * algo = 1: LU with pivoting (stable) + * algo = 2: QR (stable) + */ + const int algo = 2; + + const float h_one = 1; + const float h_zero = 0; + + printf("example of gtsv (interleaved format) \n"); + printf("choose algo = 0,1,2 to select different algorithms \n"); + printf("n = %d, batchSize = %d, algo = %d \n", n, batchSize, algo); + + /* step 1: create cusparse/cublas handle, bind a stream */ + // CHECK: cudaStat1 = hipStreamCreateWithFlags(&stream, hipStreamNonBlocking); + cudaStat1 = cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: status = hipsparseCreate(&cusparseH); + status = cusparseCreate(&cusparseH); + //CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); + assert(CUSPARSE_STATUS_SUCCESS == status); + // CHECK: status = hipsparseSetStream(cusparseH, stream); + status = cusparseSetStream(cusparseH, stream); + //CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); + assert(CUSPARSE_STATUS_SUCCESS == status); + // CHECK: cublasStat = hipblasCreate(&cublasH); + cublasStat = cublasCreate(&cublasH); + // CHECK: assert(HIPBLAS_STATUS_SUCCESS == cublasStat); + assert(CUBLAS_STATUS_SUCCESS == cublasStat); + // CHECK: cublasStat = hipblasSetStream(cublasH, stream); + cublasStat = cublasSetStream(cublasH, stream); + // CHECK: assert(HIPBLAS_STATUS_SUCCESS == cublasStat); + assert(CUBLAS_STATUS_SUCCESS == cublasStat); + + /* step 2: allocate device memory */ + // CHECK: cudaStat1 = hipMalloc((void**)&d_dl0, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_dl0, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_d0, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_d0, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_du0, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_du0, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_dl, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_dl, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_d, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_d, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_du, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_du, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_B, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_B, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_X, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_X, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + + /* step 3: prepare data in device, interleaved format */ + // CHECK: cudaStat1 = hipMemcpy(d_dl0, dl, sizeof(float)*n*batchSize, hipMemcpyHostToDevice); + cudaStat1 = cudaMemcpy(d_dl0, dl, sizeof(float)*n*batchSize, cudaMemcpyHostToDevice); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMemcpy(d_d0, d, sizeof(float)*n*batchSize, hipMemcpyHostToDevice); + cudaStat1 = cudaMemcpy(d_d0, d, sizeof(float)*n*batchSize, cudaMemcpyHostToDevice); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMemcpy(d_du0, du, sizeof(float)*n*batchSize, hipMemcpyHostToDevice); + cudaStat1 = cudaMemcpy(d_du0, du, sizeof(float)*n*batchSize, cudaMemcpyHostToDevice); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMemcpy(d_B, B, sizeof(float)*n*batchSize, hipMemcpyHostToDevice); + cudaStat1 = cudaMemcpy(d_B, B, sizeof(float)*n*batchSize, cudaMemcpyHostToDevice); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: hipDeviceSynchronize(); + cudaDeviceSynchronize(); + /* convert dl to interleaved format + * dl = transpose(dl0) + */ + // CHECK: cublasStat = hipblasSgeam( + // CHECK: HIPBLAS_OP_T, + // CHECK: HIPBLAS_OP_T, + cublasStat = cublasSgeam( + cublasH, + CUBLAS_OP_T, /* transa */ + CUBLAS_OP_T, /* transb, don't care */ + batchSize, /* number of rows of dl */ + n, /* number of columns of dl */ + &h_one, + d_dl0, /* dl0 is n-by-batchSize */ + n, /* leading dimension of dl0 */ + &h_zero, + NULL, + n, /* don't cae */ + d_dl, /* dl is batchSize-by-n */ + batchSize /* leading dimension of dl */ + ); + // CHECK: assert(HIPBLAS_STATUS_SUCCESS == cublasStat); + assert(CUBLAS_STATUS_SUCCESS == cublasStat); + /* convert d to interleaved format + * d = transpose(d0) + */ + // CHECK: cublasStat = hipblasSgeam( + // CHECK: HIPBLAS_OP_T, + // CHECK: HIPBLAS_OP_T, + cublasStat = cublasSgeam( + cublasH, + CUBLAS_OP_T, /* transa */ + CUBLAS_OP_T, /* transb, don't care */ + batchSize, /* number of rows of d */ + n, /* number of columns of d */ + &h_one, + d_d0, /* d0 is n-by-batchSize */ + n, /* leading dimension of d0 */ + &h_zero, + NULL, + n, /* don't cae */ + d_d, /* d is batchSize-by-n */ + batchSize /* leading dimension of d */ + ); + // CHECK: assert(HIPBLAS_STATUS_SUCCESS == cublasStat); + assert(CUBLAS_STATUS_SUCCESS == cublasStat); + + /* convert du to interleaved format + * du = transpose(du0) + */ + // CHECK: cublasStat = hipblasSgeam( + // CHECK: HIPBLAS_OP_T, + // CHECK: HIPBLAS_OP_T, + cublasStat = cublasSgeam( + cublasH, + CUBLAS_OP_T, /* transa */ + CUBLAS_OP_T, /* transb, don't care */ + batchSize, /* number of rows of du */ + n, /* number of columns of du */ + &h_one, + d_du0, /* du0 is n-by-batchSize */ + n, /* leading dimension of du0 */ + &h_zero, + NULL, + n, /* don't cae */ + d_du, /* du is batchSize-by-n */ + batchSize /* leading dimension of du */ + ); + // CHECK: assert(HIPBLAS_STATUS_SUCCESS == cublasStat); + assert(CUBLAS_STATUS_SUCCESS == cublasStat); + + /* convert B to interleaved format + * X = transpose(B) + */ + // CHECK: cublasStat = hipblasSgeam( + // CHECK: HIPBLAS_OP_T, + // CHECK: HIPBLAS_OP_T, + cublasStat = cublasSgeam( + cublasH, + CUBLAS_OP_T, /* transa */ + CUBLAS_OP_T, /* transb, don't care */ + batchSize, /* number of rows of X */ + n, /* number of columns of X */ + &h_one, + d_B, /* B is n-by-batchSize */ + n, /* leading dimension of B */ + &h_zero, + NULL, + n, /* don't cae */ + d_X, /* X is batchSize-by-n */ + batchSize /* leading dimension of X */ + ); + // CHECK: assert(HIPBLAS_STATUS_SUCCESS == cublasStat); + assert(CUBLAS_STATUS_SUCCESS == cublasStat); + /* step 4: prepare workspace */ + // NOTE: CUDA 10.0 + // CHECK: status = hipsparseSgtsvInterleavedBatch_bufferSizeExt( + status = cusparseSgtsvInterleavedBatch_bufferSizeExt( + cusparseH, + algo, + n, + d_dl, + d_d, + d_du, + d_X, + batchSize, + &lworkInBytes); + // CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); + assert(CUSPARSE_STATUS_SUCCESS == status); + + printf("lworkInBytes = %lld \n", (long long)lworkInBytes); + // CHECK: cudaStat1 = hipMalloc((void**)&d_work, lworkInBytes); + cudaStat1 = cudaMalloc((void**)&d_work, lworkInBytes); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + + /* step 5: solve Aj*xj = bj */ + // NOTE: CUDA 10.0 + // CHECK: status = hipsparseSgtsvInterleavedBatch( + status = cusparseSgtsvInterleavedBatch( + cusparseH, + algo, + n, + d_dl, + d_d, + d_du, + d_X, + batchSize, + d_work); + // CHECK: cudaStat1 = hipDeviceSynchronize(); + cudaStat1 = cudaDeviceSynchronize(); + // CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); + assert(CUSPARSE_STATUS_SUCCESS == status); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + + /* step 6: convert X back to aggregate format */ + /* B = transpose(X) */ + // CHECK: cublasStat = hipblasSgeam( + // CHECK: HIPBLAS_OP_T, + // CHECK: HIPBLAS_OP_T, + cublasStat = cublasSgeam( + cublasH, + CUBLAS_OP_T, /* transa */ + CUBLAS_OP_T, /* transb, don't care */ + n, /* number of rows of B */ + batchSize, /* number of columns of B */ + &h_one, + d_X, /* X is batchSize-by-n */ + batchSize, /* leading dimension of X */ + &h_zero, + NULL, + n, /* don't cae */ + d_B, /* B is n-by-batchSize */ + n /* leading dimension of B */ + ); + // CHECK: assert(HIPBLAS_STATUS_SUCCESS == cublasStat); + assert(CUBLAS_STATUS_SUCCESS == cublasStat); + // CHECK: hipDeviceSynchronize(); + cudaDeviceSynchronize(); + + /* step 7: residual evaluation */ + // CHECK: cudaStat1 = hipMemcpy(X, d_B, sizeof(float)*n*batchSize, hipMemcpyDeviceToHost); + cudaStat1 = cudaMemcpy(X, d_B, sizeof(float)*n*batchSize, cudaMemcpyDeviceToHost); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: hipDeviceSynchronize(); + cudaDeviceSynchronize(); + printf("==== x1 = inv(A1)*b1 \n"); + for (int j = 0; j < n; j++) { + printf("x1[%d] = %f\n", j, X[j]); + } + + float r1_nrminf; + residaul_eval( + n, + dl, + d, + du, + B, + X, + &r1_nrminf + ); + printf("|b1 - A1*x1| = %E\n", r1_nrminf); + + printf("\n==== x2 = inv(A2)*b2 \n"); + for (int j = 0; j < n; j++) { + printf("x2[%d] = %f\n", j, X[n + j]); + } + + float r2_nrminf; + residaul_eval( + n, + dl + n, + d + n, + du + n, + B + n, + X + n, + &r2_nrminf + ); + printf("|b2 - A2*x2| = %E\n", r2_nrminf); + + /* free resources */ + // CHECK: if (d_dl0) hipFree(d_dl0); + if (d_dl0) cudaFree(d_dl0); + // CHECK: if (d_d0) hipFree(d_d0); + if (d_d0) cudaFree(d_d0); + // CHECK: if (d_du0) hipFree(d_du0); + if (d_du0) cudaFree(d_du0); + // CHECK: if (d_dl) hipFree(d_dl); + if (d_dl) cudaFree(d_dl); + // CHECK: if (d_d) hipFree(d_d); + if (d_d) cudaFree(d_d); + // CHECK: if (d_du) hipFree(d_du); + if (d_du) cudaFree(d_du); + // CHECK: if (d_B) hipFree(d_B); + if (d_B) cudaFree(d_B); + // CHECK: if (d_X) hipFree(d_X); + if (d_X) cudaFree(d_X); + // CHECK: if (cusparseH) hipsparseDestroy(cusparseH); + if (cusparseH) cusparseDestroy(cusparseH); + // CHECK: if (cublasH) hipblasDestroy(cublasH); + if (cublasH) cublasDestroy(cublasH); + // CHECK: if (stream) hipStreamDestroy(stream); + if (stream) cudaStreamDestroy(stream); + // CHECK: hipDeviceReset(); + cudaDeviceReset(); + + return 0; +} diff --git a/projects/hip/tests/hipify-clang/lit.cfg b/projects/hip/tests/hipify-clang/lit.cfg index 98961fe166..b9ff0b2495 100644 --- a/projects/hip/tests/hipify-clang/lit.cfg +++ b/projects/hip/tests/hipify-clang/lit.cfg @@ -23,6 +23,10 @@ config.test_source_root = os.path.dirname(__file__) config.excludes = ['cmdparser.hpp'] +config.cuda_version = "@CUDA_VERSION@" +if config.cuda_version not in ['10.0']: + config.excludes.append('cuSPARSE_08.cu') + # test_exec_root: The path where tests are located (default is the test suite root). #config.test_exec_root = config.test_source_root From 98a6a2916297611f012af92b422950fc7bfaec10 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 3 Dec 2018 08:54:13 +0530 Subject: [PATCH 08/10] [ci] Add rocm-2.0.x to CI test infrastructure Change-Id: I7fc0c40d1bf50a90ce3e210f2c8e83d1f4bf6d5c [ROCm/hip commit: dc5b0a1fe00dfcf2814976d34b72c2efaa0783e4] --- projects/hip/Jenkinsfile | 45 ++++++++++++++++++++++++++++++++++++++++ 1 file changed, 45 insertions(+) diff --git a/projects/hip/Jenkinsfile b/projects/hip/Jenkinsfile index e6b60f398e..1257bfb0b4 100644 --- a/projects/hip/Jenkinsfile +++ b/projects/hip/Jenkinsfile @@ -340,6 +340,51 @@ parallel rocm_1_9: */ } }, +rocm_2_0: +{ + node('hip-rocm') + { + String hcc_ver = 'rocm-2.0.x' + String from_image = 'ci_test_nodes/rocm-2.0.x/ubuntu-16.04:latest' + String inside_args = '--device=/dev/kfd --device=/dev/dri --group-add=video' + + // Checkout source code, dependencies and version files + String source_hip_rel = checkout_and_version( hcc_ver ) + + // Create/reuse a docker image that represents the hip build environment + def hip_build_image = docker_build_image( hcc_ver, 'hip', '', source_hip_rel, from_image ) + + // Print system information for the log + hip_build_image.inside( inside_args ) + { + sh """#!/usr/bin/env bash + set -x + /opt/rocm/bin/rocm_agent_enumerator -t ALL + /opt/rocm/bin/hcc --version + """ + } + + // Conctruct a binary directory path based on build config + String build_hip_rel = build_directory_rel( build_config ); + + // Build hip inside of the build environment + docker_build_inside_image( hip_build_image, inside_args, hcc_ver, '', build_config, source_hip_rel, build_hip_rel ) + + // Clean docker build image + docker_clean_images( 'hip', docker_build_image_name( ) ) + + // After a successful build, upload a docker image of the results + /* + String hip_image_name = docker_upload_artifactory( hcc_ver, job_name, from_image, source_hip_rel, build_hip_rel ) + if( params.push_image_to_docker_hub ) + { + docker_upload_dockerhub( job_name, hip_image_name, 'rocm' ) + docker_clean_images( 'rocm', hip_image_name ) + } + docker_clean_images( job_name, hip_image_name ) + */ + } +}, rocm_head: { node('hip-rocm') From fd03a77da0d0b3cd0326357d0a515ff96225ad07 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 4 Dec 2018 19:24:29 +0300 Subject: [PATCH 09/10] [HIPIFY][SPARSE] Add 3 more CUDA 10.0 tests + lit update + fix typos [ROCm/hip commit: 72d40db35893c5c9d1346a8ceab606cfd473ba5b] --- .../markdown/CUSPARSE_API_supported_by_HIP.md | 2 +- .../src/CUDA2HIP_SPARSE_API_functions.cpp | 2 +- .../hipify-clang/cuSPARSE/cuSPARSE_08.cu | 6 +- .../hipify-clang/cuSPARSE/cuSPARSE_09.cu | 414 ++++++++++++++ .../hipify-clang/cuSPARSE/cuSPARSE_10.cu | 507 ++++++++++++++++++ .../hipify-clang/cuSPARSE/cuSPARSE_11.cu | 327 +++++++++++ projects/hip/tests/hipify-clang/lit.cfg | 3 + 7 files changed, 1256 insertions(+), 5 deletions(-) create mode 100644 projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_09.cu create mode 100644 projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_10.cu create mode 100644 projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_11.cu diff --git a/projects/hip/docs/markdown/CUSPARSE_API_supported_by_HIP.md b/projects/hip/docs/markdown/CUSPARSE_API_supported_by_HIP.md index f552246f11..4dadea2c93 100644 --- a/projects/hip/docs/markdown/CUSPARSE_API_supported_by_HIP.md +++ b/projects/hip/docs/markdown/CUSPARSE_API_supported_by_HIP.md @@ -576,7 +576,7 @@ |`cusparseXcoosortByColumn` |`hipsparseXcoosortByColumn` | |`cusparseXcsrsort_bufferSizeExt` |`hipsparseXcsrsort_bufferSizeExt` | |`cusparseXcsrsort` |`hipsparseXcsrsort` | -|`cusparseScusparseXcscsort_bufferSizeExtnnz` | | +|`cusparseXcscsort_bufferSizeExt` | | |`cusparseXcscsort` | | |`cusparseCreateCsru2csrInfo` | | |`cusparseDestroyCsru2csrInfo` | | diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp index f3b0f0eb99..29726a5048 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp @@ -554,7 +554,7 @@ const std::map CUDA_SPARSE_FUNCTION_MAP{ {"cusparseXcsrsort_bufferSizeExt", {"hipsparseXcsrsort_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseXcsrsort", {"hipsparseXcsrsort", CONV_LIB_FUNC, API_SPARSE}}, - {"cusparseScusparseXcscsort_bufferSizeExtnnz", {"hipsparseXcscsort_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseXcscsort_bufferSizeExt", {"hipsparseXcscsort_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseXcscsort", {"hipsparseXcscsort", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseCreateCsru2csrInfo", {"hipsparseCreateCsru2csrInfo", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, diff --git a/projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_08.cu b/projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_08.cu index fcfde8d3b2..cf788c8f33 100644 --- a/projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_08.cu +++ b/projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_08.cu @@ -277,8 +277,8 @@ int main(int argc, char*argv[]) assert(CUBLAS_STATUS_SUCCESS == cublasStat); /* step 4: prepare workspace */ // NOTE: CUDA 10.0 - // CHECK: status = hipsparseSgtsvInterleavedBatch_bufferSizeExt( - status = cusparseSgtsvInterleavedBatch_bufferSizeExt( + // TODO: status = hipsparseSgtsvInterleavedBatch_bufferSizeExt( + status = cusparseSgtsvInterleavedBatch_bufferSizeExt( cusparseH, algo, n, @@ -299,7 +299,7 @@ int main(int argc, char*argv[]) /* step 5: solve Aj*xj = bj */ // NOTE: CUDA 10.0 - // CHECK: status = hipsparseSgtsvInterleavedBatch( + // TODO: status = hipsparseSgtsvInterleavedBatch( status = cusparseSgtsvInterleavedBatch( cusparseH, algo, diff --git a/projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_09.cu b/projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_09.cu new file mode 100644 index 0000000000..9b1aaf623f --- /dev/null +++ b/projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_09.cu @@ -0,0 +1,414 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args +#include +#include +#include +// CHECK: #include +#include +// CHECK: #include +#include +// CHECK: #include +#include + +// NOTE: CUDA 10.0 + +/* + * compute | b - A*x|_inf + */ +void residaul_eval( + int n, + const float *dl, + const float *d, + const float *du, + const float *b, + const float *x, + float *r_nrminf_ptr) +{ + float r_nrminf = 0; + for (int i = 0; i < n; i++) { + float dot = 0; + if (i > 0) { + dot += dl[i] * x[i - 1]; + } + dot += d[i] * x[i]; + if (i < (n - 1)) { + dot += du[i] * x[i + 1]; + } + float ri = b[i] - dot; + r_nrminf = (r_nrminf > fabs(ri)) ? r_nrminf : fabs(ri); + } + + *r_nrminf_ptr = r_nrminf; +} + +int main(int argc, char*argv[]) +{ + // CHECK: hipsparseHandle_t cusparseH = NULL; + cusparseHandle_t cusparseH = NULL; + // CHECK: hipblasHandle_t cublasH = NULL; + cublasHandle_t cublasH = NULL; + // CHECK: hipStream_t stream = NULL; + cudaStream_t stream = NULL; + // CHECK: hipsparseStatus_t status = HIPSPARSE_STATUS_SUCCESS; + cusparseStatus_t status = CUSPARSE_STATUS_SUCCESS; + // CHECK: hipblasStatus_t cublasStat = HIPBLAS_STATUS_SUCCESS; + cublasStatus_t cublasStat = CUBLAS_STATUS_SUCCESS; + // CHECK: hipError_t cudaStat1 = hipSuccess; + cudaError_t cudaStat1 = cudaSuccess; + + const int n = 3; + const int batchSize = 2; + /* + * | 1 6 0 | | 1 | | -0.603960 | + * A1 =| 4 2 7 |, b1 = | 2 |, x1 = | 0.267327 | + * | 0 5 3 | | 3 | | 0.554455 | + * + * | 8 13 0 | | 4 | | -0.063291 | + * A2 =| 11 9 14 |, b2 = | 5 |, x2 = | 0.346641 | + * | 0 12 10 | | 6 | | 0.184031 | + */ + + /* + * A = (dl, d, du), B and X are in aggregate format + */ + const float dl[n * batchSize] = { 0, 4, 5, 0, 11, 12 }; + const float d[n * batchSize] = { 1, 2, 3, 8, 9, 10 }; + const float du[n * batchSize] = { 6, 7, 0, 13, 14, 0 }; + const float B[n * batchSize] = { 1, 2, 3, 4, 5, 6 }; + float X[n * batchSize]; /* Xj = Aj \ Bj */ + +/* device memory + * (d_dl0, d_d0, d_du0) is aggregate format + * (d_dl, d_d, d_du) is interleaved format + */ + float *d_dl0 = NULL; + float *d_d0 = NULL; + float *d_du0 = NULL; + float *d_dl = NULL; + float *d_d = NULL; + float *d_du = NULL; + float *d_B = NULL; + float *d_X = NULL; + + size_t lworkInBytes = 0; + char *d_work = NULL; + + /* + * algo = 0: cuThomas (unstable) + * algo = 1: LU with pivoting (stable) + * algo = 2: QR (stable) + */ + const int algo = 2; + + const float h_one = 1; + const float h_zero = 0; + + printf("example of gtsv (interleaved format) \n"); + printf("choose algo = 0,1,2 to select different algorithms \n"); + printf("n = %d, batchSize = %d, algo = %d \n", n, batchSize, algo); + + /* step 1: create cusparse/cublas handle, bind a stream */ + // CHECK: cudaStat1 = hipStreamCreateWithFlags(&stream, hipStreamNonBlocking); + cudaStat1 = cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: status = hipsparseCreate(&cusparseH); + status = cusparseCreate(&cusparseH); + // CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); + assert(CUSPARSE_STATUS_SUCCESS == status); + // CHECK: status = hipsparseSetStream(cusparseH, stream); + status = cusparseSetStream(cusparseH, stream); + // CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); + assert(CUSPARSE_STATUS_SUCCESS == status); + // CHECK: cublasStat = hipblasCreate(&cublasH); + cublasStat = cublasCreate(&cublasH); + // CHECK: assert(HIPBLAS_STATUS_SUCCESS == cublasStat); + assert(CUBLAS_STATUS_SUCCESS == cublasStat); + // CHECK: cublasStat = hipblasSetStream(cublasH, stream); + cublasStat = cublasSetStream(cublasH, stream); + // CHECK: assert(HIPBLAS_STATUS_SUCCESS == cublasStat); + assert(CUBLAS_STATUS_SUCCESS == cublasStat); + + /* step 2: allocate device memory */ + // CHECK: cudaStat1 = hipMalloc((void**)&d_dl0, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_dl0, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_d0, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_d0, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_du0, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_du0, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_dl, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_dl, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_d, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_d, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_du, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_du, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_B, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_B, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_X, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_X, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + + /* step 3: prepare data in device, interleaved format */ + // CHECK: cudaStat1 = hipMemcpy(d_dl0, dl, sizeof(float)*n*batchSize, hipMemcpyHostToDevice); + cudaStat1 = cudaMemcpy(d_dl0, dl, sizeof(float)*n*batchSize, cudaMemcpyHostToDevice); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMemcpy(d_d0, d, sizeof(float)*n*batchSize, hipMemcpyHostToDevice); + cudaStat1 = cudaMemcpy(d_d0, d, sizeof(float)*n*batchSize, cudaMemcpyHostToDevice); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMemcpy(d_du0, du, sizeof(float)*n*batchSize, hipMemcpyHostToDevice); + cudaStat1 = cudaMemcpy(d_du0, du, sizeof(float)*n*batchSize, cudaMemcpyHostToDevice); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMemcpy(d_B, B, sizeof(float)*n*batchSize, hipMemcpyHostToDevice); + cudaStat1 = cudaMemcpy(d_B, B, sizeof(float)*n*batchSize, cudaMemcpyHostToDevice); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: hipDeviceSynchronize(); + cudaDeviceSynchronize(); + /* convert dl to interleaved format + * dl = transpose(dl0) + */ + // CHECK: cublasStat = hipblasSgeam( + // CHECK: HIPBLAS_OP_T, + // CHECK: HIPBLAS_OP_T, + cublasStat = cublasSgeam( + cublasH, + CUBLAS_OP_T, /* transa */ + CUBLAS_OP_T, /* transb, don't care */ + batchSize, /* number of rows of dl */ + n, /* number of columns of dl */ + &h_one, + d_dl0, /* dl0 is n-by-batchSize */ + n, /* leading dimension of dl0 */ + &h_zero, + NULL, + n, /* don't care */ + d_dl, /* dl is batchSize-by-n */ + batchSize /* leading dimension of dl */ + ); + // CHECK: assert(HIPBLAS_STATUS_SUCCESS == cublasStat); + assert(CUBLAS_STATUS_SUCCESS == cublasStat); + /* convert d to interleaved format + * d = transpose(d0) + */ + // CHECK: cublasStat = hipblasSgeam( + // CHECK: HIPBLAS_OP_T + // CHECK: HIPBLAS_OP_T + cublasStat = cublasSgeam( + cublasH, + CUBLAS_OP_T, /* transa */ + CUBLAS_OP_T, /* transb, don't care */ + batchSize, /* number of rows of d */ + n, /* number of columns of d */ + &h_one, + d_d0, /* d0 is n-by-batchSize */ + n, /* leading dimension of d0 */ + &h_zero, + NULL, + n, /* don't cae */ + d_d, /* d is batchSize-by-n */ + batchSize /* leading dimension of d */ + ); + // CHECK: assert(HIPBLAS_STATUS_SUCCESS == cublasStat); + assert(CUBLAS_STATUS_SUCCESS == cublasStat); + + /* convert du to interleaved format + * du = transpose(du0) + */ + // CHECK: cublasStat = hipblasSgeam( + // CHECK: HIPBLAS_OP_T + // CHECK: HIPBLAS_OP_T + cublasStat = cublasSgeam( + cublasH, + CUBLAS_OP_T, /* transa */ + CUBLAS_OP_T, /* transb, don't care */ + batchSize, /* number of rows of du */ + n, /* number of columns of du */ + &h_one, + d_du0, /* du0 is n-by-batchSize */ + n, /* leading dimension of du0 */ + &h_zero, + NULL, + n, /* don't cae */ + d_du, /* du is batchSize-by-n */ + batchSize /* leading dimension of du */ + ); + // CHECK: assert(HIPBLAS_STATUS_SUCCESS == cublasStat); + assert(CUBLAS_STATUS_SUCCESS == cublasStat); + + /* convert B to interleaved format + * X = transpose(B) + */ + // CHECK: cublasStat = hipblasSgeam( + // CHECK: HIPBLAS_OP_T + // CHECK: HIPBLAS_OP_T + cublasStat = cublasSgeam( + cublasH, + CUBLAS_OP_T, /* transa */ + CUBLAS_OP_T, /* transb, don't care */ + batchSize, /* number of rows of X */ + n, /* number of columns of X */ + &h_one, + d_B, /* B is n-by-batchSize */ + n, /* leading dimension of B */ + &h_zero, + NULL, + n, /* don't cae */ + d_X, /* X is batchSize-by-n */ + batchSize /* leading dimension of X */ + ); + // CHECK: assert(HIPBLAS_STATUS_SUCCESS == cublasStat); + assert(CUBLAS_STATUS_SUCCESS == cublasStat); + /* step 4: prepare workspace */ + // NOTE: CUDA 10.0 + // TODO: status = hipsparseSgtsvInterleavedBatch_bufferSizeExt( + status = cusparseSgtsvInterleavedBatch_bufferSizeExt( + cusparseH, + algo, + n, + d_dl, + d_d, + d_du, + d_X, + batchSize, + &lworkInBytes); + // CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); + assert(CUSPARSE_STATUS_SUCCESS == status); + + printf("lworkInBytes = %lld \n", (long long)lworkInBytes); + // CHECK: cudaStat1 = hipMalloc((void**)&d_work, lworkInBytes); + cudaStat1 = cudaMalloc((void**)&d_work, lworkInBytes); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + + /* step 5: solve Aj*xj = bj */ + // NOTE: CUDA 10.0 + // TODO: status = hipsparseSgtsvInterleavedBatch( + status = cusparseSgtsvInterleavedBatch( + cusparseH, + algo, + n, + d_dl, + d_d, + d_du, + d_X, + batchSize, + d_work); + // CHECK: cudaStat1 = hipDeviceSynchronize(); + cudaStat1 = cudaDeviceSynchronize(); + // CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); + assert(CUSPARSE_STATUS_SUCCESS == status); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + + /* step 6: convert X back to aggregate format */ + /* B = transpose(X) */ + // CHECK: cublasStat = hipblasSgeam( + // CHECK: HIPBLAS_OP_T + // CHECK: HIPBLAS_OP_T + cublasStat = cublasSgeam( + cublasH, + CUBLAS_OP_T, /* transa */ + CUBLAS_OP_T, /* transb, don't care */ + n, /* number of rows of B */ + batchSize, /* number of columns of B */ + &h_one, + d_X, /* X is batchSize-by-n */ + batchSize, /* leading dimension of X */ + &h_zero, + NULL, + n, /* don't cae */ + d_B, /* B is n-by-batchSize */ + n /* leading dimension of B */ + ); + // CHECK: assert(HIPBLAS_STATUS_SUCCESS == cublasStat); + assert(CUBLAS_STATUS_SUCCESS == cublasStat); + // CHECK: hipDeviceSynchronize(); + cudaDeviceSynchronize(); + + /* step 7: residual evaluation */ + // CHECK: cudaStat1 = hipMemcpy(X, d_B, sizeof(float)*n*batchSize, hipMemcpyDeviceToHost); + cudaStat1 = cudaMemcpy(X, d_B, sizeof(float)*n*batchSize, cudaMemcpyDeviceToHost); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: hipDeviceSynchronize(); + cudaDeviceSynchronize(); + + printf("==== x1 = inv(A1)*b1 \n"); + for (int j = 0; j < n; j++) { + printf("x1[%d] = %f\n", j, X[j]); + } + + float r1_nrminf; + residaul_eval( + n, + dl, + d, + du, + B, + X, + &r1_nrminf + ); + printf("|b1 - A1*x1| = %E\n", r1_nrminf); + + printf("\n==== x2 = inv(A2)*b2 \n"); + for (int j = 0; j < n; j++) { + printf("x2[%d] = %f\n", j, X[n + j]); + } + + float r2_nrminf; + residaul_eval( + n, + dl + n, + d + n, + du + n, + B + n, + X + n, + &r2_nrminf + ); + printf("|b2 - A2*x2| = %E\n", r2_nrminf); + + /* free resources */ + // CHECK: if (d_dl0) hipFree(d_dl0); + if (d_dl0) cudaFree(d_dl0); + // CHECK: if (d_d0) hipFree(d_d0); + if (d_d0) cudaFree(d_d0); + // CHECK: if (d_du0) hipFree(d_du0); + if (d_du0) cudaFree(d_du0); + // CHECK: if (d_dl) hipFree(d_dl); + if (d_dl) cudaFree(d_dl); + // CHECK: if (d_d) hipFree(d_d); + if (d_d) cudaFree(d_d); + // CHECK: if (d_du) hipFree(d_du); + if (d_du) cudaFree(d_du); + // CHECK: if (d_B) hipFree(d_B); + if (d_B) cudaFree(d_B); + // CHECK: if (d_X) hipFree(d_X); + if (d_X) cudaFree(d_X); + // CHECK: if (cusparseH) hipsparseDestroy(cusparseH); + if (cusparseH) cusparseDestroy(cusparseH); + // CHECK: if (cublasH) hipblasDestroy(cublasH); + if (cublasH) cublasDestroy(cublasH); + // CHECK: if (stream) hipStreamDestroy(stream); + if (stream) cudaStreamDestroy(stream); + // CHECK: hipDeviceReset(); + cudaDeviceReset(); + + return 0; +} diff --git a/projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_10.cu b/projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_10.cu new file mode 100644 index 0000000000..326c231de4 --- /dev/null +++ b/projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_10.cu @@ -0,0 +1,507 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args +#include +#include +#include +// CHECK: #include +#include +// CHECK: #include +#include +// CHECK: #include +#include + +// NOTE: CUDA 10.0 + +/* + * compute | b - A*x|_inf + */ +void residaul_eval( + int n, + const float *ds, + const float *dl, + const float *d, + const float *du, + const float *dw, + const float *b, + const float *x, + float *r_nrminf_ptr) +{ + float r_nrminf = 0; + for (int i = 0; i < n; i++) { + float dot = 0; + if (i > 1) { + dot += ds[i] * x[i - 2]; + } + if (i > 0) { + dot += dl[i] * x[i - 1]; + } + dot += d[i] * x[i]; + if (i < (n - 1)) { + dot += du[i] * x[i + 1]; + } + if (i < (n - 2)) { + dot += dw[i] * x[i + 2]; + } + float ri = b[i] - dot; + r_nrminf = (r_nrminf > fabs(ri)) ? r_nrminf : fabs(ri); + } + + *r_nrminf_ptr = r_nrminf; +} + +int main(int argc, char*argv[]) +{ + // CHECK: hipsparseHandle_t cusparseH = NULL; + cusparseHandle_t cusparseH = NULL; + // CHECK: hipblasHandle_t cublasH = NULL; + cublasHandle_t cublasH = NULL; + // CHECK: hipStream_t stream = NULL; + cudaStream_t stream = NULL; + // CHECK: hipsparseStatus_t status = HIPSPARSE_STATUS_SUCCESS; + cusparseStatus_t status = CUSPARSE_STATUS_SUCCESS; + // CHECK: hipblasStatus_t cublasStat = HIPBLAS_STATUS_SUCCESS; + cublasStatus_t cublasStat = CUBLAS_STATUS_SUCCESS; + // CHECK: hipError_t cudaStat1 = hipSuccess; + cudaError_t cudaStat1 = cudaSuccess; + + const int n = 4; + const int batchSize = 2; + + /* + * | 1 8 13 0 | | 1 | | -0.0592 | + * A1 =| 5 2 9 14 |, b1 = | 2 |, x1 = | 0.3428 | + * | 11 6 3 10 | | 3 | | -0.1295 | + * | 0 12 7 4 | | 4 | | 0.1982 | + * + * | 15 22 27 0 | | 5 | | -0.0012 | + * A2 =| 19 16 23 28 |, b2 = | 6 |, x2 = | 0.2792 | + * | 25 20 17 24 | | 7 | | -0.0416 | + * | 0 26 21 18 | | 8 | | 0.0898 | + */ + + /* + * A = (ds, dl, d, du, dw), B and X are in aggregate format + */ + const float ds[n * batchSize] = { 0, 0, 11, 12, 0, 0, 25, 26 }; + const float dl[n * batchSize] = { 0, 5, 6, 7, 0, 19, 20, 21 }; + const float d[n * batchSize] = { 1, 2, 3, 4, 15, 16, 17, 18 }; + const float du[n * batchSize] = { 8, 9, 10, 0, 22, 23, 24, 0 }; + const float dw[n * batchSize] = { 13,14, 0, 0, 27, 28, 0, 0 }; + const float B[n * batchSize] = { 1, 2, 3, 4, 5, 6, 7, 8 }; + float X[n * batchSize]; /* Xj = Aj \ Bj */ + +/* device memory + * (d_ds0, d_dl0, d_d0, d_du0, d_dw0) is aggregate format + * (d_ds, d_dl, d_d, d_du, d_dw) is interleaved format + */ + float *d_ds0 = NULL; + float *d_dl0 = NULL; + float *d_d0 = NULL; + float *d_du0 = NULL; + float *d_dw0 = NULL; + float *d_ds = NULL; + float *d_dl = NULL; + float *d_d = NULL; + float *d_du = NULL; + float *d_dw = NULL; + float *d_B = NULL; + float *d_X = NULL; + + size_t lworkInBytes = 0; + char *d_work = NULL; + + const float h_one = 1; + const float h_zero = 0; + + int algo = 0; /* QR factorization */ + + printf("example of gpsv (interleaved format) \n"); + printf("n = %d, batchSize = %d\n", n, batchSize); + + /* step 1: create cusparse/cublas handle, bind a stream */ + // CHECK: cudaStat1 = hipStreamCreateWithFlags(&stream, hipStreamNonBlocking); + cudaStat1 = cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: status = hipsparseCreate(&cusparseH); + status = cusparseCreate(&cusparseH); + // CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); + assert(CUSPARSE_STATUS_SUCCESS == status); + // CHECK: status = hipsparseSetStream(cusparseH, stream); + status = cusparseSetStream(cusparseH, stream); + // CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); + assert(CUSPARSE_STATUS_SUCCESS == status); + // CHECK: cublasStat = hipblasCreate(&cublasH); + cublasStat = cublasCreate(cublasH); + // CHECK: assert(HIPBLAS_STATUS_SUCCESS == cublasStat); + assert(CUBLAS_STATUS_SUCCESS == cublasStat); + // CHECK: cublasStat = hipblasSetStream(cublasH, stream); + cublasStat = cublasSetStream(cublasH, stream); + // CHECK: assert(HIPBLAS_STATUS_SUCCESS == cublasStat); + assert(CUBLAS_STATUS_SUCCESS == cublasStat); + /* step 2: allocate device memory */ + // CHECK: cudaStat1 = hipMalloc((void**)&d_ds0, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_ds0, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_dl0, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_dl0, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_d0, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_d0, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_du0, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_du0, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_dw0, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_dw0, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_ds, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_ds, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_dl, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_dl, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_d, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_d, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_du, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_du, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_dw, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_dw, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_B, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_B, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_X, sizeof(float)*n*batchSize); + cudaStat1 = cudaMalloc((void**)&d_X, sizeof(float)*n*batchSize); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + /* step 3: prepare data in device, interleaved format */ + // CHECK: cudaStat1 = hipMemcpy(d_ds0, ds, sizeof(float)*n*batchSize, hipMemcpyHostToDevice); + cudaStat1 = cudaMemcpy(d_ds0, ds, sizeof(float)*n*batchSize, cudaMemcpyHostToDevice); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMemcpy(d_dl0, dl, sizeof(float)*n*batchSize, hipMemcpyHostToDevice); + cudaStat1 = cudaMemcpy(d_dl0, dl, sizeof(float)*n*batchSize, cudaMemcpyHostToDevice); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMemcpy(d_d0, d, sizeof(float)*n*batchSize, hipMemcpyHostToDevice); + cudaStat1 = cudaMemcpy(d_d0, d, sizeof(float)*n*batchSize, cudaMemcpyHostToDevice); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMemcpy(d_du0, du, sizeof(float)*n*batchSize, hipMemcpyHostToDevice); + cudaStat1 = cudaMemcpy(d_du0, du, sizeof(float)*n*batchSize, cudaMemcpyHostToDevice); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMemcpy(d_dw0, dw, sizeof(float)*n*batchSize, hipMemcpyHostToDevice); + cudaStat1 = cudaMemcpy(d_dw0, dw, sizeof(float)*n*batchSize, cudaMemcpyHostToDevice); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMemcpy(d_B, B, sizeof(float)*n*batchSize, hipMemcpyHostToDevice); + cudaStat1 = cudaMemcpy(d_B, B, sizeof(float)*n*batchSize, cudaMemcpyHostToDevice); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: hipDeviceSynchronize(); + cudaDeviceSynchronize(); + /* convert ds to interleaved format + * ds = transpose(ds0) */ + // CHECK: cublasStat = hipblasSgeam( + // CHECK: HIPBLAS_OP_T, + // CHECK: HIPBLAS_OP_T, + cublasStat = cublasSgeam( + cublasH, + CUBLAS_OP_T, /* transa */ + CUBLAS_OP_T, /* transb, don't care */ + batchSize, /* number of rows of ds */ + n, /* number of columns of ds */ + &h_one, + d_ds0, /* ds0 is n-by-batchSize */ + n, /* leading dimension of ds0 */ + &h_zero, + NULL, + n, /* don't cae */ + d_ds, /* ds is batchSize-by-n */ + batchSize); /* leading dimension of ds */ + // CHECK: assert(HIPBLAS_STATUS_SUCCESS == cublasStat); + assert(CUBLAS_STATUS_SUCCESS == cublasStat); + /* convert dl to interleaved format + * dl = transpose(dl0) + */ + // CHECK: cublasStat = hipblasSgeam( + // CHECK: HIPBLAS_OP_T, + // CHECK: HIPBLAS_OP_T, + cublasStat = cublasSgeam( + cublasH, + CUBLAS_OP_T, /* transa */ + CUBLAS_OP_T, /* transb, don't care */ + batchSize, /* number of rows of dl */ + n, /* number of columns of dl */ + &h_one, + d_dl0, /* dl0 is n-by-batchSize */ + n, /* leading dimension of dl0 */ + &h_zero, + NULL, + n, /* don't cae */ + d_dl, /* dl is batchSize-by-n */ + batchSize /* leading dimension of dl */ + ); + // CHECK: assert(HIPBLAS_STATUS_SUCCESS == cublasStat); + assert(CUBLAS_STATUS_SUCCESS == cublasStat); + + /* convert d to interleaved format + * d = transpose(d0) + */ + // CHECK: cublasStat = hipblasSgeam( + // CHECK: HIPBLAS_OP_T, + // CHECK: HIPBLAS_OP_T, + cublasStat = cublasSgeam( + cublasH, + CUBLAS_OP_T, /* transa */ + CUBLAS_OP_T, /* transb, don't care */ + batchSize, /* number of rows of d */ + n, /* number of columns of d */ + &h_one, + d_d0, /* d0 is n-by-batchSize */ + n, /* leading dimension of d0 */ + &h_zero, + NULL, + n, /* don't cae */ + d_d, /* d is batchSize-by-n */ + batchSize /* leading dimension of d */ + ); + // CHECK: assert(HIPBLAS_STATUS_SUCCESS == cublasStat); + assert(CUBLAS_STATUS_SUCCESS == cublasStat); + + /* convert du to interleaved format + * du = transpose(du0) + */ + // CHECK: cublasStat = hipblasSgeam( + // CHECK: HIPBLAS_OP_T, + // CHECK: HIPBLAS_OP_T, + cublasStat = cublasSgeam( + cublasH, + CUBLAS_OP_T, /* transa */ + CUBLAS_OP_T, /* transb, don't care */ + batchSize, /* number of rows of du */ + n, /* number of columns of du */ + &h_one, + d_du0, /* du0 is n-by-batchSize */ + n, /* leading dimension of du0 */ + &h_zero, + NULL, + n, /* don't cae */ + d_du, /* du is batchSize-by-n */ + batchSize /* leading dimension of du */ + ); + // CHECK: assert(HIPBLAS_STATUS_SUCCESS == cublasStat); + assert(CUBLAS_STATUS_SUCCESS == cublasStat); + /* convert dw to interleaved format + * dw = transpose(dw0) + */ + // CHECK: cublasStat = hipblasSgeam( + // CHECK: HIPBLAS_OP_T, + // CHECK: HIPBLAS_OP_T, + cublasStat = cublasSgeam( + cublasH, + CUBLAS_OP_T, /* transa */ + CUBLAS_OP_T, /* transb, don't care */ + batchSize, /* number of rows of dw */ + n, /* number of columns of dw */ + &h_one, + d_dw0, /* dw0 is n-by-batchSize */ + n, /* leading dimension of dw0 */ + &h_zero, + NULL, + n, /* don't cae */ + d_dw, /* dw is batchSize-by-n */ + batchSize /* leading dimension of dw */ + ); + // CHECK: assert(HIPBLAS_STATUS_SUCCESS == cublasStat); + assert(CUBLAS_STATUS_SUCCESS == cublasStat); + + /* convert B to interleaved format + * X = transpose(B) + */ + // CHECK: cublasStat = hipblasSgeam( + // CHECK: HIPBLAS_OP_T, + // CHECK: HIPBLAS_OP_T, + cublasStat = cublasSgeam( + cublasH, + CUBLAS_OP_T, /* transa */ + CUBLAS_OP_T, /* transb, don't care */ + batchSize, /* number of rows of X */ + n, /* number of columns of X */ + &h_one, + d_B, /* B is n-by-batchSize */ + n, /* leading dimension of B */ + &h_zero, + NULL, + n, /* don't cae */ + d_X, /* X is batchSize-by-n */ + batchSize /* leading dimension of X */ + ); + // CHECK: assert(HIPBLAS_STATUS_SUCCESS == cublasStat); + assert(CUBLAS_STATUS_SUCCESS == cublasStat); + + /* step 4: prepare workspace */ + // NOTE: CUDA 10.0 + // TODO: status = hipsparseSgpsvInterleavedBatch_bufferSizeExt( + status = cusparseSgpsvInterleavedBatch_bufferSizeExt( + cusparseH, + algo, + n, + d_ds, + d_dl, + d_d, + d_du, + d_dw, + d_X, + batchSize, + &lworkInBytes); + // CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); + assert(CUSPARSE_STATUS_SUCCESS == status); + + printf("lworkInBytes = %lld \n", (long long)lworkInBytes); + // CHECK: cudaStat1 = hipMalloc((void**)&d_work, lworkInBytes); + cudaStat1 = cudaMalloc((void**)&d_work, lworkInBytes); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + /* step 5: solve Aj*xj = bj */ + // NOTE: CUDA 10.0 + // TODO: status = hipsparseSgpsvInterleavedBatch( + status = cusparseSgpsvInterleavedBatch( + cusparseH, + algo, + n, + d_ds, + d_dl, + d_d, + d_du, + d_dw, + d_X, + batchSize, + d_work); + // CHECK: cudaStat1 = hipDeviceSynchronize(); + cudaStat1 = cudaDeviceSynchronize(); + // CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); + assert(CUSPARSE_STATUS_SUCCESS == status); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + + /* step 6: convert X back to aggregate format */ + /* B = transpose(X) */ + // CHECK: cublasStat = hipblasSgeam( + // CHECK: HIPBLAS_OP_T, + // CHECK: HIPBLAS_OP_T, + cublasStat = cublasSgeam( + cublasH, + CUBLAS_OP_T, /* transa */ + CUBLAS_OP_T, /* transb, don't care */ + n, /* number of rows of B */ + batchSize, /* number of columns of B */ + &h_one, + d_X, /* X is batchSize-by-n */ + batchSize, /* leading dimension of X */ + &h_zero, + NULL, + n, /* don't cae */ + d_B, /* B is n-by-batchSize */ + n /* leading dimension of B */ + ); + // CHECK: assert(HIPBLAS_STATUS_SUCCESS == cublasStat); + assert(CUBLAS_STATUS_SUCCESS == cublasStat); + // CHECK: hipDeviceSynchronize(); + cudaDeviceSynchronize(); + + /* step 7: residual evaluation */ + // CHECK: cudaStat1 = hipMemcpy(X, d_B, sizeof(float)*n*batchSize, hipMemcpyDeviceToHost); + cudaStat1 = cudaMemcpy(X, d_B, sizeof(float)*n*batchSize, cudaMemcpyDeviceToHost); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: hipDeviceSynchronize(); + cudaDeviceSynchronize(); + + printf("==== x1 = inv(A1)*b1 \n"); + for (int j = 0; j < n; j++) { + printf("x1[%d] = %f\n", j, X[j]); + } + + float r1_nrminf; + residaul_eval( + n, + ds, + dl, + d, + du, + dw, + B, + X, + &r1_nrminf + ); + printf("|b1 - A1*x1| = %E\n", r1_nrminf); + printf("\n==== x2 = inv(A2)*b2 \n"); + for (int j = 0; j < n; j++) { + printf("x2[%d] = %f\n", j, X[n + j]); + } + + float r2_nrminf; + residaul_eval( + n, + ds + n, + dl + n, + d + n, + du + n, + dw + n, + B + n, + X + n, + &r2_nrminf + ); + printf("|b2 - A2*x2| = %E\n", r2_nrminf); + + /* free resources */ + // CHECK: if (d_ds0) hipFree(d_ds0); + if (d_ds0) cudaFree(d_ds0); + // CHECK: if (d_dl0) hipFree(d_dl0); + if (d_dl0) cudaFree(d_dl0); + // CHECK: if (d_d0) hipFree(d_d0); + if (d_d0) cudaFree(d_d0); + // CHECK: if (d_du0) hipFree(d_du0); + if (d_du0) cudaFree(d_du0); + // CHECK: if (d_dw0) hipFree(d_dw0); + if (d_dw0) cudaFree(d_dw0); + // CHECK: if (d_ds) hipFree(d_ds); + if (d_ds) cudaFree(d_ds); + // CHECK: if (d_dl) hipFree(d_dl); + if (d_dl) cudaFree(d_dl); + // CHECK: if (d_d) hipFree(d_d); + if (d_d) cudaFree(d_d); + // CHECK: if (d_du) hipFree(d_du); + if (d_du) cudaFree(d_du); + // CHECK: if (d_dw) hipFree(d_dw); + if (d_dw) cudaFree(d_dw); + // CHECK: if (d_B) hipFree(d_B); + if (d_B) cudaFree(d_B); + // CHECK: if (d_X) hipFree(d_X); + if (d_X) cudaFree(d_X); + // CHECK: if (cusparseH) hipsparseDestroy(cusparseH); + if (cusparseH) cusparseDestroy(cusparseH); + // CHECK: if (cublasH) hipblasDestroy(cublasH); + if (cublasH) cublasDestroy(cublasH); + // CHECK: if (stream) hipStreamDestroy(stream); + if (stream) cudaStreamDestroy(stream); + // CHECK: hipDeviceReset(); + cudaDeviceReset(); + + return 0; +} diff --git a/projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_11.cu b/projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_11.cu new file mode 100644 index 0000000000..2d905bcf6d --- /dev/null +++ b/projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_11.cu @@ -0,0 +1,327 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args +#include +#include +#include +// CHECK: #include +#include +// CHECK: #include +#include + +// NOTE: CUDA 10.0 + +/* compute | b - A*x|_inf */ +void residaul_eval( + int n, + // CHECK: const hipsparseMatDescr_t descrA, + const cusparseMatDescr_t descrA, + const float *csrVal, + const int *csrRowPtr, + const int *csrColInd, + const float *b, + const float *x, + float *r_nrminf_ptr) +{ + // CHECK: const int base = (hipsparseGetMatIndexBase(descrA) != HIPSPARSE_INDEX_BASE_ONE) ? 0 : 1; + const int base = (cusparseGetMatIndexBase(descrA) != CUSPARSE_INDEX_BASE_ONE) ? 0 : 1; + // CHECK: const int lower = (HIPSPARSE_FILL_MODE_LOWER == hipsparseGetMatFillMode(descrA)) ? 1 : 0; + const int lower = (CUSPARSE_FILL_MODE_LOWER == cusparseGetMatFillMode(descrA)) ? 1 : 0; + // CHECK: const int unit = (HIPSPARSE_DIAG_TYPE_UNIT == hipsparseGetMatDiagType(descrA)) ? 1 : 0; + const int unit = (CUSPARSE_DIAG_TYPE_UNIT == cusparseGetMatDiagType(descrA)) ? 1 : 0; + + float r_nrminf = 0; + for (int row = 0; row < n; row++) { + const int start = csrRowPtr[row] - base; + const int end = csrRowPtr[row + 1] - base; + float dot = 0; + for (int colidx = start; colidx < end; colidx++) { + const int col = csrColInd[colidx] - base; + float Aij = csrVal[colidx]; + float xj = x[col]; + if ((row == col) && unit) { + Aij = 1.0; + } + int valid = (row >= col) && lower || + (row <= col) && !lower; + if (valid) { + dot += Aij * xj; + } + } + float ri = b[row] - dot; + r_nrminf = (r_nrminf > fabs(ri)) ? r_nrminf : fabs(ri); + } + *r_nrminf_ptr = r_nrminf; +} + +int main(int argc, char*argv[]) +{ + // CHECK: hipsparseHandle_t handle = NULL; + cusparseHandle_t handle = NULL; + // CHECK: hipStream_t stream = NULL; + cudaStream_t stream = NULL; + // CHECK: hipsparseMatDescr_t descrA = NULL; + cusparseMatDescr_t descrA = NULL; + // NOTE: CUDA 10.0 + // TODO: csrsm2Info_t info = NULL; + csrsm2Info_t info = NULL; + // CHECK: hipsparseStatus_t status = HIPSPARSE_STATUS_SUCCESS; + cusparseStatus_t status = CUSPARSE_STATUS_SUCCESS; + // CHECK: hipError_t cudaStat1 = hipSuccess; + cudaError_t cudaStat1 = cudaSuccess; + const int nrhs = 2; + const int n = 4; + const int nnzA = 9; + // CHECK: const hipsparseSolvePolicy_t policy = HIPSPARSE_SOLVE_POLICY_NO_LEVEL; + const cusparseSolvePolicy_t policy = CUSPARSE_SOLVE_POLICY_NO_LEVEL; + const float h_one = 1.0; + /* + * | 1 0 2 -3 | + * | 0 4 0 0 | + * A = | 5 0 6 7 | + * | 0 8 0 9 | + * + * Regard A as a lower triangle matrix L with non-unit diagonal. + * | 1 5 | | 1 5 | + * Given B = | 2 6 |, X = L \ B = | 0.5 1.5 | + * | 3 7 | | -0.3333 -3 | + * | 4 8 | | 0 -0.4444 | + */ + const int csrRowPtrA[n + 1] = { 1, 4, 5, 8, 10 }; + const int csrColIndA[nnzA] = { 1, 3, 4, 2, 1, 3, 4, 2, 4 }; + const float csrValA[nnzA] = { 1, 2, -3, 4, 5, 6, 7, 8, 9 }; + const float B[n*nrhs] = { 1,2,3,4,5,6,7,8 }; + float X[n*nrhs]; + + int *d_csrRowPtrA = NULL; + int *d_csrColIndA = NULL; + float *d_csrValA = NULL; + float *d_B = NULL; + + size_t lworkInBytes = 0; + char *d_work = NULL; + + const int algo = 0; /* non-block version */ + + printf("example of csrsm2 \n"); + + /* step 1: create cusparse handle, bind a stream */ + // CHECK: cudaStat1 = hipStreamCreateWithFlags(&stream, hipStreamNonBlocking); + cudaStat1 = cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: status = hipsparseCreate(&handle); + status = cusparseCreate(&handle); + // CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); + assert(CUSPARSE_STATUS_SUCCESS == status); + + status = cusparseSetStream(handle, stream); + // CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); + assert(CUSPARSE_STATUS_SUCCESS == status); + + // NOTE: CUDA 10.0 + // TODO: status = hipsparseCreateCsrsm2Info(&info); + status = cusparseCreateCsrsm2Info(&info); + // CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); + assert(CUSPARSE_STATUS_SUCCESS == status); + + /* step 2: configuration of matrix A */ + status = cusparseCreateMatDescr(&descrA); + // CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); + assert(CUSPARSE_STATUS_SUCCESS == status); + /* A is base-1*/ + // CHECK: hipsparseSetMatIndexBase(descrA, HIPSPARSE_INDEX_BASE_ONE); + cusparseSetMatIndexBase(descrA, CUSPARSE_INDEX_BASE_ONE); + // CHECK: hipsparseSetMatType(descrA, HIPSPARSE_MATRIX_TYPE_GENERAL); + cusparseSetMatType(descrA, CUSPARSE_MATRIX_TYPE_GENERAL); + /* A is lower triangle */ + // CHECK: hipsparseSetMatFillMode(descrA, HIPSPARSE_FILL_MODE_LOWER); + cusparseSetMatFillMode(descrA, CUSPARSE_FILL_MODE_LOWER); + /* A has non unit diagonal */ + // CHECK: hipsparseSetMatDiagType(descrA, HIPSPARSE_DIAG_TYPE_NON_UNIT); + cusparseSetMatDiagType(descrA, CUSPARSE_DIAG_TYPE_NON_UNIT); + // CHECK: cudaStat1 = hipMalloc((void**)&d_csrRowPtrA, sizeof(int)*(n + 1)); + cudaStat1 = cudaMalloc((void**)&d_csrRowPtrA, sizeof(int)*(n + 1)); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_csrColIndA, sizeof(int)*nnzA); + cudaStat1 = cudaMalloc((void**)&d_csrColIndA, sizeof(int)*nnzA); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_csrValA, sizeof(float)*nnzA); + cudaStat1 = cudaMalloc((void**)&d_csrValA, sizeof(float)*nnzA); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMalloc((void**)&d_B, sizeof(float)*n*nrhs); + cudaStat1 = cudaMalloc((void**)&d_B, sizeof(float)*n*nrhs); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMemcpy(d_csrRowPtrA, csrRowPtrA, sizeof(int)*(n + 1), hipMemcpyHostToDevice); + cudaStat1 = cudaMemcpy(d_csrRowPtrA, csrRowPtrA, sizeof(int)*(n + 1), cudaMemcpyHostToDevice); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMemcpy(d_csrColIndA, csrColIndA, sizeof(int)*nnzA, hipMemcpyHostToDevice); + cudaStat1 = cudaMemcpy(d_csrColIndA, csrColIndA, sizeof(int)*nnzA, cudaMemcpyHostToDevice); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMemcpy(d_csrValA, csrValA, sizeof(float)*nnzA, hipMemcpyHostToDevice); + cudaStat1 = cudaMemcpy(d_csrValA, csrValA, sizeof(float)*nnzA, cudaMemcpyHostToDevice); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: cudaStat1 = hipMemcpy(d_B, B, sizeof(float)*n*nrhs, hipMemcpyHostToDevice); + cudaStat1 = cudaMemcpy(d_B, B, sizeof(float)*n*nrhs, cudaMemcpyHostToDevice); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + + /* step 3: query workspace */ + // NOTE: CUDA 10.0 + // TODO: status = hipsparseScsrsm2_bufferSizeExt( + // CHECK: HIPSPARSE_OPERATION_NON_TRANSPOSE, + // CHECK: HIPSPARSE_OPERATION_NON_TRANSPOSE, + status = cusparseScsrsm2_bufferSizeExt( + handle, + algo, + CUSPARSE_OPERATION_NON_TRANSPOSE, /* transA */ + CUSPARSE_OPERATION_NON_TRANSPOSE, /* transB */ + n, + nrhs, + nnzA, + &h_one, + descrA, + d_csrValA, + d_csrRowPtrA, + d_csrColIndA, + d_B, + n, /* ldb */ + info, + policy, + &lworkInBytes); + // CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); + assert(CUSPARSE_STATUS_SUCCESS == status); + + printf("lworkInBytes = %lld \n", (long long)lworkInBytes); + // CHECK: if (NULL != d_work) { hipFree(d_work); } + if (NULL != d_work) { cudaFree(d_work); } + // CHECK: cudaStat1 = hipMalloc((void**)&d_work, lworkInBytes); + cudaStat1 = cudaMalloc((void**)&d_work, lworkInBytes); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + + /* step 4: analysis */ + // NOTE: CUDA 10.0 + // TODO: status = hipsparseScsrsm2_analysis( + // CHECK: HIPSPARSE_OPERATION_NON_TRANSPOSE, + // CHECK: HIPSPARSE_OPERATION_NON_TRANSPOSE, + status = cusparseScsrsm2_analysis( + handle, + algo, + CUSPARSE_OPERATION_NON_TRANSPOSE, /* transA */ + CUSPARSE_OPERATION_NON_TRANSPOSE, /* transB */ + n, + nrhs, + nnzA, + &h_one, + descrA, + d_csrValA, + d_csrRowPtrA, + d_csrColIndA, + d_B, + n, /* ldb */ + info, + policy, + d_work); + // CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); + assert(CUSPARSE_STATUS_SUCCESS == status); + /* step 5: solve L * X = B */ + // NOTE: CUDA 10.0 + // TODO: status = hipsparseScsrsm2_solve( + // CHECK: HIPSPARSE_OPERATION_NON_TRANSPOSE, + // CHECK: HIPSPARSE_OPERATION_NON_TRANSPOSE, + status = cusparseScsrsm2_solve( + handle, + algo, + CUSPARSE_OPERATION_NON_TRANSPOSE, /* transA */ + CUSPARSE_OPERATION_NON_TRANSPOSE, /* transB */ + n, + nrhs, + nnzA, + &h_one, + descrA, + d_csrValA, + d_csrRowPtrA, + d_csrColIndA, + d_B, + n, /* ldb */ + info, + policy, + d_work); + // CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); + assert(CUSPARSE_STATUS_SUCCESS == status); + // CHECK: cudaStat1 = hipDeviceSynchronize(); + cudaStat1 = cudaDeviceSynchronize(); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + + /* step 6:measure residual B - A*X */ + // CHECK: cudaStat1 = hipMemcpy(X, d_B, sizeof(float)*n*nrhs, hipMemcpyDeviceToHost); + cudaStat1 = cudaMemcpy(X, d_B, sizeof(float)*n*nrhs, cudaMemcpyDeviceToHost); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // CHECK: hipDeviceSynchronize(); + cudaDeviceSynchronize(); + + printf("==== x1 = inv(A)*b1 \n"); + for (int j = 0; j < n; j++) { + printf("x1[%d] = %f\n", j, X[j]); + } + float r1_nrminf; + residaul_eval( + n, + descrA, + csrValA, + csrRowPtrA, + csrColIndA, + B, + X, + &r1_nrminf + ); + printf("|b1 - A*x1| = %E\n", r1_nrminf); + + printf("==== x2 = inv(A)*b2 \n"); + for (int j = 0; j < n; j++) { + printf("x2[%d] = %f\n", j, X[n + j]); + } + float r2_nrminf; + residaul_eval( + n, + descrA, + csrValA, + csrRowPtrA, + csrColIndA, + B + n, + X + n, + &r2_nrminf + ); + printf("|b2 - A*x2| = %E\n", r2_nrminf); + + /* free resources */ + // CHECK: if (d_csrRowPtrA) hipFree(d_csrRowPtrA); + if (d_csrRowPtrA) cudaFree(d_csrRowPtrA); + // CHECK: if (d_csrColIndA) hipFree(d_csrColIndA); + if (d_csrColIndA) cudaFree(d_csrColIndA); + // CHECK: if (d_csrValA) hipFree(d_csrValA); + if (d_csrValA) cudaFree(d_csrValA); + // CHECK: if (d_B) hipFree(d_B); + if (d_B) cudaFree(d_B); + // CHECK: if (handle) hipsparseDestroy(handle); + if (handle) cusparseDestroy(handle); + // CHECK: if (stream) hipStreamDestroy(stream); + if (stream) cudaStreamDestroy(stream); + // CHECK: if (descrA) hipsparseDestroyMatDescr(descrA); + if (descrA) cusparseDestroyMatDescr(descrA); + // NOTE: CUDA 10.0 + // TODO: if (info) hipsparseDestroyCsrsm2Info(info); + if (info) cusparseDestroyCsrsm2Info(info); + // CHECK: hipDeviceReset(); + cudaDeviceReset(); + + return 0; +} diff --git a/projects/hip/tests/hipify-clang/lit.cfg b/projects/hip/tests/hipify-clang/lit.cfg index b9ff0b2495..f959147cd6 100644 --- a/projects/hip/tests/hipify-clang/lit.cfg +++ b/projects/hip/tests/hipify-clang/lit.cfg @@ -26,6 +26,9 @@ config.excludes = ['cmdparser.hpp'] config.cuda_version = "@CUDA_VERSION@" if config.cuda_version not in ['10.0']: config.excludes.append('cuSPARSE_08.cu') + config.excludes.append('cuSPARSE_09.cu') + config.excludes.append('cuSPARSE_10.cu') + config.excludes.append('cuSPARSE_11.cu') # test_exec_root: The path where tests are located (default is the test suite root). #config.test_exec_root = config.test_source_root From 5cae606a0faef4d7ed33676cb8e3e543d58ae9c7 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 4 Dec 2018 20:47:34 +0300 Subject: [PATCH 10/10] [HIPIFY][tests] Reorganization [ROCm/hip commit: 09612ac03fe70aae340967e95ac91801373a4643] --- .../hipify-clang/{ => unit_tests/headers}/headers_test_01.cu | 0 .../hipify-clang/{ => unit_tests/headers}/headers_test_02.cu | 0 .../hipify-clang/{ => unit_tests/headers}/headers_test_03.cu | 0 .../hipify-clang/{ => unit_tests/headers}/headers_test_04.cu | 0 .../hipify-clang/{ => unit_tests/headers}/headers_test_05.cu | 0 .../hipify-clang/{ => unit_tests/headers}/headers_test_06.cu | 0 .../hipify-clang/{ => unit_tests/headers}/headers_test_07.cu | 0 .../hipify-clang/{ => unit_tests/headers}/headers_test_08.cu | 0 .../hipify-clang/{ => unit_tests/headers}/headers_test_09.cu | 0 .../{ => unit_tests/libraries}/cuBLAS/cublas_0_based_indexing.cu | 0 .../{ => unit_tests/libraries}/cuBLAS/cublas_1_based_indexing.cu | 0 .../libraries}/cuBLAS/cublas_sgemm_matrix_multiplication.cu | 0 .../{ => unit_tests/libraries}/cuComplex/cuComplex_Julia.cu | 0 .../{ => unit_tests/libraries}/cuDNN/cudnn_convolution_forward.cu | 0 .../{ => unit_tests/libraries}/cuDNN/cudnn_softmax.cu | 0 .../hipify-clang/{ => unit_tests/libraries}/cuFFT/simple_cufft.cu | 0 .../libraries}/cuRAND/benchmark_curand_generate.cpp | 0 .../{ => unit_tests/libraries}/cuRAND/benchmark_curand_kernel.cpp | 0 .../hipify-clang/{ => unit_tests/libraries}/cuRAND/cmdparser.hpp | 0 .../{ => unit_tests/libraries}/cuRAND/poisson_api_example.cu | 0 .../{ => unit_tests/libraries}/cuSPARSE/cuSPARSE_01.cu | 0 .../{ => unit_tests/libraries}/cuSPARSE/cuSPARSE_02.cu | 0 .../{ => unit_tests/libraries}/cuSPARSE/cuSPARSE_03.cu | 0 .../{ => unit_tests/libraries}/cuSPARSE/cuSPARSE_04.cu | 0 .../{ => unit_tests/libraries}/cuSPARSE/cuSPARSE_05.cu | 0 .../{ => unit_tests/libraries}/cuSPARSE/cuSPARSE_06.cu | 0 .../{ => unit_tests/libraries}/cuSPARSE/cuSPARSE_07.cu | 0 .../{ => unit_tests/libraries}/cuSPARSE/cuSPARSE_08.cu | 0 .../{ => unit_tests/libraries}/cuSPARSE/cuSPARSE_09.cu | 0 .../{ => unit_tests/libraries}/cuSPARSE/cuSPARSE_10.cu | 0 .../{ => unit_tests/libraries}/cuSPARSE/cuSPARSE_11.cu | 0 .../hip/tests/hipify-clang/{ => unit_tests/samples}/allocators.cu | 0 projects/hip/tests/hipify-clang/{ => unit_tests/samples}/axpy.cu | 0 .../hip/tests/hipify-clang/{ => unit_tests/samples}/coalescing.cu | 0 .../tests/hipify-clang/{ => unit_tests/samples}/cudaRegister.cu | 0 .../{ => unit_tests/samples}/dynamic_shared_memory.cu | 0 projects/hip/tests/hipify-clang/{ => unit_tests/samples}/intro.cu | 0 .../hip/tests/hipify-clang/{ => unit_tests/samples}/square.cu | 0 .../hipify-clang/{ => unit_tests/samples}/static_shared_memory.cu | 0 .../hip/tests/hipify-clang/{ => unit_tests/samples}/vec_add.cu | 0 40 files changed, 0 insertions(+), 0 deletions(-) rename projects/hip/tests/hipify-clang/{ => unit_tests/headers}/headers_test_01.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/headers}/headers_test_02.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/headers}/headers_test_03.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/headers}/headers_test_04.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/headers}/headers_test_05.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/headers}/headers_test_06.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/headers}/headers_test_07.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/headers}/headers_test_08.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/headers}/headers_test_09.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/libraries}/cuBLAS/cublas_0_based_indexing.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/libraries}/cuBLAS/cublas_1_based_indexing.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/libraries}/cuBLAS/cublas_sgemm_matrix_multiplication.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/libraries}/cuComplex/cuComplex_Julia.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/libraries}/cuDNN/cudnn_convolution_forward.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/libraries}/cuDNN/cudnn_softmax.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/libraries}/cuFFT/simple_cufft.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/libraries}/cuRAND/benchmark_curand_generate.cpp (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/libraries}/cuRAND/benchmark_curand_kernel.cpp (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/libraries}/cuRAND/cmdparser.hpp (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/libraries}/cuRAND/poisson_api_example.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/libraries}/cuSPARSE/cuSPARSE_01.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/libraries}/cuSPARSE/cuSPARSE_02.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/libraries}/cuSPARSE/cuSPARSE_03.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/libraries}/cuSPARSE/cuSPARSE_04.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/libraries}/cuSPARSE/cuSPARSE_05.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/libraries}/cuSPARSE/cuSPARSE_06.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/libraries}/cuSPARSE/cuSPARSE_07.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/libraries}/cuSPARSE/cuSPARSE_08.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/libraries}/cuSPARSE/cuSPARSE_09.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/libraries}/cuSPARSE/cuSPARSE_10.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/libraries}/cuSPARSE/cuSPARSE_11.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/samples}/allocators.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/samples}/axpy.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/samples}/coalescing.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/samples}/cudaRegister.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/samples}/dynamic_shared_memory.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/samples}/intro.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/samples}/square.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/samples}/static_shared_memory.cu (100%) rename projects/hip/tests/hipify-clang/{ => unit_tests/samples}/vec_add.cu (100%) diff --git a/projects/hip/tests/hipify-clang/headers_test_01.cu b/projects/hip/tests/hipify-clang/unit_tests/headers/headers_test_01.cu similarity index 100% rename from projects/hip/tests/hipify-clang/headers_test_01.cu rename to projects/hip/tests/hipify-clang/unit_tests/headers/headers_test_01.cu diff --git a/projects/hip/tests/hipify-clang/headers_test_02.cu b/projects/hip/tests/hipify-clang/unit_tests/headers/headers_test_02.cu similarity index 100% rename from projects/hip/tests/hipify-clang/headers_test_02.cu rename to projects/hip/tests/hipify-clang/unit_tests/headers/headers_test_02.cu diff --git a/projects/hip/tests/hipify-clang/headers_test_03.cu b/projects/hip/tests/hipify-clang/unit_tests/headers/headers_test_03.cu similarity index 100% rename from projects/hip/tests/hipify-clang/headers_test_03.cu rename to projects/hip/tests/hipify-clang/unit_tests/headers/headers_test_03.cu diff --git a/projects/hip/tests/hipify-clang/headers_test_04.cu b/projects/hip/tests/hipify-clang/unit_tests/headers/headers_test_04.cu similarity index 100% rename from projects/hip/tests/hipify-clang/headers_test_04.cu rename to projects/hip/tests/hipify-clang/unit_tests/headers/headers_test_04.cu diff --git a/projects/hip/tests/hipify-clang/headers_test_05.cu b/projects/hip/tests/hipify-clang/unit_tests/headers/headers_test_05.cu similarity index 100% rename from projects/hip/tests/hipify-clang/headers_test_05.cu rename to projects/hip/tests/hipify-clang/unit_tests/headers/headers_test_05.cu diff --git a/projects/hip/tests/hipify-clang/headers_test_06.cu b/projects/hip/tests/hipify-clang/unit_tests/headers/headers_test_06.cu similarity index 100% rename from projects/hip/tests/hipify-clang/headers_test_06.cu rename to projects/hip/tests/hipify-clang/unit_tests/headers/headers_test_06.cu diff --git a/projects/hip/tests/hipify-clang/headers_test_07.cu b/projects/hip/tests/hipify-clang/unit_tests/headers/headers_test_07.cu similarity index 100% rename from projects/hip/tests/hipify-clang/headers_test_07.cu rename to projects/hip/tests/hipify-clang/unit_tests/headers/headers_test_07.cu diff --git a/projects/hip/tests/hipify-clang/headers_test_08.cu b/projects/hip/tests/hipify-clang/unit_tests/headers/headers_test_08.cu similarity index 100% rename from projects/hip/tests/hipify-clang/headers_test_08.cu rename to projects/hip/tests/hipify-clang/unit_tests/headers/headers_test_08.cu diff --git a/projects/hip/tests/hipify-clang/headers_test_09.cu b/projects/hip/tests/hipify-clang/unit_tests/headers/headers_test_09.cu similarity index 100% rename from projects/hip/tests/hipify-clang/headers_test_09.cu rename to projects/hip/tests/hipify-clang/unit_tests/headers/headers_test_09.cu diff --git a/projects/hip/tests/hipify-clang/cuBLAS/cublas_0_based_indexing.cu b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuBLAS/cublas_0_based_indexing.cu similarity index 100% rename from projects/hip/tests/hipify-clang/cuBLAS/cublas_0_based_indexing.cu rename to projects/hip/tests/hipify-clang/unit_tests/libraries/cuBLAS/cublas_0_based_indexing.cu diff --git a/projects/hip/tests/hipify-clang/cuBLAS/cublas_1_based_indexing.cu b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuBLAS/cublas_1_based_indexing.cu similarity index 100% rename from projects/hip/tests/hipify-clang/cuBLAS/cublas_1_based_indexing.cu rename to projects/hip/tests/hipify-clang/unit_tests/libraries/cuBLAS/cublas_1_based_indexing.cu diff --git a/projects/hip/tests/hipify-clang/cuBLAS/cublas_sgemm_matrix_multiplication.cu b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuBLAS/cublas_sgemm_matrix_multiplication.cu similarity index 100% rename from projects/hip/tests/hipify-clang/cuBLAS/cublas_sgemm_matrix_multiplication.cu rename to projects/hip/tests/hipify-clang/unit_tests/libraries/cuBLAS/cublas_sgemm_matrix_multiplication.cu diff --git a/projects/hip/tests/hipify-clang/cuComplex/cuComplex_Julia.cu b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuComplex/cuComplex_Julia.cu similarity index 100% rename from projects/hip/tests/hipify-clang/cuComplex/cuComplex_Julia.cu rename to projects/hip/tests/hipify-clang/unit_tests/libraries/cuComplex/cuComplex_Julia.cu diff --git a/projects/hip/tests/hipify-clang/cuDNN/cudnn_convolution_forward.cu b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu similarity index 100% rename from projects/hip/tests/hipify-clang/cuDNN/cudnn_convolution_forward.cu rename to projects/hip/tests/hipify-clang/unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu diff --git a/projects/hip/tests/hipify-clang/cuDNN/cudnn_softmax.cu b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuDNN/cudnn_softmax.cu similarity index 100% rename from projects/hip/tests/hipify-clang/cuDNN/cudnn_softmax.cu rename to projects/hip/tests/hipify-clang/unit_tests/libraries/cuDNN/cudnn_softmax.cu diff --git a/projects/hip/tests/hipify-clang/cuFFT/simple_cufft.cu b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuFFT/simple_cufft.cu similarity index 100% rename from projects/hip/tests/hipify-clang/cuFFT/simple_cufft.cu rename to projects/hip/tests/hipify-clang/unit_tests/libraries/cuFFT/simple_cufft.cu diff --git a/projects/hip/tests/hipify-clang/cuRAND/benchmark_curand_generate.cpp b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp similarity index 100% rename from projects/hip/tests/hipify-clang/cuRAND/benchmark_curand_generate.cpp rename to projects/hip/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp diff --git a/projects/hip/tests/hipify-clang/cuRAND/benchmark_curand_kernel.cpp b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp similarity index 100% rename from projects/hip/tests/hipify-clang/cuRAND/benchmark_curand_kernel.cpp rename to projects/hip/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp diff --git a/projects/hip/tests/hipify-clang/cuRAND/cmdparser.hpp b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuRAND/cmdparser.hpp similarity index 100% rename from projects/hip/tests/hipify-clang/cuRAND/cmdparser.hpp rename to projects/hip/tests/hipify-clang/unit_tests/libraries/cuRAND/cmdparser.hpp diff --git a/projects/hip/tests/hipify-clang/cuRAND/poisson_api_example.cu b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuRAND/poisson_api_example.cu similarity index 100% rename from projects/hip/tests/hipify-clang/cuRAND/poisson_api_example.cu rename to projects/hip/tests/hipify-clang/unit_tests/libraries/cuRAND/poisson_api_example.cu diff --git a/projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_01.cu b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_01.cu similarity index 100% rename from projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_01.cu rename to projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_01.cu diff --git a/projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_02.cu b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_02.cu similarity index 100% rename from projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_02.cu rename to projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_02.cu diff --git a/projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_03.cu b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu similarity index 100% rename from projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_03.cu rename to projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu diff --git a/projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_04.cu b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_04.cu similarity index 100% rename from projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_04.cu rename to projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_04.cu diff --git a/projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_05.cu b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_05.cu similarity index 100% rename from projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_05.cu rename to projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_05.cu diff --git a/projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_06.cu b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_06.cu similarity index 100% rename from projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_06.cu rename to projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_06.cu diff --git a/projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_07.cu b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_07.cu similarity index 100% rename from projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_07.cu rename to projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_07.cu diff --git a/projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_08.cu b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_08.cu similarity index 100% rename from projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_08.cu rename to projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_08.cu diff --git a/projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_09.cu b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_09.cu similarity index 100% rename from projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_09.cu rename to projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_09.cu diff --git a/projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_10.cu b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu similarity index 100% rename from projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_10.cu rename to projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu diff --git a/projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_11.cu b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_11.cu similarity index 100% rename from projects/hip/tests/hipify-clang/cuSPARSE/cuSPARSE_11.cu rename to projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_11.cu diff --git a/projects/hip/tests/hipify-clang/allocators.cu b/projects/hip/tests/hipify-clang/unit_tests/samples/allocators.cu similarity index 100% rename from projects/hip/tests/hipify-clang/allocators.cu rename to projects/hip/tests/hipify-clang/unit_tests/samples/allocators.cu diff --git a/projects/hip/tests/hipify-clang/axpy.cu b/projects/hip/tests/hipify-clang/unit_tests/samples/axpy.cu similarity index 100% rename from projects/hip/tests/hipify-clang/axpy.cu rename to projects/hip/tests/hipify-clang/unit_tests/samples/axpy.cu diff --git a/projects/hip/tests/hipify-clang/coalescing.cu b/projects/hip/tests/hipify-clang/unit_tests/samples/coalescing.cu similarity index 100% rename from projects/hip/tests/hipify-clang/coalescing.cu rename to projects/hip/tests/hipify-clang/unit_tests/samples/coalescing.cu diff --git a/projects/hip/tests/hipify-clang/cudaRegister.cu b/projects/hip/tests/hipify-clang/unit_tests/samples/cudaRegister.cu similarity index 100% rename from projects/hip/tests/hipify-clang/cudaRegister.cu rename to projects/hip/tests/hipify-clang/unit_tests/samples/cudaRegister.cu diff --git a/projects/hip/tests/hipify-clang/dynamic_shared_memory.cu b/projects/hip/tests/hipify-clang/unit_tests/samples/dynamic_shared_memory.cu similarity index 100% rename from projects/hip/tests/hipify-clang/dynamic_shared_memory.cu rename to projects/hip/tests/hipify-clang/unit_tests/samples/dynamic_shared_memory.cu diff --git a/projects/hip/tests/hipify-clang/intro.cu b/projects/hip/tests/hipify-clang/unit_tests/samples/intro.cu similarity index 100% rename from projects/hip/tests/hipify-clang/intro.cu rename to projects/hip/tests/hipify-clang/unit_tests/samples/intro.cu diff --git a/projects/hip/tests/hipify-clang/square.cu b/projects/hip/tests/hipify-clang/unit_tests/samples/square.cu similarity index 100% rename from projects/hip/tests/hipify-clang/square.cu rename to projects/hip/tests/hipify-clang/unit_tests/samples/square.cu diff --git a/projects/hip/tests/hipify-clang/static_shared_memory.cu b/projects/hip/tests/hipify-clang/unit_tests/samples/static_shared_memory.cu similarity index 100% rename from projects/hip/tests/hipify-clang/static_shared_memory.cu rename to projects/hip/tests/hipify-clang/unit_tests/samples/static_shared_memory.cu diff --git a/projects/hip/tests/hipify-clang/vec_add.cu b/projects/hip/tests/hipify-clang/unit_tests/samples/vec_add.cu similarity index 100% rename from projects/hip/tests/hipify-clang/vec_add.cu rename to projects/hip/tests/hipify-clang/unit_tests/samples/vec_add.cu