From 271c1447e21aeb9dea8ff7a90b6e0f83fdf26611 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 13 Jan 2020 18:34:10 +0300 Subject: [PATCH 1/5] [HIPIFY] Sync with hipSPARSE + Update hipify-perl and doc accordingly [ROCm/clr commit: 41beed4839bfe7a86a80b844f6fd08ed86a5aee4] --- projects/clr/hipamd/bin/hipify-perl | 46 ++++++++++ .../markdown/CUSPARSE_API_supported_by_HIP.md | 92 +++++++++---------- .../src/CUDA2HIP_SPARSE_API_functions.cpp | 92 +++++++++---------- 3 files changed, 138 insertions(+), 92 deletions(-) diff --git a/projects/clr/hipamd/bin/hipify-perl b/projects/clr/hipamd/bin/hipify-perl index dc0f79df2c..4fe9b80349 100755 --- a/projects/clr/hipamd/bin/hipify-perl +++ b/projects/clr/hipamd/bin/hipify-perl @@ -737,6 +737,28 @@ sub simpleSubstitutions { $ft{'library'} += s/\bcurandSetPseudoRandomGeneratorSeed\b/hiprandSetPseudoRandomGeneratorSeed/g; $ft{'library'} += s/\bcurandSetQuasiRandomGeneratorDimensions\b/hiprandSetQuasiRandomGeneratorDimensions/g; $ft{'library'} += s/\bcurandSetStream\b/hiprandSetStream/g; + $ft{'library'} += s/\bcusparseCaxpyi\b/hipsparseCaxpyi/g; + $ft{'library'} += s/\bcusparseCcsr2csc\b/hipsparseCcsr2csc/g; + $ft{'library'} += s/\bcusparseCcsr2hyb\b/hipsparseCcsr2hyb/g; + $ft{'library'} += s/\bcusparseCcsrgemm\b/hipsparseCcsrgemm/g; + $ft{'library'} += s/\bcusparseCcsrgemm2\b/hipsparseCcsrgemm2/g; + $ft{'library'} += s/\bcusparseCcsrgemm2_bufferSizeExt\b/hipsparseCcsrgemm2_bufferSizeExt/g; + $ft{'library'} += s/\bcusparseCcsrilu02\b/hipsparseCcsrilu02/g; + $ft{'library'} += s/\bcusparseCcsrilu02_analysis\b/hipsparseCcsrilu02_analysis/g; + $ft{'library'} += s/\bcusparseCcsrilu02_bufferSize\b/hipsparseCcsrilu02_bufferSize/g; + $ft{'library'} += s/\bcusparseCcsrilu02_bufferSizeExt\b/hipsparseCcsrilu02_bufferSizeExt/g; + $ft{'library'} += s/\bcusparseCcsrmm\b/hipsparseCcsrmm/g; + $ft{'library'} += s/\bcusparseCcsrmm2\b/hipsparseCcsrmm2/g; + $ft{'library'} += s/\bcusparseCcsrmv\b/hipsparseCcsrmv/g; + $ft{'library'} += s/\bcusparseCcsrsv2_analysis\b/hipsparseCcsrsv2_analysis/g; + $ft{'library'} += s/\bcusparseCcsrsv2_bufferSize\b/hipsparseCcsrsv2_bufferSize/g; + $ft{'library'} += s/\bcusparseCcsrsv2_bufferSizeExt\b/hipsparseCcsrsv2_bufferSizeExt/g; + $ft{'library'} += s/\bcusparseCcsrsv2_solve\b/hipsparseCcsrsv2_solve/g; + $ft{'library'} += s/\bcusparseCdotci\b/hipsparseCdotci/g; + $ft{'library'} += s/\bcusparseCdoti\b/hipsparseCdoti/g; + $ft{'library'} += s/\bcusparseCgthr\b/hipsparseCgthr/g; + $ft{'library'} += s/\bcusparseCgthrz\b/hipsparseCgthrz/g; + $ft{'library'} += s/\bcusparseChybmv\b/hipsparseChybmv/g; $ft{'library'} += s/\bcusparseCreate\b/hipsparseCreate/g; $ft{'library'} += s/\bcusparseCreateCsrgemm2Info\b/hipsparseCreateCsrgemm2Info/g; $ft{'library'} += s/\bcusparseCreateCsrilu02Info\b/hipsparseCreateCsrilu02Info/g; @@ -744,6 +766,7 @@ sub simpleSubstitutions { $ft{'library'} += s/\bcusparseCreateHybMat\b/hipsparseCreateHybMat/g; $ft{'library'} += s/\bcusparseCreateIdentityPermutation\b/hipsparseCreateIdentityPermutation/g; $ft{'library'} += s/\bcusparseCreateMatDescr\b/hipsparseCreateMatDescr/g; + $ft{'library'} += s/\bcusparseCsctr\b/hipsparseCsctr/g; $ft{'library'} += s/\bcusparseDaxpyi\b/hipsparseDaxpyi/g; $ft{'library'} += s/\bcusparseDcsr2csc\b/hipsparseDcsr2csc/g; $ft{'library'} += s/\bcusparseDcsr2hyb\b/hipsparseDcsr2hyb/g; @@ -823,6 +846,29 @@ sub simpleSubstitutions { $ft{'library'} += s/\bcusparseXcsrsort\b/hipsparseXcsrsort/g; $ft{'library'} += s/\bcusparseXcsrsort_bufferSizeExt\b/hipsparseXcsrsort_bufferSizeExt/g; $ft{'library'} += s/\bcusparseXcsrsv2_zeroPivot\b/hipsparseXcsrsv2_zeroPivot/g; + $ft{'library'} += s/\bcusparseZaxpyi\b/hipsparseZaxpyi/g; + $ft{'library'} += s/\bcusparseZcsr2csc\b/hipsparseZcsr2csc/g; + $ft{'library'} += s/\bcusparseZcsr2hyb\b/hipsparseZcsr2hyb/g; + $ft{'library'} += s/\bcusparseZcsrgemm\b/hipsparseZcsrgemm/g; + $ft{'library'} += s/\bcusparseZcsrgemm2\b/hipsparseZcsrgemm2/g; + $ft{'library'} += s/\bcusparseZcsrgemm2_bufferSizeExt\b/hipsparseZcsrgemm2_bufferSizeExt/g; + $ft{'library'} += s/\bcusparseZcsrilu02\b/hipsparseZcsrilu02/g; + $ft{'library'} += s/\bcusparseZcsrilu02_analysis\b/hipsparseZcsrilu02_analysis/g; + $ft{'library'} += s/\bcusparseZcsrilu02_bufferSize\b/hipsparseZcsrilu02_bufferSize/g; + $ft{'library'} += s/\bcusparseZcsrilu02_bufferSizeExt\b/hipsparseZcsrilu02_bufferSizeExt/g; + $ft{'library'} += s/\bcusparseZcsrmm\b/hipsparseZcsrmm/g; + $ft{'library'} += s/\bcusparseZcsrmm2\b/hipsparseZcsrmm2/g; + $ft{'library'} += s/\bcusparseZcsrmv\b/hipsparseZcsrmv/g; + $ft{'library'} += s/\bcusparseZcsrsv2_analysis\b/hipsparseZcsrsv2_analysis/g; + $ft{'library'} += s/\bcusparseZcsrsv2_bufferSize\b/hipsparseZcsrsv2_bufferSize/g; + $ft{'library'} += s/\bcusparseZcsrsv2_bufferSizeExt\b/hipsparseZcsrsv2_bufferSizeExt/g; + $ft{'library'} += s/\bcusparseZcsrsv2_solve\b/hipsparseZcsrsv2_solve/g; + $ft{'library'} += s/\bcusparseZdotci\b/hipsparseZdotci/g; + $ft{'library'} += s/\bcusparseZdoti\b/hipsparseZdoti/g; + $ft{'library'} += s/\bcusparseZgthr\b/hipsparseZgthr/g; + $ft{'library'} += s/\bcusparseZgthrz\b/hipsparseZgthrz/g; + $ft{'library'} += s/\bcusparseZhybmv\b/hipsparseZhybmv/g; + $ft{'library'} += s/\bcusparseZsctr\b/hipsparseZsctr/g; $ft{'device_library'} += s/\bcurand\b/hiprand/g; $ft{'device_library'} += s/\bcurand_discrete\b/hiprand_discrete/g; $ft{'device_library'} += s/\bcurand_discrete4\b/hiprand_discrete4/g; diff --git a/projects/clr/hipamd/docs/markdown/CUSPARSE_API_supported_by_HIP.md b/projects/clr/hipamd/docs/markdown/CUSPARSE_API_supported_by_HIP.md index 03f7cc2902..fc7a8ee8cd 100644 --- a/projects/clr/hipamd/docs/markdown/CUSPARSE_API_supported_by_HIP.md +++ b/projects/clr/hipamd/docs/markdown/CUSPARSE_API_supported_by_HIP.md @@ -176,28 +176,28 @@ |-----------------------------------------------------------|-------------------------------------------------|:----------------:| |`cusparseSaxpyi` |`hipsparseSaxpyi` | |`cusparseDaxpyi` |`hipsparseDaxpyi` | -|`cusparseCaxpyi` | | -|`cusparseZaxpyi` | | +|`cusparseCaxpyi` |`hipsparseCaxpyi` | +|`cusparseZaxpyi` |`hipsparseZaxpyi` | |`cusparseSdoti` |`hipsparseSdoti` | |`cusparseDdoti` |`hipsparseDdoti` | -|`cusparseCdoti` | | -|`cusparseZdoti` | | -|`cusparseCdotci` | | -|`cusparseZdotci` | | +|`cusparseCdoti` |`hipsparseCdoti` | +|`cusparseZdoti` |`hipsparseZdoti` | +|`cusparseCdotci` |`hipsparseCdotci` | +|`cusparseZdotci` |`hipsparseZdotci` | |`cusparseSgthr` |`hipsparseSgthr` | |`cusparseDgthr` |`hipsparseDgthr` | -|`cusparseCgthr` | | -|`cusparseZgthr` | | +|`cusparseCgthr` |`hipsparseCgthr` | +|`cusparseZgthr` |`hipsparseZgthr` | |`cusparseSgthrz` |`hipsparseSgthrz` | |`cusparseDgthrz` |`hipsparseDgthrz` | -|`cusparseCgthrz` | | -|`cusparseZgthrz` | | +|`cusparseCgthrz` |`hipsparseCgthrz` | +|`cusparseZgthrz` |`hipsparseZgthrz` | |`cusparseSroti` |`hipsparseSroti` | |`cusparseDroti` |`hipsparseDroti` | |`cusparseSsctr` |`hipsparseSsctr` | |`cusparseDsctr` |`hipsparseDsctr` | -|`cusparseCsctr` | | -|`cusparseZsctr` | | +|`cusparseCsctr` |`hipsparseCsctr` | +|`cusparseZsctr` |`hipsparseZsctr` | ## **4. cuSPARSE Level 2 Function Reference** @@ -213,8 +213,8 @@ |`cusparseZbsrxmv` | | |`cusparseScsrmv` |`hipsparseScsrmv` | |`cusparseDcsrmv` |`hipsparseDcsrmv` | -|`cusparseCcsrmv` | | -|`cusparseZcsrmv` | | +|`cusparseCcsrmv` |`hipsparseCcsrmv` | +|`cusparseZcsrmv` |`hipsparseZcsrmv` | |`cusparseCsrmvEx` | | 8.0 | |`cusparseCsrmvEx_bufferSize` | | 8.0 | |`cusparseScsrmv_mp` | | 8.0 | @@ -260,23 +260,23 @@ |`cusparseScsrsv2_bufferSizeExt` |`hipsparseScsrsv2_bufferSizeExt` | |`cusparseDcsrsv2_bufferSize` |`hipsparseDcsrsv2_bufferSize` | |`cusparseDcsrsv2_bufferSizeExt` |`hipsparseDcsrsv2_bufferSizeExt` | -|`cusparseCcsrsv2_bufferSize` | | -|`cusparseCcsrsv2_bufferSizeExt` | | -|`cusparseZcsrsv2_bufferSize` | | -|`cusparseZcsrsv2_bufferSizeExt` | | +|`cusparseCcsrsv2_bufferSize` |`hipsparseCcsrsv2_bufferSize` | +|`cusparseCcsrsv2_bufferSizeExt` |`hipsparseCcsrsv2_bufferSizeExt` | +|`cusparseZcsrsv2_bufferSize` |`hipsparseZcsrsv2_bufferSize` | +|`cusparseZcsrsv2_bufferSizeExt` |`hipsparseZcsrsv2_bufferSizeExt` | |`cusparseScsrsv2_analysis` |`hipsparseScsrsv2_analysis` | |`cusparseDcsrsv2_analysis` |`hipsparseDcsrsv2_analysis` | -|`cusparseCcsrsv2_analysis` | | -|`cusparseZcsrsv2_analysis` | | +|`cusparseCcsrsv2_analysis` |`hipsparseCcsrsv2_analysis` | +|`cusparseZcsrsv2_analysis` |`hipsparseZcsrsv2_analysis` | |`cusparseScsrsv2_solve` |`hipsparseScsrsv2_solve` | |`cusparseDcsrsv2_solve` |`hipsparseDcsrsv2_solve` | -|`cusparseCcsrsv2_solve` | | -|`cusparseZcsrsv2_solve` | | +|`cusparseCcsrsv2_solve` |`hipsparseCcsrsv2_solve` | +|`cusparseZcsrsv2_solve` |`hipsparseZcsrsv2_solve` | |`cusparseXcsrsv2_zeroPivot` |`hipsparseXcsrsv2_zeroPivot` | |`cusparseShybmv` |`hipsparseShybmv` | |`cusparseDhybmv` |`hipsparseDhybmv` | -|`cusparseChybmv` | | -|`cusparseZhybmv` | | +|`cusparseChybmv` |`hipsparseChybmv` | +|`cusparseZhybmv` |`hipsparseZhybmv` | |`cusparseShybsv_analysis` | | |`cusparseDhybsv_analysis` | | |`cusparseChybsv_analysis` | | @@ -292,12 +292,12 @@ |-----------------------------------------------------------|-------------------------------------------------|:----------------:| |`cusparseScsrmm` |`hipsparseScsrmm` | |`cusparseDcsrmm` |`hipsparseDcsrmm` | -|`cusparseCcsrmm` | | -|`cusparseZcsrmm` | | +|`cusparseCcsrmm` |`hipsparseCcsrmm` | +|`cusparseZcsrmm` |`hipsparseZcsrmm` | |`cusparseScsrmm2` |`hipsparseScsrmm2` | |`cusparseDcsrmm2` |`hipsparseDcsrmm2` | -|`cusparseCcsrmm2` | | -|`cusparseZcsrmm2` | | +|`cusparseCcsrmm2` |`hipsparseCcsrmm2` | +|`cusparseZcsrmm2` |`hipsparseZcsrmm2` | |`cusparseScsrsm_analysis` | | |`cusparseDcsrsm_analysis` | | |`cusparseCcsrsm_analysis` | | @@ -366,17 +366,17 @@ |`cusparseXcsrgemmNnz` |`hipsparseXcsrgemmNnz` | |`cusparseScsrgemm` |`hipsparseScsrgemm` | |`cusparseDcsrgemm` |`hipsparseDcsrgemm` | -|`cusparseCcsrgemm` | | -|`cusparseZcsrgemm` | | +|`cusparseCcsrgemm` |`hipsparseCcsrgemm` | +|`cusparseZcsrgemm` |`hipsparseZcsrgemm` | |`cusparseXcsrgemm2Nnz` |`hipsparseXcsrgemm2Nnz` | |`cusparseScsrgemm2` |`hipsparseScsrgemm2` | |`cusparseDcsrgemm2` |`hipsparseDcsrgemm2` | -|`cusparseCcsrgemm2` | | -|`cusparseZcsrgemm2` | | +|`cusparseCcsrgemm2` |`hipsparseCcsrgemm2` | +|`cusparseZcsrgemm2` |`hipsparseZcsrgemm2` | |`cusparseScsrgemm2_bufferSizeExt` |`hipsparseScsrgemm2_bufferSizeExt` | |`cusparseDcsrgemm2_bufferSizeExt` |`hipsparseDcsrgemm2_bufferSizeExt` | -|`cusparseCcsrgemm2_bufferSizeExt` | | -|`cusparseZcsrgemm2_bufferSizeExt` | | +|`cusparseCcsrgemm2_bufferSizeExt` |`hipsparseCcsrgemm2_bufferSizeExt` | +|`cusparseZcsrgemm2_bufferSizeExt` |`hipsparseZcsrgemm2_bufferSizeExt` | ## **7. cuSPARSE Preconditioners Reference** @@ -442,18 +442,18 @@ |`cusparseScsrilu02_bufferSizeExt` |`hipsparseScsrilu02_bufferSizeExt` | |`cusparseDcsrilu02_bufferSize` |`hipsparseDcsrilu02_bufferSize` | |`cusparseDcsrilu02_bufferSizeExt` |`hipsparseDcsrilu02_bufferSizeExt` | -|`cusparseCcsrilu02_bufferSize` | | -|`cusparseCcsrilu02_bufferSizeExt` | | -|`cusparseZcsrilu02_bufferSize` | | -|`cusparseZcsrilu02_bufferSizeExt` | | +|`cusparseCcsrilu02_bufferSize` |`hipsparseCcsrilu02_bufferSize` | +|`cusparseCcsrilu02_bufferSizeExt` |`hipsparseCcsrilu02_bufferSizeExt` | +|`cusparseZcsrilu02_bufferSize` |`hipsparseZcsrilu02_bufferSize` | +|`cusparseZcsrilu02_bufferSizeExt` |`hipsparseZcsrilu02_bufferSizeExt` | |`cusparseScsrilu02_analysis` |`hipsparseScsrilu02_analysis` | |`cusparseDcsrilu02_analysis` |`hipsparseDcsrilu02_analysis` | -|`cusparseCcsrilu02_analysis` | | -|`cusparseZcsrilu02_analysis` | | +|`cusparseCcsrilu02_analysis` |`hipsparseCcsrilu02_analysis` | +|`cusparseZcsrilu02_analysis` |`hipsparseZcsrilu02_analysis` | |`cusparseScsrilu02` |`hipsparseScsrilu02` | |`cusparseDcsrilu02` |`hipsparseDcsrilu02` | -|`cusparseCcsrilu02` | | -|`cusparseZcsrilu02` | | +|`cusparseCcsrilu02` |`hipsparseCcsrilu02` | +|`cusparseZcsrilu02` |`hipsparseZcsrilu02` | |`cusparseXbsric02_zeroPivot` |`hipsparseXcsrilu02_zeroPivot` | |`cusparseSbsrilu02_numericBoost` | | |`cusparseDbsrilu02_numericBoost` | | @@ -621,8 +621,8 @@ |`cusparseXcsr2coo` |`hipsparseXcsr2coo` | |`cusparseScsr2csc` |`hipsparseScsr2csc` | |`cusparseDcsr2csc` |`hipsparseDcsr2csc` | -|`cusparseCcsr2csc` | | -|`cusparseZcsr2csc` | | +|`cusparseCcsr2csc` |`hipsparseCcsr2csc` | +|`cusparseZcsr2csc` |`hipsparseZcsr2csc` | |`cusparseCsr2cscEx` | | 8.0 | |`cusparseCsr2cscEx2` | | 10.1 | |`cusparseCsr2cscEx2_bufferSize` | | 10.1 | @@ -636,8 +636,8 @@ |`cusparseZcsr2csr_compress` | | 8.0 | |`cusparseScsr2hyb` |`hipsparseScsr2hyb` | |`cusparseDcsr2hyb` |`hipsparseDcsr2hyb` | -|`cusparseCcsr2hyb` | | -|`cusparseZcsr2hyb` | | +|`cusparseCcsr2hyb` |`hipsparseCcsr2hyb` | +|`cusparseZcsr2hyb` |`hipsparseZcsr2hyb` | |`cusparseSdense2csc` | | |`cusparseDdense2csc` | | |`cusparseCdense2csc` | | diff --git a/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp b/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp index bc8cfe46a5..1df1bb9cba 100644 --- a/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp +++ b/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp @@ -71,34 +71,34 @@ const std::map CUDA_SPARSE_FUNCTION_MAP{ // 6. cuSPARSE Level 1 Function Reference {"cusparseSaxpyi", {"hipsparseSaxpyi", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDaxpyi", {"hipsparseDaxpyi", "", CONV_LIB_FUNC, API_SPARSE}}, - {"cusparseCaxpyi", {"hipsparseCaxpyi", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseZaxpyi", {"hipsparseZaxpyi", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCaxpyi", {"hipsparseCaxpyi", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseZaxpyi", {"hipsparseZaxpyi", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseSdoti", {"hipsparseSdoti", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDdoti", {"hipsparseDdoti", "", CONV_LIB_FUNC, API_SPARSE}}, - {"cusparseCdoti", {"hipsparseCdoti", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseZdoti", {"hipsparseZdoti", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCdoti", {"hipsparseCdoti", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseZdoti", {"hipsparseZdoti", "", CONV_LIB_FUNC, API_SPARSE}}, - {"cusparseCdotci", {"hipsparseCdotci", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseZdotci", {"hipsparseZdotci", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCdotci", {"hipsparseCdotci", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseZdotci", {"hipsparseZdotci", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseSgthr", {"hipsparseSgthr", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDgthr", {"hipsparseDgthr", "", CONV_LIB_FUNC, API_SPARSE}}, - {"cusparseCgthr", {"hipsparseCgthr", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseZgthr", {"hipsparseZgthr", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCgthr", {"hipsparseCgthr", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseZgthr", {"hipsparseZgthr", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseSgthrz", {"hipsparseSgthrz", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDgthrz", {"hipsparseDgthrz", "", CONV_LIB_FUNC, API_SPARSE}}, - {"cusparseCgthrz", {"hipsparseCgthrz", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseZgthrz", {"hipsparseZgthrz", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCgthrz", {"hipsparseCgthrz", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseZgthrz", {"hipsparseZgthrz", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseSroti", {"hipsparseSroti", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDroti", {"hipsparseDroti", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseSsctr", {"hipsparseSsctr", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDsctr", {"hipsparseDsctr", "", CONV_LIB_FUNC, API_SPARSE}}, - {"cusparseCsctr", {"hipsparseCsctr", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseZsctr", {"hipsparseZsctr", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCsctr", {"hipsparseCsctr", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseZsctr", {"hipsparseZsctr", "", CONV_LIB_FUNC, API_SPARSE}}, // 7. cuSPARSE Level 2 Function Reference {"cusparseSbsrmv", {"hipsparseSbsrmv", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, @@ -113,8 +113,8 @@ const std::map CUDA_SPARSE_FUNCTION_MAP{ {"cusparseScsrmv", {"hipsparseScsrmv", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDcsrmv", {"hipsparseDcsrmv", "", CONV_LIB_FUNC, API_SPARSE}}, - {"cusparseCcsrmv", {"hipsparseCcsrmv", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseZcsrmv", {"hipsparseZcsrmv", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsrmv", {"hipsparseCcsrmv", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseZcsrmv", {"hipsparseZcsrmv", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseCsrmvEx", {"hipsparseCsrmvEx", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseCsrmvEx_bufferSize", {"hipsparseCsrmvEx_bufferSize", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, @@ -171,26 +171,26 @@ const std::map CUDA_SPARSE_FUNCTION_MAP{ {"cusparseScsrsv2_bufferSizeExt", {"hipsparseScsrsv2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDcsrsv2_bufferSize", {"hipsparseDcsrsv2_bufferSize", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDcsrsv2_bufferSizeExt", {"hipsparseDcsrsv2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}}, - {"cusparseCcsrsv2_bufferSize", {"hipsparseCcsrsv2_bufferSize", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseCcsrsv2_bufferSizeExt", {"hipsparseCcsrsv2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseZcsrsv2_bufferSize", {"hipsparseZcsrsv2_bufferSize", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseZcsrsv2_bufferSizeExt", {"hipsparseZcsrsv2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsrsv2_bufferSize", {"hipsparseCcsrsv2_bufferSize", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseCcsrsv2_bufferSizeExt", {"hipsparseCcsrsv2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseZcsrsv2_bufferSize", {"hipsparseZcsrsv2_bufferSize", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseZcsrsv2_bufferSizeExt", {"hipsparseZcsrsv2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseScsrsv2_analysis", {"hipsparseScsrsv2_analysis", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDcsrsv2_analysis", {"hipsparseDcsrsv2_analysis", "", CONV_LIB_FUNC, API_SPARSE}}, - {"cusparseCcsrsv2_analysis", {"hipsparseCcsrsv2_analysis", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseZcsrsv2_analysis", {"hipsparseZcsrsv2_analysis", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsrsv2_analysis", {"hipsparseCcsrsv2_analysis", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseZcsrsv2_analysis", {"hipsparseZcsrsv2_analysis", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseScsrsv2_solve", {"hipsparseScsrsv2_solve", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDcsrsv2_solve", {"hipsparseDcsrsv2_solve", "", CONV_LIB_FUNC, API_SPARSE}}, - {"cusparseCcsrsv2_solve", {"hipsparseCcsrsv2_solve", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseZcsrsv2_solve", {"hipsparseZcsrsv2_solve", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsrsv2_solve", {"hipsparseCcsrsv2_solve", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseZcsrsv2_solve", {"hipsparseZcsrsv2_solve", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseXcsrsv2_zeroPivot", {"hipsparseXcsrsv2_zeroPivot", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseShybmv", {"hipsparseShybmv", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDhybmv", {"hipsparseDhybmv", "", CONV_LIB_FUNC, API_SPARSE}}, - {"cusparseChybmv", {"hipsparseChybmv", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseZhybmv", {"hipsparseZhybmv", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseChybmv", {"hipsparseChybmv", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseZhybmv", {"hipsparseZhybmv", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseShybsv_analysis", {"hipsparseShybsv_analysis", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseDhybsv_analysis", {"hipsparseDhybsv_analysis", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, @@ -205,13 +205,13 @@ const std::map CUDA_SPARSE_FUNCTION_MAP{ // 8. cuSPARSE Level 3 Function Reference {"cusparseScsrmm", {"hipsparseScsrmm", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDcsrmm", {"hipsparseDcsrmm", "", CONV_LIB_FUNC, API_SPARSE}}, - {"cusparseCcsrmm", {"hipsparseCcsrmm", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseZcsrmm", {"hipsparseZcsrmm", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsrmm", {"hipsparseCcsrmm", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseZcsrmm", {"hipsparseZcsrmm", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseScsrmm2", {"hipsparseScsrmm2", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDcsrmm2", {"hipsparseDcsrmm2", "", CONV_LIB_FUNC, API_SPARSE}}, - {"cusparseCcsrmm2", {"hipsparseCcsrmm2", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseZcsrmm2", {"hipsparseZcsrmm2", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsrmm2", {"hipsparseCcsrmm2", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseZcsrmm2", {"hipsparseZcsrmm2", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseScsrsm_analysis", {"hipsparseScsrsm_analysis", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseDcsrsm_analysis", {"hipsparseDcsrsm_analysis", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, @@ -292,19 +292,19 @@ const std::map CUDA_SPARSE_FUNCTION_MAP{ {"cusparseXcsrgemmNnz", {"hipsparseXcsrgemmNnz", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseScsrgemm", {"hipsparseScsrgemm", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDcsrgemm", {"hipsparseDcsrgemm", "", CONV_LIB_FUNC, API_SPARSE}}, - {"cusparseCcsrgemm", {"hipsparseCcsrgemm", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseZcsrgemm", {"hipsparseZcsrgemm", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsrgemm", {"hipsparseCcsrgemm", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseZcsrgemm", {"hipsparseZcsrgemm", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseXcsrgemm2Nnz", {"hipsparseXcsrgemm2Nnz", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseScsrgemm2", {"hipsparseScsrgemm2", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDcsrgemm2", {"hipsparseDcsrgemm2", "", CONV_LIB_FUNC, API_SPARSE}}, - {"cusparseCcsrgemm2", {"hipsparseCcsrgemm2", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseZcsrgemm2", {"hipsparseZcsrgemm2", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsrgemm2", {"hipsparseCcsrgemm2", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseZcsrgemm2", {"hipsparseZcsrgemm2", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseScsrgemm2_bufferSizeExt", {"hipsparseScsrgemm2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDcsrgemm2_bufferSizeExt", {"hipsparseDcsrgemm2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}}, - {"cusparseCcsrgemm2_bufferSizeExt", {"hipsparseCcsrgemm2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseZcsrgemm2_bufferSizeExt", {"hipsparseZcsrgemm2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsrgemm2_bufferSizeExt", {"hipsparseCcsrgemm2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseZcsrgemm2_bufferSizeExt", {"hipsparseZcsrgemm2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}}, // 10. cuSPARSE Preconditioners Reference // 10.1. Incomplete Cholesky Factorization : level 0 @@ -372,20 +372,20 @@ const std::map CUDA_SPARSE_FUNCTION_MAP{ {"cusparseScsrilu02_bufferSizeExt", {"hipsparseScsrilu02_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDcsrilu02_bufferSize", {"hipsparseDcsrilu02_bufferSize", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDcsrilu02_bufferSizeExt", {"hipsparseDcsrilu02_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}}, - {"cusparseCcsrilu02_bufferSize", {"hipsparseCcsrilu02_bufferSize", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseCcsrilu02_bufferSizeExt", {"hipsparseCcsrilu02_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseZcsrilu02_bufferSize", {"hipsparseZcsrilu02_bufferSize", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseZcsrilu02_bufferSizeExt", {"hipsparseZcsrilu02_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsrilu02_bufferSize", {"hipsparseCcsrilu02_bufferSize", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseCcsrilu02_bufferSizeExt", {"hipsparseCcsrilu02_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseZcsrilu02_bufferSize", {"hipsparseZcsrilu02_bufferSize", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseZcsrilu02_bufferSizeExt", {"hipsparseZcsrilu02_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseScsrilu02_analysis", {"hipsparseScsrilu02_analysis", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDcsrilu02_analysis", {"hipsparseDcsrilu02_analysis", "", CONV_LIB_FUNC, API_SPARSE}}, - {"cusparseCcsrilu02_analysis", {"hipsparseCcsrilu02_analysis", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseZcsrilu02_analysis", {"hipsparseZcsrilu02_analysis", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsrilu02_analysis", {"hipsparseCcsrilu02_analysis", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseZcsrilu02_analysis", {"hipsparseZcsrilu02_analysis", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseScsrilu02", {"hipsparseScsrilu02", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDcsrilu02", {"hipsparseDcsrilu02", "", CONV_LIB_FUNC, API_SPARSE}}, - {"cusparseCcsrilu02", {"hipsparseCcsrilu02", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseZcsrilu02", {"hipsparseZcsrilu02", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsrilu02", {"hipsparseCcsrilu02", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseZcsrilu02", {"hipsparseZcsrilu02", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseXbsric02_zeroPivot", {"hipsparseXcsrilu02_zeroPivot", "", CONV_LIB_FUNC, API_SPARSE}}, @@ -567,8 +567,8 @@ const std::map CUDA_SPARSE_FUNCTION_MAP{ {"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}}, + {"cusparseCcsr2csc", {"hipsparseCcsr2csc", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseZcsr2csc", {"hipsparseZcsr2csc", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseCsr2cscEx", {"hipsparseCsr2cscEx", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseCsr2cscEx2", {"hipsparseCsr2cscEx2", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, @@ -586,8 +586,8 @@ const std::map CUDA_SPARSE_FUNCTION_MAP{ {"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}}, + {"cusparseCcsr2hyb", {"hipsparseCcsr2hyb", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseZcsr2hyb", {"hipsparseZcsr2hyb", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseSdense2csc", {"hipsparseSdense2csc", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseDdense2csc", {"hipsparseDdense2csc", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, From d7647096f6556125814df020de6c12c190e1e80f Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 13 Jan 2020 18:49:28 +0300 Subject: [PATCH 2/5] [HIPIFY][fix] Add "--" separator in-place if needed [ROCm/clr commit: df563c966ca849b51ca627d9d789e1cc50c3542b] --- projects/clr/hipamd/hipify-clang/src/main.cpp | 18 +++++++----------- 1 file changed, 7 insertions(+), 11 deletions(-) diff --git a/projects/clr/hipamd/hipify-clang/src/main.cpp b/projects/clr/hipamd/hipify-clang/src/main.cpp index fe798091c4..c75c24b066 100644 --- a/projects/clr/hipamd/hipify-clang/src/main.cpp +++ b/projects/clr/hipamd/hipify-clang/src/main.cpp @@ -158,16 +158,6 @@ void appendArgumentsAdjusters(ct::RefactoringTool &Tool, const std::string &sSou Tool.appendArgumentsAdjuster(ct::getClangSyntaxOnlyAdjuster()); } -void addSeparatorIfNeeded(int& argc, const char** argv) { - std::vector new_argv(argv, argv + argc); - if (std::find(new_argv.begin(), new_argv.end(), std::string("--")) == new_argv.end()) { - new_argv.push_back("--"); - new_argv.push_back(nullptr); - argv = new_argv.data(); - argc++; - } -} - bool generatePython() { bool bToRoc = TranslateToRoc; TranslateToRoc = true; @@ -177,7 +167,13 @@ bool generatePython() { } int main(int argc, const char **argv) { - addSeparatorIfNeeded(argc, argv); + std::vector new_argv(argv, argv + argc); + if (std::find(new_argv.begin(), new_argv.end(), std::string("--")) == new_argv.end()) { + new_argv.push_back("--"); + new_argv.push_back(nullptr); + argv = new_argv.data(); + argc++; + } llcompat::PrintStackTraceOnErrorSignal(); ct::CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::ZeroOrMore); if (!llcompat::CheckCompatibility()) { From 8bf421009533c1375e2bded0f1a28c7c0205fffd Mon Sep 17 00:00:00 2001 From: kpyzhov Date: Tue, 14 Jan 2020 14:30:21 -0500 Subject: [PATCH 3/5] Add missing constructors for Scalar_accessor class. (#1792) [ROCm/clr commit: 09186146153fb5d1b2e28d93e42fa76d40c23b1b] --- .../clr/hipamd/include/hip/hcc_detail/hip_vector_types.h | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h index 2079bb7e44..453e92ebfc 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h @@ -105,6 +105,14 @@ THE SOFTWARE. // Idea from https://t0rakka.silvrback.com/simd-scalar-accessor Vector data; + __host__ __device__ + Scalar_accessor() = default; + + __host__ __device__ + Scalar_accessor(T x) { + data[idx] = x; + } + __host__ __device__ operator T() const noexcept { return data[idx]; } __host__ __device__ From 2d617bba3bbe148d3ba7066d2dc2965e4029abac Mon Sep 17 00:00:00 2001 From: kpyzhov Date: Wed, 15 Jan 2020 21:17:08 -0500 Subject: [PATCH 4/5] Don't use accelerated vector element access for hip-clang. (#1796) [ROCm/clr commit: 1d6a2fb0646f83064435bfbf5ea8536597c81d7d] --- .../include/hip/hcc_detail/hip_vector_types.h | 30 ++++++++++++++----- 1 file changed, 22 insertions(+), 8 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h index 453e92ebfc..9ef8e32516 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h @@ -105,14 +105,6 @@ THE SOFTWARE. // Idea from https://t0rakka.silvrback.com/simd-scalar-accessor Vector data; - __host__ __device__ - Scalar_accessor() = default; - - __host__ __device__ - Scalar_accessor(T x) { - data[idx] = x; - } - __host__ __device__ operator T() const noexcept { return data[idx]; } __host__ __device__ @@ -302,7 +294,13 @@ THE SOFTWARE. union { Native_vec_ data; +#if __HIP_CLANG_ONLY__ + struct { + T x; + }; +#else hip_impl::Scalar_accessor x; +#endif }; using value_type = T; @@ -325,8 +323,15 @@ THE SOFTWARE. union { Native_vec_ data; +#if __HIP_CLANG_ONLY__ + struct { + T x; + T y; + }; +#else hip_impl::Scalar_accessor x; hip_impl::Scalar_accessor y; +#endif }; using value_type = T; @@ -506,10 +511,19 @@ THE SOFTWARE. union { Native_vec_ data; +#if __HIP_CLANG_ONLY__ + struct { + T x; + T y; + T z; + T w; + }; +#else hip_impl::Scalar_accessor x; hip_impl::Scalar_accessor y; hip_impl::Scalar_accessor z; hip_impl::Scalar_accessor w; +#endif }; using value_type = T; From 7f788181ec4d037169772b08ff7f30fe00dd2dec Mon Sep 17 00:00:00 2001 From: ansurya <50609411+ansurya@users.noreply.github.com> Date: Thu, 16 Jan 2020 08:54:30 +0530 Subject: [PATCH 5/5] Fix texture tests, always pass row pitch to HSA API for linear layout images (#1785) => New ROCr calculates pitch as per HSA specification and addrlib is used to check whether HW can support that configuration. Hence few texture tests are failing with HSA_EXT_STATUS_ERROR_IMAGE_PITCH_UNSUPPORTED. => Determine pitch for linear images and always pass rowpitch to HSA API's. [ROCm/clr commit: 688b6e6af07104c9427298c1e06ba1503dfaf7a7] --- projects/clr/hipamd/src/hip_hcc_internal.h | 51 +++++++++++++++ projects/clr/hipamd/src/hip_memory.cpp | 75 +++++++--------------- projects/clr/hipamd/src/hip_texture.cpp | 15 ++++- 3 files changed, 86 insertions(+), 55 deletions(-) diff --git a/projects/clr/hipamd/src/hip_hcc_internal.h b/projects/clr/hipamd/src/hip_hcc_internal.h index 6787bc7e5e..0510015c42 100644 --- a/projects/clr/hipamd/src/hip_hcc_internal.h +++ b/projects/clr/hipamd/src/hip_hcc_internal.h @@ -91,6 +91,57 @@ extern int HIP_DUMP_CODE_OBJECT; // TODO - remove when this is standard behavior. extern int HCC_OPT_FLUSH; +#define IMAGE_PITCH_ALIGNMENT 256 +template inline T alignDown(T value, size_t alignment) { + return (T)(value & ~(alignment - 1)); +} + +template inline T* alignDown(T* value, size_t alignment) { + return (T*)alignDown((intptr_t)value, alignment); +} + +template inline T alignUp(T value, size_t alignment) { + return alignDown((T)(value + alignment - 1), alignment); +} + +template inline T* alignUp(T* value, size_t alignment) { + return (T*)alignDown((intptr_t)(value + alignment - 1), alignment); +} + +size_t getNumChannels(hsa_ext_image_channel_order_t channelOrder) { + switch (channelOrder) { + case HSA_EXT_IMAGE_CHANNEL_ORDER_RG: + return 2; + case HSA_EXT_IMAGE_CHANNEL_ORDER_RGB: + return 3; + case HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA: + return 4; + case HSA_EXT_IMAGE_CHANNEL_ORDER_R: + default: + return 1; + } +} + +size_t getElementSize(hsa_ext_image_channel_order_t channelOrder, hsa_ext_image_channel_type_t channelType) { + size_t bytesPerPixel = getNumChannels(channelOrder); + switch (channelType) { + case HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8: + case HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8: + break; + + case HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32: + case HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32: + case HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT: + bytesPerPixel *= 4; + break; + + default: + bytesPerPixel *= 2; + break; + } + return bytesPerPixel; +} + // Class to assign a short TID to each new thread, for HIP debugging purposes. class TidInfo { public: diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 43399bfdb5..444e41107a 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -37,6 +37,7 @@ __device__ uint32_t __hip_device_page_flag[__HIP_NUM_PAGES]; namespace hip_internal { namespace { + inline const char* hsa_to_string(hsa_status_t err) noexcept { @@ -791,21 +792,13 @@ hipError_t allocImage(TlsData* tls,hsa_ext_image_geometry_t geometry, int width, hsa_amd_memory_pool_t* allocRegion = static_cast(acc.get_hsa_am_region()); hsa_amd_memory_pool_get_info(*allocRegion, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &allocGranularity); - hsa_ext_image_descriptor_t imageDescriptor; - imageDescriptor.geometry = geometry; - imageDescriptor.width = width; - imageDescriptor.height = height; - imageDescriptor.depth = depth; - imageDescriptor.array_size = array_size; - imageDescriptor.format.channel_order = channelOrder; - imageDescriptor.format.channel_type = channelType; + size_t rowPitch = getElementSize(channelOrder, channelType) * alignUp(width, IMAGE_PITCH_ALIGNMENT); + if(HSA_EXT_IMAGE_GEOMETRY_2DA == geometry) + imageInfo.size = rowPitch * (height == 0 ? 1 : height) * (array_size == 0 ? 1 : array_size) ; + else + imageInfo.size = rowPitch * (height == 0 ? 1 : height) * (depth == 0 ? 1 : depth) ; - hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW; - hsa_status_t status = - hsa_ext_image_data_get_info_with_layout(*agent, &imageDescriptor, permission, HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, 0, 0, &imageInfo); - if(imageInfo.size == 0 || HSA_STATUS_SUCCESS != status){ - return hipErrorRuntimeOther; - } + imageInfo.alignment = IMAGE_PITCH_ALIGNMENT; size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment; const unsigned am_flags = 0; *ptr = hip_internal::allocAndSharePtr("device_array", imageInfo.size, ctx, @@ -898,55 +891,31 @@ extern void getChannelOrderAndType(const hipChannelFormatDesc& desc, hipError_t GetImageInfo(hsa_ext_image_geometry_t geometry,int width, int height, int depth, hipChannelFormatDesc desc, hsa_ext_image_data_info_t &imageInfo,int array_size __dparm(0)) { - hsa_ext_image_descriptor_t imageDescriptor; - imageDescriptor.geometry = geometry; - imageDescriptor.width = width; - imageDescriptor.height = height; - imageDescriptor.depth = depth; - imageDescriptor.array_size = array_size; hsa_ext_image_channel_order_t channelOrder; hsa_ext_image_channel_type_t channelType; getChannelOrderAndType(desc, hipReadModeElementType, &channelOrder, &channelType); - imageDescriptor.format.channel_order = channelOrder; - imageDescriptor.format.channel_type = channelType; - - hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW; - // Get the current device agent. - hc::accelerator acc; - hsa_agent_t* agent = static_cast(acc.get_hsa_agent()); - if (!agent) - return hipErrorInvalidHandle; - hsa_status_t status = - hsa_ext_image_data_get_info_with_layout(*agent, &imageDescriptor, permission, HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, 0, 0, &imageInfo); - if(HSA_STATUS_SUCCESS != status){ - return hipErrorRuntimeOther; - } + size_t rowPitch = getElementSize(channelOrder, channelType) * alignUp(width, IMAGE_PITCH_ALIGNMENT); + if(HSA_EXT_IMAGE_GEOMETRY_2DA == geometry) + imageInfo.size = rowPitch * (height == 0 ? 1 : height) * (array_size == 0 ? 1 : array_size); + else + imageInfo.size = rowPitch * (height == 0 ? 1 : height) * (depth == 0 ? 1 : depth); + imageInfo.alignment = IMAGE_PITCH_ALIGNMENT; return hipSuccess; } hipError_t GetImageInfo(hsa_ext_image_geometry_t geometry,size_t width, size_t height, size_t depth, hsa_ext_image_channel_order_t channelOrder, hsa_ext_image_channel_type_t channelType, hsa_ext_image_data_info_t &imageInfo,size_t array_size __dparm(0)) { - hsa_ext_image_descriptor_t imageDescriptor; - imageDescriptor.geometry = geometry; - imageDescriptor.width = width; - imageDescriptor.height = height; - imageDescriptor.depth = depth; - imageDescriptor.array_size = array_size; - imageDescriptor.format.channel_order = channelOrder; - imageDescriptor.format.channel_type = channelType; - // Get the current device agent. - hc::accelerator acc; - hsa_agent_t* agent = static_cast(acc.get_hsa_agent()); - if (!agent) - return hipErrorInvalidResourceHandle; - hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW; - hsa_status_t status = - hsa_ext_image_data_get_info_with_layout(*agent, &imageDescriptor, permission, HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, 0, 0, &imageInfo); - if(HSA_STATUS_SUCCESS != status){ - return hipErrorRuntimeOther; - } + size_t rowPitch = getElementSize(channelOrder, channelType) * alignUp(width, IMAGE_PITCH_ALIGNMENT); + + if(HSA_EXT_IMAGE_GEOMETRY_2DA == geometry) + imageInfo.size = rowPitch * (height == 0 ? 1 : height) * (array_size == 0 ? 1 : array_size); + else + imageInfo.size = rowPitch * (height == 0 ? 1 : height) * (depth == 0 ? 1 : depth); + + imageInfo.alignment = IMAGE_PITCH_ALIGNMENT; + return hipSuccess; } diff --git a/projects/clr/hipamd/src/hip_texture.cpp b/projects/clr/hipamd/src/hip_texture.cpp index 7ea8da8624..81eb5ad272 100644 --- a/projects/clr/hipamd/src/hip_texture.cpp +++ b/projects/clr/hipamd/src/hip_texture.cpp @@ -295,6 +295,10 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou pTexDesc->normalizedCoords); hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW; + + if(hipResourceTypePitch2D != pResDesc->resType) + pitch = getElementSize(channelOrder, channelType) * alignUp(imageDescriptor.width, IMAGE_PITCH_ALIGNMENT); + if (HSA_STATUS_SUCCESS != hsa_ext_image_create_with_layout( *agent, &imageDescriptor, devPtr, permission, HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, pitch, 0, &(pTexture->image)) || @@ -430,9 +434,11 @@ hipError_t ihipBindTextureImpl(TlsData *tls_, int dim, enum hipTextureReadMode r hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW; + size_t rowPitch = getElementSize(channelOrder, channelType) * alignUp(size, IMAGE_PITCH_ALIGNMENT); + if (HSA_STATUS_SUCCESS != hsa_ext_image_create_with_layout( *agent, &imageDescriptor, devPtr, permission, - HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, 0, 0, &(pTexture->image)) || + HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, rowPitch, 0, &(pTexture->image)) || HSA_STATUS_SUCCESS != hsa_ext_sampler_create(*agent, &samplerDescriptor, &(pTexture->sampler))) { return hipErrorRuntimeOther; @@ -503,6 +509,9 @@ hipError_t ihipBindTexture2DImpl(TlsData *tls, int dim, enum hipTextureReadMode hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW; + if( 0 == pitch) + pitch = getElementSize(channelOrder, channelType) * alignUp(width, IMAGE_PITCH_ALIGNMENT); + if (HSA_STATUS_SUCCESS != hsa_ext_image_create_with_layout( *agent, &imageDescriptor, devPtr, permission, HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, pitch, 0, &(pTexture->image)) || @@ -606,9 +615,11 @@ hipError_t ihipBindTextureToArrayImpl(TlsData *tls_, int dim, enum hipTextureRea hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW; + size_t rowPitch = getElementSize(channelOrder, channelType) * alignUp(imageDescriptor.width, IMAGE_PITCH_ALIGNMENT); + if (HSA_STATUS_SUCCESS != hsa_ext_image_create_with_layout( *agent, &imageDescriptor, array->data, permission, - HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, 0, 0, &(pTexture->image)) || + HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, rowPitch, 0, &(pTexture->image)) || HSA_STATUS_SUCCESS != hsa_ext_sampler_create(*agent, &samplerDescriptor, &(pTexture->sampler))) { return hipErrorRuntimeOther;