Merge 'master' into 'amd-master'
Change-Id: I4cb9bd9cacaf55727096fd55d05c6e17ffd4e1b0
[ROCm/hip commit: 294fe43ef0]
This commit is contained in:
Jenkins
Vendored
@@ -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')
|
||||
|
||||
@@ -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` |
|
||||
|`cusparseXcscsort_bufferSizeExt` | |
|
||||
|`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` | |
|
||||
|
||||
@@ -245,6 +245,7 @@ const std::map<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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}},
|
||||
|
||||
{"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}},
|
||||
{"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}},
|
||||
};
|
||||
|
||||
@@ -33,7 +33,6 @@ THE SOFTWARE.
|
||||
|
||||
#include <cstddef>
|
||||
#include <cstdint>
|
||||
#include <cstring>
|
||||
#include <functional>
|
||||
#include <iostream>
|
||||
#include <mutex>
|
||||
@@ -57,9 +56,7 @@ template <
|
||||
typename... Ts,
|
||||
typename std::enable_if<n == sizeof...(Ts)>::type* = nullptr>
|
||||
inline std::vector<std::uint8_t> make_kernarg(
|
||||
const std::tuple<Ts...>&,
|
||||
const std::vector<std::pair<std::size_t, std::size_t>>&,
|
||||
std::vector<std::uint8_t> kernarg) {
|
||||
std::vector<std::uint8_t> kernarg, const std::tuple<Ts...>&) {
|
||||
return kernarg;
|
||||
}
|
||||
|
||||
@@ -68,9 +65,7 @@ template <
|
||||
typename... Ts,
|
||||
typename std::enable_if<n != sizeof...(Ts)>::type* = nullptr>
|
||||
inline std::vector<std::uint8_t> make_kernarg(
|
||||
const std::tuple<Ts...>& formals,
|
||||
const std::vector<std::pair<std::size_t, std::size_t>>& size_align,
|
||||
std::vector<std::uint8_t> kernarg) {
|
||||
std::vector<std::uint8_t> kernarg, const std::tuple<Ts...>& formals) {
|
||||
using T = typename std::tuple_element<n, std::tuple<Ts...>>::type;
|
||||
|
||||
static_assert(
|
||||
@@ -85,44 +80,24 @@ inline std::vector<std::uint8_t> 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<n>(formals),
|
||||
size_align[n].first);
|
||||
new (kernarg.data() + kernarg.size() - sizeof(T)) T{std::get<n>(formals)};
|
||||
|
||||
return make_kernarg<n + 1>(formals, size_align, std::move(kernarg));
|
||||
return make_kernarg<n + 1>(std::move(kernarg), formals);
|
||||
}
|
||||
|
||||
template <typename... Formals, typename... Actuals>
|
||||
inline std::vector<std::uint8_t> make_kernarg(
|
||||
void (*kernel)(Formals...), std::tuple<Actuals...> actuals) {
|
||||
void (*)(Formals...), std::tuple<Actuals...> actuals) {
|
||||
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<std::uintptr_t>(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<Formals...> to_formals{std::move(actuals)};
|
||||
std::vector<std::uint8_t> 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,
|
||||
|
||||
@@ -99,8 +99,6 @@ const std::unordered_map<std::uintptr_t, std::vector<std::pair<hsa_agent_t, Kern
|
||||
functions(bool rebuild = false);
|
||||
const std::unordered_map<std::uintptr_t, std::string>& function_names(bool rebuild = false);
|
||||
std::unordered_map<std::string, void*>& globals(bool rebuild = false);
|
||||
std::unordered_map<
|
||||
std::string, std::vector<std::pair<std::size_t, std::size_t>>>& kernargs();
|
||||
|
||||
hsa_executable_t load_executable(const std::string& file, hsa_executable_t executable,
|
||||
hsa_agent_t agent);
|
||||
|
||||
@@ -312,8 +312,8 @@ const unordered_map<string, vector<hsa_executable_symbol_t>>& 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,90 +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<pair<size_t, size_t>>& size_align) {
|
||||
if (f == l) return f;
|
||||
if (!size_align.empty()) return l;
|
||||
|
||||
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<string, vector<pair<size_t, size_t>>>& 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<char*>(desc), static_cast<char*>(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_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]);
|
||||
} while (true);
|
||||
}
|
||||
}
|
||||
} // namespace
|
||||
|
||||
namespace hip_impl {
|
||||
@@ -585,27 +501,6 @@ unordered_map<string, void*>& globals(bool rebuild) {
|
||||
return r;
|
||||
}
|
||||
|
||||
unordered_map<string, vector<pair<size_t, size_t>>>& kernargs() {
|
||||
static unordered_map<string, vector<pair<size_t, size_t>>> r;
|
||||
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()}};
|
||||
|
||||
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;
|
||||
|
||||
@@ -23,6 +23,13 @@ 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')
|
||||
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
|
||||
|
||||
|
||||
@@ -0,0 +1,413 @@
|
||||
// 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>
|
||||
// CHECK: #include <hipblas.h>
|
||||
#include <cublas_v2.h>
|
||||
|
||||
// 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
|
||||
// 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;
|
||||
}
|
||||
@@ -0,0 +1,414 @@
|
||||
// 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>
|
||||
// CHECK: #include <hipblas.h>
|
||||
#include <cublas_v2.h>
|
||||
|
||||
// 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;
|
||||
}
|
||||
@@ -0,0 +1,507 @@
|
||||
// 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>
|
||||
// CHECK: #include <hipblas.h>
|
||||
#include <cublas_v2.h>
|
||||
|
||||
// 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;
|
||||
}
|
||||
@@ -0,0 +1,327 @@
|
||||
// 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>
|
||||
|
||||
// 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;
|
||||
}
|
||||
Verwijs in nieuw issue
Block a user