Merge 'master' into 'amd-master'
Change-Id: I90191bfc6db938b1ab25d0878756fbddf9bcc1bf
Bu işleme şunda yer alıyor:
Jenkins
@@ -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
|
||||
|
||||
+12
-4
@@ -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;
|
||||
|
||||
@@ -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` | |
|
||||
|
||||
@@ -174,4 +174,90 @@ const std::map<llvm::StringRef, hipCounter> 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}},
|
||||
};
|
||||
|
||||
@@ -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<hipModule_t>*
|
||||
__hipRegisterFatBinary(const void* data)
|
||||
{
|
||||
@@ -108,21 +83,13 @@ __hipRegisterFatBinary(const void* data)
|
||||
|
||||
std::string image{reinterpret_cast<const char*>(
|
||||
reinterpret_cast<uintptr_t>(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<int> 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();
|
||||
|
||||
@@ -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 <string>
|
||||
#include <fstream>
|
||||
|
||||
#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<int> 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<const __ClangOffloadBundleHeader*>(data);
|
||||
std::string magic(reinterpret_cast<const char*>(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<const __ClangOffloadBundleDesc*>(
|
||||
reinterpret_cast<uintptr_t>(&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<const char*>(
|
||||
reinterpret_cast<uintptr_t>(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;
|
||||
}
|
||||
|
||||
@@ -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
|
||||
@@ -48,6 +48,7 @@ THE SOFTWARE.
|
||||
#include <utility>
|
||||
#include <vector>
|
||||
#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);
|
||||
|
||||
|
||||
@@ -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);
|
||||
|
||||
@@ -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);
|
||||
|
||||
@@ -0,0 +1,229 @@
|
||||
// RUN: %run_test hipify "%s" "%t" %cuda_args
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#include <assert.h>
|
||||
// CHECK: #include <hip/hip_runtime.h>
|
||||
#include <cuda_runtime.h>
|
||||
// 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;
|
||||
}
|
||||
@@ -0,0 +1,261 @@
|
||||
// RUN: %run_test hipify "%s" "%t" %cuda_args
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#include <assert.h>
|
||||
// CHECK: #include <hip/hip_runtime.h>
|
||||
#include <cuda_runtime.h>
|
||||
// CHECK: #include <hipsparse.h>
|
||||
#include <cusparse.h>
|
||||
|
||||
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;
|
||||
}
|
||||
Yeni konuda referans
Bir kullanıcı engelle