diff --git a/CMakeLists.txt b/CMakeLists.txt index 45d7640643..7730950ad5 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -213,6 +213,7 @@ if(HIP_PLATFORM STREQUAL "hcc") src/hip_device.cpp src/hip_error.cpp src/hip_event.cpp + src/hip_fatbin.cpp src/hip_memory.cpp src/hip_peer.cpp src/hip_stream.cpp diff --git a/bin/hipcc b/bin/hipcc index 68e4a96721..a717a0bf92 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -95,15 +95,19 @@ $HIP_VERSION= `$HIP_PATH/bin/hipconfig --version`; ($HIP_VERSION_MAJOR, $HIP_VERSION_MINOR, $HIP_VERSION_PATCH) = split(/\./, $HIP_VERSION); if (defined $HIP_VDI_HOME) { + my $bits = ""; + if (-d "$HIP_VDI_HOME/bin/x86_64") { + $bits = "/x86_64"; + } if (!defined $HIP_CLANG_PATH) { - $HIP_CLANG_PATH = "$HIP_VDI_HOME/bin/x86_64"; + $HIP_CLANG_PATH = "$HIP_VDI_HOME/bin" . $bits; } if (!defined $DEVICE_LIB_PATH) { - $DEVICE_LIB_PATH = "$HIP_VDI_HOME/lib/x86_64/bitcode"; + $DEVICE_LIB_PATH = "$HIP_VDI_HOME/lib" . $bits . "/bitcode"; } $HIP_CLANG_INCLUDE_PATH = "$HIP_VDI_HOME/include/clang"; $HIP_INCLUDE_PATH = "$HIP_VDI_HOME/include"; - $HIP_LIB_PATH = "$HIP_VDI_HOME/lib/x86_64"; + $HIP_LIB_PATH = "$HIP_VDI_HOME/lib" . $bits; } if (defined $HIP_CLANG_PATH) { @@ -334,7 +338,7 @@ foreach $arg (@ARGV) $trimarg = $arg; $trimarg =~ s/^\s+|\s+$//g; # Remive whitespace my $swallowArg = 0; - if ($arg eq '-c') { + if ($arg eq '-c' or $arg eq '--genco') { $compileOnly = 1; $needCXXFLAGS = 1; $needLDFLAGS = 0; @@ -386,6 +390,10 @@ foreach $arg (@ARGV) $swallowArg = 1; } + if (($arg =~ /--genco/) and $HIP_PLATFORM eq 'clang' ) { + $arg = "--cuda-device-only"; + } + if(($trimarg eq '-stdlib=libstdc++') and ($setStdLib eq 0)) { $HIPCXXFLAGS .= $HCC_WA_FLAGS; diff --git a/docs/markdown/CUSPARSE_API_supported_by_HIP.md b/docs/markdown/CUSPARSE_API_supported_by_HIP.md index d3707573a5..be82f35310 100644 --- a/docs/markdown/CUSPARSE_API_supported_by_HIP.md +++ b/docs/markdown/CUSPARSE_API_supported_by_HIP.md @@ -1,6 +1,6 @@ # CUSPARSE API supported by HIP -## **1. CUSPARSE Data types** +## **1. cuSPARSE Data types** | **type** | **CUDA** | **HIP** | |-------------:|---------------------------------------------------------------|------------------------------------------------------------| @@ -83,7 +83,7 @@ | struct |`pruneInfo` | | | typedef |`pruneInfo_t` | | -## **2.cuSPARSE Helper Function Reference** +## **2. cuSPARSE Helper Function Reference** | **CUDA** | **HIP** | |-----------------------------------------------------------|-------------------------------------------------| @@ -129,7 +129,7 @@ |`cusparseCreatePruneInfo` | | |`cusparseDestroyPruneInfo` | | -## **3.cuSPARSE Level 1 Function Reference** +## **3. cuSPARSE Level 1 Function Reference** | **CUDA** | **HIP** | |-----------------------------------------------------------|-------------------------------------------------| @@ -158,7 +158,7 @@ |`cusparseCsctr` | | |`cusparseZsctr` | | -## **4.cuSPARSE Level 2 Function Reference** +## **4. cuSPARSE Level 2 Function Reference** | **CUDA** | **HIP** | |-----------------------------------------------------------|-------------------------------------------------| @@ -236,3 +236,79 @@ |`cusparseDhybsv_solve` | | |`cusparseChybsv_solve` | | |`cusparseZhybsv_solve` | | + +## **5. cuSPARSE Level 3 Function Reference** + +| **CUDA** | **HIP** | +|-----------------------------------------------------------|-------------------------------------------------| +|`cusparseScsrmm` |`hipsparseScsrmm` | +|`cusparseDcsrmm` |`hipsparseDcsrmm` | +|`cusparseCcsrmm` | | +|`cusparseZcsrmm` | | +|`cusparseScsrmm2` |`hipsparseScsrmm2` | +|`cusparseDcsrmm2` |`hipsparseDcsrmm2` | +|`cusparseCcsrmm2` | | +|`cusparseZcsrmm2` | | +|`cusparseScsrsm_analysis` | | +|`cusparseDcsrsm_analysis` | | +|`cusparseCcsrsm_analysis` | | +|`cusparseZcsrsm_analysis` | | +|`cusparseScsrsm_solve` | | +|`cusparseDcsrsm_solve` | | +|`cusparseCcsrsm_solve` | | +|`cusparseZcsrsm_solve` | | +|`cusparseScsrsm2_bufferSizeExt` | | +|`cusparseDcsrsm2_bufferSizeExt` | | +|`cusparseCcsrsm2_bufferSizeExt` | | +|`cusparseZcsrsm2_bufferSizeExt` | | +|`cusparseScsrsm2_analysis` | | +|`cusparseDcsrsm2_analysis` | | +|`cusparseCcsrsm2_analysis` | | +|`cusparseZcsrsm2_analysis` | | +|`cusparseScsrsm2_solve` | | +|`cusparseDcsrsm2_solve` | | +|`cusparseCcsrsm2_solve` | | +|`cusparseZcsrsm2_solve` | | +|`cusparseXcsrsm2_zeroPivot` | | +|`cusparseSbsrmm` | | +|`cusparseDbsrmm` | | +|`cusparseCbsrmm` | | +|`cusparseZbsrmm` | | +|`cusparseSbsrsm2_bufferSize` | | +|`cusparseDbsrsm2_bufferSize` | | +|`cusparseCbsrsm2_bufferSize` | | +|`cusparseZbsrsm2_bufferSize` | | +|`cusparseSbsrsm2_analysis` | | +|`cusparseDbsrsm2_analysis` | | +|`cusparseCbsrsm2_analysis` | | +|`cusparseZbsrsm2_analysis` | | +|`cusparseSbsrsm2_solve` | | +|`cusparseDbsrsm2_solve` | | +|`cusparseCbsrsm2_solve` | | +|`cusparseZbsrsm2_solve` | | +|`cusparseXbsrsm2_zeroPivot` | | +|`cusparseSgemmi` | | +|`cusparseDgemmi` | | +|`cusparseCgemmi` | | +|`cusparseZgemmi` | | + +## **6. cuSPARSE Extra Function Reference** + +| **CUDA** | **HIP** | +|-----------------------------------------------------------|-------------------------------------------------| +|`cusparseXcsrgeamNnz` | | +|`cusparseScsrgeam` | | +|`cusparseDcsrgeam` | | +|`cusparseCcsrgeam` | | +|`cusparseScsrgeam2_bufferSizeExt` | | +|`cusparseDcsrgeam2_bufferSizeExt` | | +|`cusparseCcsrgeam2_bufferSizeExt` | | +|`cusparseZcsrgeam2_bufferSizeExt` | | +|`cusparseXcsrgemmNnz` | | +|`cusparseScsrgemm` | | +|`cusparseDcsrgemm` | | +|`cusparseCcsrgemm` | | +|`cusparseScsrgemm2_bufferSizeExt` | | +|`cusparseDcsrgemm2_bufferSizeExt` | | +|`cusparseCcsrgemm2_bufferSizeExt` | | +|`cusparseZcsrgemm2_bufferSizeExt` | | diff --git a/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp b/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp index 429c8d0ea2..ed46fc5a54 100644 --- a/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp +++ b/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp @@ -174,4 +174,90 @@ const std::map CUDA_SPARSE_FUNCTION_MAP{ {"cusparseDhybsv_solve", {"hipsparseDhybsv_solve", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseChybsv_solve", {"hipsparseChybsv_solve", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZhybsv_solve", {"hipsparseZhybsv_solve", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + // 8. cuSPARSE Level 3 Function Reference + {"cusparseScsrmm", {"hipsparseScsrmm", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseDcsrmm", {"hipsparseDcsrmm", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseCcsrmm", {"hipsparseCcsrmm", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZcsrmm", {"hipsparseZcsrmm", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseScsrmm2", {"hipsparseScsrmm2", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseDcsrmm2", {"hipsparseDcsrmm2", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseCcsrmm2", {"hipsparseCcsrmm2", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZcsrmm2", {"hipsparseZcsrmm2", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseScsrsm_analysis", {"hipsparseScsrsm_analysis", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDcsrsm_analysis", {"hipsparseDcsrsm_analysis", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsrsm_analysis", {"hipsparseCcsrsm_analysis", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZcsrsm_analysis", {"hipsparseZcsrsm_analysis", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseScsrsm_solve", {"hipsparseScsrsm_solve", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDcsrsm_solve", {"hipsparseDcsrsm_solve", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsrsm_solve", {"hipsparseCcsrsm_solve", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZcsrsm_solve", {"hipsparseZcsrsm_solve", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseScsrsm2_bufferSizeExt", {"hipsparseScsrsm2_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDcsrsm2_bufferSizeExt", {"hipsparseDcsrsm2_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsrsm2_bufferSizeExt", {"hipsparseCcsrsm2_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZcsrsm2_bufferSizeExt", {"hipsparseZcsrsm2_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseScsrsm2_analysis", {"hipsparseScsrsm2_analysis", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDcsrsm2_analysis", {"hipsparseDcsrsm2_analysis", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsrsm2_analysis", {"hipsparseCcsrsm2_analysis", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZcsrsm2_analysis", {"hipsparseZcsrsm2_analysis", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseScsrsm2_solve", {"hipsparseScsrsm2_solve", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDcsrsm2_solve", {"hipsparseDcsrsm2_solve", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsrsm2_solve", {"hipsparseCcsrsm2_solve", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZcsrsm2_solve", {"hipsparseZcsrsm2_solve", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseXcsrsm2_zeroPivot", {"hipsparseXcsrsm2_zeroPivot", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseSbsrmm", {"hipsparseSbsrmm", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDbsrmm", {"hipsparseDbsrmm", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCbsrmm", {"hipsparseCbsrmm", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZbsrmm", {"hipsparseZbsrmm", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseSbsrsm2_bufferSize", {"hipsparseCbsrsm2_bufferSize", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDbsrsm2_bufferSize", {"hipsparseDbsrsm2_bufferSize", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCbsrsm2_bufferSize", {"hipsparseCbsrsm2_bufferSize", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZbsrsm2_bufferSize", {"hipsparseZbsrsm2_bufferSize", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseSbsrsm2_analysis", {"hipsparseSbsrsm2_analysis", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDbsrsm2_analysis", {"hipsparseDbsrsm2_analysis", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCbsrsm2_analysis", {"hipsparseCbsrsm2_analysis", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZbsrsm2_analysis", {"hipsparseZbsrsm2_analysis", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseSbsrsm2_solve", {"hipsparseSbsrsm2_solve", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDbsrsm2_solve", {"hipsparseDbsrsm2_solve", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCbsrsm2_solve", {"hipsparseCbsrsm2_solve", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZbsrsm2_solve", {"hipsparseZbsrsm2_solve", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseXbsrsm2_zeroPivot", {"hipsparseXbsrsm2_zeroPivot", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseSgemmi", {"hipsparseSgemmi", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDgemmi", {"hipsparseDgemmi", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCgemmi", {"hipsparseCgemmi", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZgemmi", {"hipsparseZgemmi", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + // 9. cuSPARSE Extra Function Reference + {"cusparseXcsrgeamNnz", {"hipsparseXcsrgeamNnz", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"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}}, + + {"cusparseScsrgeam2_bufferSizeExt", {"hipsparseScsrgeam2_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDcsrgeam2_bufferSizeExt", {"hipsparseDcsrgeam2_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsrgeam2_bufferSizeExt", {"hipsparseCcsrgeam2_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZcsrgeam2_bufferSizeExt", {"hipsparseZcsrgeam2_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseXcsrgemmNnz", {"hipsparseXcsrgemmNnz", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"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}}, + + {"cusparseScsrgemm2_bufferSizeExt", {"hipsparseScsrgemm2_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDcsrgemm2_bufferSizeExt", {"hipsparseDcsrgemm2_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsrgemm2_bufferSizeExt", {"hipsparseCcsrgemm2_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseZcsrgemm2_bufferSizeExt", {"hipsparseZcsrgemm2_bufferSizeExt", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, }; diff --git a/src/hip_clang.cpp b/src/hip_clang.cpp index fe08bbe45c..ef8d456103 100644 --- a/src/hip_clang.cpp +++ b/src/hip_clang.cpp @@ -26,34 +26,9 @@ THE SOFTWARE. #include "hip/hip_runtime.h" #include "hip_hcc_internal.h" +#include "hip_fatbin.h" #include "trace_helper.h" -constexpr unsigned __hipFatMAGIC2 = 0x48495046; // "HIPF" - -#define CLANG_OFFLOAD_BUNDLER_MAGIC "__CLANG_OFFLOAD_BUNDLE__" -#define AMDGCN_AMDHSA_TRIPLE "hip-amdgcn-amd-amdhsa" - -struct __ClangOffloadBundleDesc { - uint64_t offset; - uint64_t size; - uint64_t tripleSize; - const char triple[1]; -}; - -struct __ClangOffloadBundleHeader { - const char magic[sizeof(CLANG_OFFLOAD_BUNDLER_MAGIC) - 1]; - uint64_t numBundles; - __ClangOffloadBundleDesc desc[1]; -}; - -struct __CudaFatBinaryWrapper { - unsigned int magic; - unsigned int version; - __ClangOffloadBundleHeader* binary; - void* unused; -}; - - extern "C" std::vector* __hipRegisterFatBinary(const void* data) { @@ -108,21 +83,13 @@ __hipRegisterFatBinary(const void* data) std::string image{reinterpret_cast( reinterpret_cast(header) + desc->offset), desc->size}; + if (HIP_DUMP_CODE_OBJECT) + __hipDumpCodeObject(image); module->executable = hip_impl::load_executable(image, module->executable, agent); if (module->executable.handle) { modules->at(deviceId) = module; tprintf(DB_FB, "Loaded code object for %s\n", name); - if (HIP_DUMP_CODE_OBJECT) { - char fname[30]; - static std::atomic index; - sprintf(fname, "__hip_dump_code_object%04d.o", index++); - tprintf(DB_FB, "Dump code object %s\n", fname); - std::ofstream ofs; - ofs.open(fname, std::ios::binary); - ofs << image; - ofs.close(); - } } else { fprintf(stderr, "Failed to load code object for %s\n", name); abort(); diff --git a/src/hip_fatbin.cpp b/src/hip_fatbin.cpp new file mode 100644 index 0000000000..8fe7740ed7 --- /dev/null +++ b/src/hip_fatbin.cpp @@ -0,0 +1,91 @@ +/* +Copyright (c) 2018 - present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#include +#include + +#include "hip_fatbin.h" +#include "hip/hip_runtime.h" +#include "hip_hcc_internal.h" +#include "trace_helper.h" + +void __hipDumpCodeObject(const std::string& image) { + char fname[30]; + static std::atomic index; + sprintf(fname, "__hip_dump_code_object%04d.o", index++); + tprintf(DB_FB, "Dump code object %s\n", fname); + std::ofstream ofs; + ofs.open(fname, std::ios::binary); + ofs << image; + ofs.close(); +} + +// Returns a pointer to the code object in the fatbin. The pointer should not +// be freed. +const void* __hipExtractCodeObjectFromFatBinary(const void* data, + const char* agent_name) +{ + HIP_INIT(); + + tprintf(DB_FB, "Enter __hipExtractCodeObjectFromFatBinary(%p, \"%s\")\n", + data, agent_name); + + const __ClangOffloadBundleHeader* header + = reinterpret_cast(data); + std::string magic(reinterpret_cast(header), + sizeof(CLANG_OFFLOAD_BUNDLER_MAGIC) - 1); + if (magic.compare(CLANG_OFFLOAD_BUNDLER_MAGIC)) { + return nullptr; + } + + const __ClangOffloadBundleDesc* desc = &header->desc[0]; + for (uint64_t i = 0; i < header->numBundles; ++i, + desc = reinterpret_cast( + reinterpret_cast(&desc->triple[0]) + desc->tripleSize)) { + + std::string triple{&desc->triple[0], sizeof(AMDGCN_AMDHSA_TRIPLE) - 1}; + if (triple.compare(AMDGCN_AMDHSA_TRIPLE)) + continue; + + std::string target{&desc->triple[sizeof(AMDGCN_AMDHSA_TRIPLE)], + desc->tripleSize - sizeof(AMDGCN_AMDHSA_TRIPLE)}; + tprintf(DB_FB, "Found hip-clang bundle for %s\n", target.c_str()); + if (target.compare(agent_name)) { + continue; + } + + auto *codeobj = reinterpret_cast( + reinterpret_cast(header) + desc->offset); + if (HIP_DUMP_CODE_OBJECT) + __hipDumpCodeObject(std::string{codeobj, desc->size}); + + tprintf(DB_FB, "__hipExtractCodeObjectFromFatBinary succeeds and returns %p\n", + codeobj); + return codeobj; + } + + // hipcc --genco for HCC generates fat binaries with different triple strings. + // It will reach here and return a null pointer. The fat binary itself will + // be handled in a different place. + tprintf(DB_FB, "No hip-clang device code bundle for %s\n", agent_name); + return nullptr; +} + diff --git a/src/hip_fatbin.h b/src/hip_fatbin.h new file mode 100644 index 0000000000..7b4a063c68 --- /dev/null +++ b/src/hip_fatbin.h @@ -0,0 +1,58 @@ +/* +Copyright (c) 2018 - present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#ifndef HIP_SRC_HIP_FATBIN_H +#define HIP_SRC_HIP_FATBIN_H + +#include "hip/hip_runtime.h" +#include "hip_hcc_internal.h" + +// hip-clang fatbin format +constexpr unsigned __hipFatMAGIC2 = 0x48495046; // "HIPF" + +#define CLANG_OFFLOAD_BUNDLER_MAGIC "__CLANG_OFFLOAD_BUNDLE__" +#define AMDGCN_AMDHSA_TRIPLE "hip-amdgcn-amd-amdhsa" + +struct __ClangOffloadBundleDesc { + uint64_t offset; + uint64_t size; + uint64_t tripleSize; + const char triple[1]; +}; + +struct __ClangOffloadBundleHeader { + const char magic[sizeof(CLANG_OFFLOAD_BUNDLER_MAGIC) - 1]; + uint64_t numBundles; + __ClangOffloadBundleDesc desc[1]; +}; + +struct __CudaFatBinaryWrapper { + unsigned int magic; + unsigned int version; + __ClangOffloadBundleHeader* binary; + void* unused; +}; + +const void* __hipExtractCodeObjectFromFatBinary(const void* data, + const char* agent_name); +void __hipDumpCodeObject(const std::string& image); + +#endif // HIP_SRC_HIP_FATBIN_H diff --git a/src/hip_module.cpp b/src/hip_module.cpp index 2f1e4ee322..657fb06b5e 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -48,6 +48,7 @@ THE SOFTWARE. #include #include #include "../include/hip/hcc_detail/code_object_bundle.hpp" +#include "hip_fatbin.h" // TODO Use Pool APIs from HCC to get memory regions. using namespace ELFIO; @@ -556,6 +557,12 @@ hipError_t ihipModuleLoadData(hipModule_t* module, const void* image) { auto ctx = ihipGetTlsDefaultCtx(); if (!ctx) return hipErrorInvalidContext; + // try extracting code object from image as fatbin. + char name[64] = {}; + hsa_agent_get_info(this_agent(), HSA_AGENT_INFO_NAME, name); + if (auto *code_obj = __hipExtractCodeObjectFromFatBinary(image, name)) + image = code_obj; + hsa_executable_create_alt(HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, nullptr, &(*module)->executable); diff --git a/tests/hipify-clang/cuSPARSE/cuSPARSE_01.cu b/tests/hipify-clang/cuSPARSE/cuSPARSE_01.cu index 5ef7c188ee..df2499c041 100644 --- a/tests/hipify-clang/cuSPARSE/cuSPARSE_01.cu +++ b/tests/hipify-clang/cuSPARSE/cuSPARSE_01.cu @@ -261,7 +261,7 @@ int main(){ return 2; } /* exercise Level 1 routines (scatter vector elements) */ - // TODO: status= hipsparseDsctr(handle, nnz_vector, xVal, xInd, + // CHECK: status= hipsparseDsctr(handle, nnz_vector, xVal, xInd, // CHECK: &y[n], HIPSPARSE_INDEX_BASE_ZERO); status= cusparseDsctr(handle, nnz_vector, xVal, xInd, &y[n], CUSPARSE_INDEX_BASE_ZERO); @@ -299,7 +299,7 @@ int main(){ CLEANUP("Memset on Device failed"); return 1; } - // TODO: status= hipsparseDcsrmm(handle, HIPSPARSE_OPERATION_NON_TRANSPOSE, n, 2, n, + // CHECK: status= hipsparseDcsrmm(handle, HIPSPARSE_OPERATION_NON_TRANSPOSE, n, 2, n, status= cusparseDcsrmm(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, n, 2, n, nnz, &dfive, descr, cooVal, csrRowPtr, cooColIndex, y, n, &dzero, z, n+1); diff --git a/tests/hipify-clang/cuSPARSE/cuSPARSE_02.cu b/tests/hipify-clang/cuSPARSE/cuSPARSE_02.cu index ab2defefe7..de7629367f 100644 --- a/tests/hipify-clang/cuSPARSE/cuSPARSE_02.cu +++ b/tests/hipify-clang/cuSPARSE/cuSPARSE_02.cu @@ -116,7 +116,7 @@ int main(int argc, char*argv[]) assert(CUSPARSE_STATUS_SUCCESS == cusparseStat); // step 2: configuration of matrix A - // cusparseStat = hipsparseCreateMatDescr(&descrA); + // CHECK: cusparseStat = hipsparseCreateMatDescr(&descrA); cusparseStat = cusparseCreateMatDescr(&descrA); // assert(HIPSPARSE_STATUS_SUCCESS == cusparseStat); assert(CUSPARSE_STATUS_SUCCESS == cusparseStat); diff --git a/tests/hipify-clang/cuSPARSE/cuSPARSE_03.cu b/tests/hipify-clang/cuSPARSE/cuSPARSE_03.cu new file mode 100644 index 0000000000..ef52072576 --- /dev/null +++ b/tests/hipify-clang/cuSPARSE/cuSPARSE_03.cu @@ -0,0 +1,229 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args +#include +#include +#include +// CHECK: #include +#include +// CHECK: #include "hipsparse.h" +#include "cusparse.h" + +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: hipsparseStatus_t status = HIPSPARSE_STATUS_SUCCESS; + cusparseStatus_t status = CUSPARSE_STATUS_SUCCESS; + // CHECK: hipError_t cudaStat1 = hipSuccess; + // CHECK: hipError_t cudaStat2 = hipSuccess; + // CHECK: hipError_t cudaStat3 = hipSuccess; + // CHECK: hipError_t cudaStat4 = hipSuccess; + // CHECK: hipError_t cudaStat5 = hipSuccess; + // CHECK: hipError_t cudaStat6 = hipSuccess; + cudaError_t cudaStat1 = cudaSuccess; + cudaError_t cudaStat2 = cudaSuccess; + cudaError_t cudaStat3 = cudaSuccess; + cudaError_t cudaStat4 = cudaSuccess; + cudaError_t cudaStat5 = cudaSuccess; + cudaError_t cudaStat6 = cudaSuccess; + + /* + * A is a 3x3 sparse matrix + * | 1 2 0 | + * A = | 0 5 0 | + * | 0 8 0 | + */ + const int m = 3; + const int n = 3; + const int nnz = 4; + +#if 0 + /* index starts at 0 */ + int h_cooRows[nnz] = { 2, 1, 0, 0 }; + int h_cooCols[nnz] = { 1, 1, 0, 1 }; +#else + /* index starts at -2 */ + int h_cooRows[nnz] = { 0, -1, -2, -2 }; + int h_cooCols[nnz] = { -1, -1, -2, -1 }; +#endif + double h_cooVals[nnz] = { 8.0, 5.0, 1.0, 2.0 }; + int h_P[nnz]; + + int *d_cooRows = NULL; + int *d_cooCols = NULL; + int *d_P = NULL; + double *d_cooVals = NULL; + double *d_cooVals_sorted = NULL; + size_t pBufferSizeInBytes = 0; + void *pBuffer = NULL; + + printf("m = %d, n = %d, nnz=%d \n", m, n, nnz); + + /* 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); + // CHECK: status = hipsparseSetStream(handle, stream); + status = cusparseSetStream(handle, stream); + // CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); + assert(CUSPARSE_STATUS_SUCCESS == status); + + /* step 2: allocate buffer */ + // TODO: status = hipsparseXcoosort_bufferSizeExt( + status = cusparseXcoosort_bufferSizeExt( + handle, + m, + n, + nnz, + d_cooRows, + d_cooCols, + &pBufferSizeInBytes + ); + // CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); + assert(CUSPARSE_STATUS_SUCCESS == status); + + printf("pBufferSizeInBytes = %lld bytes \n", (long long)pBufferSizeInBytes); + + // CHECK: cudaStat1 = hipMalloc(&d_cooRows, sizeof(int)*nnz); + cudaStat1 = cudaMalloc(&d_cooRows, sizeof(int)*nnz); + // CHECK: cudaStat2 = hipMalloc(&d_cooCols, sizeof(int)*nnz); + cudaStat2 = cudaMalloc(&d_cooCols, sizeof(int)*nnz); + // CHECK: cudaStat3 = hipMalloc(&d_P, sizeof(int)*nnz); + cudaStat3 = cudaMalloc(&d_P, sizeof(int)*nnz); + // CHECK: cudaStat4 = hipMalloc(&d_cooVals, sizeof(double)*nnz); + cudaStat4 = cudaMalloc(&d_cooVals, sizeof(double)*nnz); + // CHECK: cudaStat5 = hipMalloc(&d_cooVals_sorted, sizeof(double)*nnz); + cudaStat5 = cudaMalloc(&d_cooVals_sorted, sizeof(double)*nnz); + // CHECK: cudaStat6 = hipMalloc(&pBuffer, sizeof(char)* pBufferSizeInBytes); + cudaStat6 = cudaMalloc(&pBuffer, sizeof(char)* pBufferSizeInBytes); + + // CHECK: assert(hipSuccess == cudaStat1); + // CHECK: assert(hipSuccess == cudaStat2); + // CHECK: assert(hipSuccess == cudaStat3); + // CHECK: assert(hipSuccess == cudaStat4); + // CHECK: assert(hipSuccess == cudaStat5); + // CHECK: assert(hipSuccess == cudaStat6); + assert(cudaSuccess == cudaStat1); + assert(cudaSuccess == cudaStat2); + assert(cudaSuccess == cudaStat3); + assert(cudaSuccess == cudaStat4); + assert(cudaSuccess == cudaStat5); + assert(cudaSuccess == cudaStat6); + + // CHECK: cudaStat1 = hipMemcpy(d_cooRows, h_cooRows, sizeof(int)*nnz, hipMemcpyHostToDevice); + cudaStat1 = cudaMemcpy(d_cooRows, h_cooRows, sizeof(int)*nnz, cudaMemcpyHostToDevice); + // CHECK: cudaStat2 = hipMemcpy(d_cooCols, h_cooCols, sizeof(int)*nnz, hipMemcpyHostToDevice); + cudaStat2 = cudaMemcpy(d_cooCols, h_cooCols, sizeof(int)*nnz, cudaMemcpyHostToDevice); + // CHECK: cudaStat3 = hipMemcpy(d_cooVals, h_cooVals, sizeof(double)*nnz, hipMemcpyHostToDevice); + cudaStat3 = cudaMemcpy(d_cooVals, h_cooVals, sizeof(double)*nnz, cudaMemcpyHostToDevice); + // CHECK: cudaStat4 = hipDeviceSynchronize(); + cudaStat4 = cudaDeviceSynchronize(); + + // CHECK: assert(hipSuccess == cudaStat1); + // CHECK: assert(hipSuccess == cudaStat2); + // CHECK: assert(hipSuccess == cudaStat3); + // CHECK: assert(hipSuccess == cudaStat4); + assert(cudaSuccess == cudaStat1); + assert(cudaSuccess == cudaStat2); + assert(cudaSuccess == cudaStat3); + assert(cudaSuccess == cudaStat4); + + /* step 3: setup permutation vector P to identity */ + // TODO: status = hipsparseCreateIdentityPermutation( + status = cusparseCreateIdentityPermutation( + handle, + nnz, + d_P); + // CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); + assert(CUSPARSE_STATUS_SUCCESS == status); + + /* step 4: sort COO format by Row */ + // TODO: status = hipsparseXcoosortByRow( + status = cusparseXcoosortByRow( + handle, + m, + n, + nnz, + d_cooRows, + d_cooCols, + d_P, + pBuffer + ); + // CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); + assert(CUSPARSE_STATUS_SUCCESS == status); + + /* step 5: gather sorted cooVals */ + // CHECK: status = hipsparseDgthr( + // CHECK: HIPSPARSE_INDEX_BASE_ZERO + status = cusparseDgthr( + handle, + nnz, + d_cooVals, + d_cooVals_sorted, + d_P, + CUSPARSE_INDEX_BASE_ZERO + ); + // CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); + assert(CUSPARSE_STATUS_SUCCESS == status); + /* wait until the computation is done */ + // CHECK: cudaStat1 = hipDeviceSynchronize(); + cudaStat1 = cudaDeviceSynchronize(); + // CHECK: cudaStat2 = hipMemcpy(h_cooRows, d_cooRows, sizeof(int)*nnz, hipMemcpyDeviceToHost); + cudaStat2 = cudaMemcpy(h_cooRows, d_cooRows, sizeof(int)*nnz, cudaMemcpyDeviceToHost); + // CHECK: cudaStat3 = hipMemcpy(h_cooCols, d_cooCols, sizeof(int)*nnz, hipMemcpyDeviceToHost); + cudaStat3 = cudaMemcpy(h_cooCols, d_cooCols, sizeof(int)*nnz, cudaMemcpyDeviceToHost); + // CHECK: cudaStat4 = hipMemcpy(h_P, d_P, sizeof(int)*nnz, hipMemcpyDeviceToHost); + cudaStat4 = cudaMemcpy(h_P, d_P, sizeof(int)*nnz, cudaMemcpyDeviceToHost); + // CHECK: cudaStat5 = hipMemcpy(h_cooVals, d_cooVals_sorted, sizeof(double)*nnz, hipMemcpyDeviceToHost); + cudaStat5 = cudaMemcpy(h_cooVals, d_cooVals_sorted, sizeof(double)*nnz, cudaMemcpyDeviceToHost); + // CHECK: cudaStat6 = hipDeviceSynchronize(); + cudaStat6 = cudaDeviceSynchronize(); + // CHECK: assert(hipSuccess == cudaStat1); + // CHECK: assert(hipSuccess == cudaStat2); + // CHECK: assert(hipSuccess == cudaStat3); + // CHECK: assert(hipSuccess == cudaStat4); + // CHECK: assert(hipSuccess == cudaStat5); + // CHECK: assert(hipSuccess == cudaStat6); + assert(cudaSuccess == cudaStat1); + assert(cudaSuccess == cudaStat2); + assert(cudaSuccess == cudaStat3); + assert(cudaSuccess == cudaStat4); + assert(cudaSuccess == cudaStat5); + assert(cudaSuccess == cudaStat6); + + printf("sorted coo: \n"); + for (int j = 0; j < nnz; j++) { + printf("(%d, %d, %f) \n", h_cooRows[j], h_cooCols[j], h_cooVals[j]); + } + + for (int j = 0; j < nnz; j++) { + printf("P[%d] = %d \n", j, h_P[j]); + } + + /* free resources */ + // CHECK: if (d_cooRows) hipFree(d_cooRows); + if (d_cooRows) cudaFree(d_cooRows); + // CHECK: if (d_cooCols) hipFree(d_cooCols); + if (d_cooCols) cudaFree(d_cooCols); + // CHECK: if (d_P) hipFree(d_P); + if (d_P) cudaFree(d_P); + // CHECK: if (d_cooVals) hipFree(d_cooVals); + if (d_cooVals) cudaFree(d_cooVals); + // CHECK: if (d_cooVals_sorted) hipFree(d_cooVals_sorted); + if (d_cooVals_sorted) cudaFree(d_cooVals_sorted); + // CHECK: if (pBuffer) hipFree(pBuffer); + if (pBuffer) cudaFree(pBuffer); + // if (handle) hipsparseDestroy(handle); + if (handle) cusparseDestroy(handle); + // CHECK: if (stream) hipStreamDestroy(stream); + if (stream) cudaStreamDestroy(stream); + // CHECK: hipDeviceReset(); + cudaDeviceReset(); + return 0; +} diff --git a/tests/hipify-clang/cuSPARSE/cuSPARSE_04.cu b/tests/hipify-clang/cuSPARSE/cuSPARSE_04.cu new file mode 100644 index 0000000000..32760f4fa7 --- /dev/null +++ b/tests/hipify-clang/cuSPARSE/cuSPARSE_04.cu @@ -0,0 +1,261 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args +#include +#include +#include +// CHECK: #include +#include +// CHECK: #include +#include + +void printMatrix(int m, int n, const float*A, int lda, const char* name) +{ + for (int row = 0; row < m; row++) { + for (int col = 0; col < n; col++) { + float Areg = A[row + col * lda]; + printf("%s(%d,%d) = %f\n", name, row + 1, col + 1, Areg); + } + } +} + +void printCsr( + int m, + int n, + int nnz, + // CHECK: const hipsparseMatDescr_t descrA, + const cusparseMatDescr_t descrA, + const float *csrValA, + const int *csrRowPtrA, + const int *csrColIndA, + const char* name) +{ + // CHECK: const int base = (hipsparseGetMatIndexBase(descrA) != HIPSPARSE_INDEX_BASE_ONE) ? 0 : 1; + const int base = (cusparseGetMatIndexBase(descrA) != CUSPARSE_INDEX_BASE_ONE) ? 0 : 1; + + printf("matrix %s is %d-by-%d, nnz=%d, base=%d\n", name, m, n, nnz, base); + for (int row = 0; row < m; row++) { + const int start = csrRowPtrA[row] - base; + const int end = csrRowPtrA[row + 1] - base; + for (int colidx = start; colidx < end; colidx++) { + const int col = csrColIndA[colidx] - base; + const float Areg = csrValA[colidx]; + printf("%s(%d,%d) = %f\n", name, row + 1, col + 1, Areg); + } + } +} + +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 descrC = NULL; + cusparseMatDescr_t descrC = NULL; + // CHECK: hipsparseStatus_t status = HIPSPARSE_STATUS_SUCCESS; + cusparseStatus_t status = CUSPARSE_STATUS_SUCCESS; + // CHECK: hipError_t cudaStat1 = hipSuccess; + // CHECK: hipError_t cudaStat2 = hipSuccess; + // CHECK: hipError_t cudaStat3 = hipSuccess; + // CHECK: hipError_t cudaStat4 = hipSuccess; + // CHECK: hipError_t cudaStat5 = hipSuccess; + cudaError_t cudaStat1 = cudaSuccess; + cudaError_t cudaStat2 = cudaSuccess; + cudaError_t cudaStat3 = cudaSuccess; + cudaError_t cudaStat4 = cudaSuccess; + cudaError_t cudaStat5 = cudaSuccess; + const int m = 4; + const int n = 4; + const int lda = m; + /* + * | 1 0 2 -3 | + * | 0 4 0 0 | + * A = | 5 0 6 7 | + * | 0 8 0 9 | + * + */ + const float A[lda*n] = { 1, 0, 5, 0, 0, 4, 0, 8, 2, 0, 6, 0, -3, 0, 7, 9 }; + int* csrRowPtrC = NULL; + int* csrColIndC = NULL; + float* csrValC = NULL; + + float *d_A = NULL; + int *d_csrRowPtrC = NULL; + int *d_csrColIndC = NULL; + float *d_csrValC = NULL; + + size_t lworkInBytes = 0; + char *d_work = NULL; + + int nnzC = 0; + + float threshold = 4.1; /* remove Aij <= 4.1 */ +// float threshold = 0; /* remove zeros */ + + printf("example of pruneDense2csr \n"); + + printf("prune |A(i,j)| <= threshold \n"); + printf("threshold = %E \n", threshold); + + printMatrix(m, n, A, lda, "A"); + + /* 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); + // CHECK: status = hipsparseSetStream(handle, stream); + status = cusparseSetStream(handle, stream); + // CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); + assert(CUSPARSE_STATUS_SUCCESS == status); + + /* step 2: configuration of matrix C */ + // CHECK: status = hipsparseCreateMatDescr(&descrC); + status = cusparseCreateMatDescr(&descrC); + // CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); + assert(CUSPARSE_STATUS_SUCCESS == status); + // CHECK: hipsparseSetMatIndexBase(descrC, HIPSPARSE_INDEX_BASE_ZERO); + cusparseSetMatIndexBase(descrC, CUSPARSE_INDEX_BASE_ZERO); + // CHECK: hipsparseSetMatType(descrC, HIPSPARSE_MATRIX_TYPE_GENERAL); + cusparseSetMatType(descrC, CUSPARSE_MATRIX_TYPE_GENERAL); + // CHECK: cudaStat1 = hipMalloc((void**)&d_A, sizeof(float)*lda*n); + cudaStat1 = cudaMalloc((void**)&d_A, sizeof(float)*lda*n); + // CHECK: cudaStat2 = hipMalloc((void**)&d_csrRowPtrC, sizeof(int)*(m + 1)); + cudaStat2 = cudaMalloc((void**)&d_csrRowPtrC, sizeof(int)*(m + 1)); + // CHECK: assert(hipSuccess == cudaStat1); + // CHECK: assert(hipSuccess == cudaStat2); + assert(cudaSuccess == cudaStat1); + assert(cudaSuccess == cudaStat2); + + /* step 3: query workspace */ + // CHECK: cudaStat1 = hipMemcpy(d_A, A, sizeof(float)*lda*n, hipMemcpyHostToDevice); + cudaStat1 = cudaMemcpy(d_A, A, sizeof(float)*lda*n, cudaMemcpyHostToDevice); + // CHECK: assert(hipSuccess == cudaStat1); + assert(cudaSuccess == cudaStat1); + // TODO: status = hipsparseSpruneDense2csr_bufferSizeExt( + status = cusparseSpruneDense2csr_bufferSizeExt( + handle, + m, + n, + d_A, + lda, + &threshold, + descrC, + d_csrValC, + d_csrRowPtrC, + d_csrColIndC, + &lworkInBytes); + // CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); + assert(CUSPARSE_STATUS_SUCCESS == status); + + printf("lworkInBytes (prune) = %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: compute csrRowPtrC and nnzC */ + // TODO: status = hipsparseSpruneDense2csrNnz( + status = cusparseSpruneDense2csrNnz( + handle, + m, + n, + d_A, + lda, + &threshold, + descrC, + d_csrRowPtrC, + &nnzC, /* host */ + 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); + + printf("nnzC = %d\n", nnzC); + if (0 == nnzC) { + printf("C is empty \n"); + return 0; + } + + /* step 5: compute csrColIndC and csrValC */ + // CHECK: cudaStat1 = hipMalloc((void**)&d_csrColIndC, sizeof(int) * nnzC); + cudaStat1 = cudaMalloc((void**)&d_csrColIndC, sizeof(int) * nnzC); + // CHECK: cudaStat2 = hipMalloc((void**)&d_csrValC, sizeof(float) * nnzC); + cudaStat2 = cudaMalloc((void**)&d_csrValC, sizeof(float) * nnzC); + // CHECK: assert(hipSuccess == cudaStat1); + // CHECK: assert(hipSuccess == cudaStat2); + assert(cudaSuccess == cudaStat1); + assert(cudaSuccess == cudaStat2); + // TODO: status = hipsparseSpruneDense2csr( + status = cusparseSpruneDense2csr( + handle, + m, + n, + d_A, + lda, + &threshold, + descrC, + d_csrValC, + d_csrRowPtrC, + d_csrColIndC, + 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: output C */ + csrRowPtrC = (int*)malloc(sizeof(int)*(m + 1)); + csrColIndC = (int*)malloc(sizeof(int)*nnzC); + csrValC = (float*)malloc(sizeof(float)*nnzC); + assert(NULL != csrRowPtrC); + assert(NULL != csrColIndC); + assert(NULL != csrValC); + // CHECK: cudaStat1 = hipMemcpy(csrRowPtrC, d_csrRowPtrC, sizeof(int)*(m + 1), hipMemcpyDeviceToHost); + cudaStat1 = cudaMemcpy(csrRowPtrC, d_csrRowPtrC, sizeof(int)*(m + 1), cudaMemcpyDeviceToHost); + // CHECK: cudaStat2 = hipMemcpy(csrColIndC, d_csrColIndC, sizeof(int)*nnzC, hipMemcpyDeviceToHost); + cudaStat2 = cudaMemcpy(csrColIndC, d_csrColIndC, sizeof(int)*nnzC, cudaMemcpyDeviceToHost); + // CHECK: cudaStat3 = hipMemcpy(csrValC, d_csrValC, sizeof(float)*nnzC, hipMemcpyDeviceToHost); + cudaStat3 = cudaMemcpy(csrValC, d_csrValC, sizeof(float)*nnzC, cudaMemcpyDeviceToHost); + // CHECK: assert(hipSuccess == cudaStat1); + // CHECK: assert(hipSuccess == cudaStat2); + // CHECK: assert(hipSuccess == cudaStat3); + assert(cudaSuccess == cudaStat1); + assert(cudaSuccess == cudaStat2); + assert(cudaSuccess == cudaStat3); + + printCsr(m, n, nnzC, descrC, csrValC, csrRowPtrC, csrColIndC, "C"); + + /* free resources */ + // CHECK: if (d_A) hipFree(d_A); + if (d_A) cudaFree(d_A); + // CHECK: if (d_csrRowPtrC) hipFree(d_csrRowPtrC); + if (d_csrRowPtrC) cudaFree(d_csrRowPtrC); + // CHECK: if (d_csrColIndC) hipFree(d_csrColIndC); + if (d_csrColIndC) cudaFree(d_csrColIndC); + // CHECK: if (d_csrValC) hipFree(d_csrValC); + if (d_csrValC) cudaFree(d_csrValC); + + if (csrRowPtrC) free(csrRowPtrC); + if (csrColIndC) free(csrColIndC); + if (csrValC) free(csrValC); + // CHECK: if (handle) hipsparseDestroy(handle); + if (handle) cusparseDestroy(handle); + // CHECK: if (stream) hipStreamDestroy(stream); + if (stream) cudaStreamDestroy(stream); + // CHECK: if (descrC) hipsparseDestroyMatDescr(descrC); + if (descrC) cusparseDestroyMatDescr(descrC); + // CHECK: hipDeviceReset(); + cudaDeviceReset(); + return 0; +}