AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge 'master' into 'amd-master'

Browse Source 
Change-Id: I35af993320b9694efb88ba6a8567c58b62213fb9


[ROCm/clr commit: 5d0b11543b]
This commit is contained in:
Jenkins
2020-01-19 05:27:16 -05:00
parent f538e3b581 7f788181ec
commit e6f3f63636
8 changed files with 253 additions and 158 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/bin/hipify-perl
+46
View File
@@ -737,6 +737,28 @@ sub simpleSubstitutions {
$ft{'library'} += s/\bcurandSetPseudoRandomGeneratorSeed\b/hiprandSetPseudoRandomGeneratorSeed/g; $ft{'library'} += s/\bcurandSetPseudoRandomGeneratorSeed\b/hiprandSetPseudoRandomGeneratorSeed/g;
$ft{'library'} += s/\bcurandSetQuasiRandomGeneratorDimensions\b/hiprandSetQuasiRandomGeneratorDimensions/g; $ft{'library'} += s/\bcurandSetQuasiRandomGeneratorDimensions\b/hiprandSetQuasiRandomGeneratorDimensions/g;
$ft{'library'} += s/\bcurandSetStream\b/hiprandSetStream/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/\bcusparseCreate\b/hipsparseCreate/g;
$ft{'library'} += s/\bcusparseCreateCsrgemm2Info\b/hipsparseCreateCsrgemm2Info/g; $ft{'library'} += s/\bcusparseCreateCsrgemm2Info\b/hipsparseCreateCsrgemm2Info/g;
$ft{'library'} += s/\bcusparseCreateCsrilu02Info\b/hipsparseCreateCsrilu02Info/g; $ft{'library'} += s/\bcusparseCreateCsrilu02Info\b/hipsparseCreateCsrilu02Info/g;
@@ -744,6 +766,7 @@ sub simpleSubstitutions {
$ft{'library'} += s/\bcusparseCreateHybMat\b/hipsparseCreateHybMat/g; $ft{'library'} += s/\bcusparseCreateHybMat\b/hipsparseCreateHybMat/g;
$ft{'library'} += s/\bcusparseCreateIdentityPermutation\b/hipsparseCreateIdentityPermutation/g; $ft{'library'} += s/\bcusparseCreateIdentityPermutation\b/hipsparseCreateIdentityPermutation/g;
$ft{'library'} += s/\bcusparseCreateMatDescr\b/hipsparseCreateMatDescr/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/\bcusparseDaxpyi\b/hipsparseDaxpyi/g;
$ft{'library'} += s/\bcusparseDcsr2csc\b/hipsparseDcsr2csc/g; $ft{'library'} += s/\bcusparseDcsr2csc\b/hipsparseDcsr2csc/g;
$ft{'library'} += s/\bcusparseDcsr2hyb\b/hipsparseDcsr2hyb/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\b/hipsparseXcsrsort/g;
$ft{'library'} += s/\bcusparseXcsrsort_bufferSizeExt\b/hipsparseXcsrsort_bufferSizeExt/g; $ft{'library'} += s/\bcusparseXcsrsort_bufferSizeExt\b/hipsparseXcsrsort_bufferSizeExt/g;
$ft{'library'} += s/\bcusparseXcsrsv2_zeroPivot\b/hipsparseXcsrsv2_zeroPivot/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\b/hiprand/g;
$ft{'device_library'} += s/\bcurand_discrete\b/hiprand_discrete/g; $ft{'device_library'} += s/\bcurand_discrete\b/hiprand_discrete/g;
$ft{'device_library'} += s/\bcurand_discrete4\b/hiprand_discrete4/g; $ft{'device_library'} += s/\bcurand_discrete4\b/hiprand_discrete4/g;
projects/clr/hipamd/docs/markdown/CUSPARSE_API_supported_by_HIP.md
+46 -46
View File
@@ -176,28 +176,28 @@
|-----------------------------------------------------------|-------------------------------------------------|:----------------:| |-----------------------------------------------------------|-------------------------------------------------|:----------------:|
|`cusparseSaxpyi` |`hipsparseSaxpyi` | |`cusparseSaxpyi` |`hipsparseSaxpyi` |
|`cusparseDaxpyi` |`hipsparseDaxpyi` | |`cusparseDaxpyi` |`hipsparseDaxpyi` |
|`cusparseCaxpyi` | | |`cusparseCaxpyi` |`hipsparseCaxpyi` |
|`cusparseZaxpyi` | | |`cusparseZaxpyi` |`hipsparseZaxpyi` |
|`cusparseSdoti` |`hipsparseSdoti` | |`cusparseSdoti` |`hipsparseSdoti` |
|`cusparseDdoti` |`hipsparseDdoti` | |`cusparseDdoti` |`hipsparseDdoti` |
|`cusparseCdoti` | | |`cusparseCdoti` |`hipsparseCdoti` |
|`cusparseZdoti` | | |`cusparseZdoti` |`hipsparseZdoti` |
|`cusparseCdotci` | | |`cusparseCdotci` |`hipsparseCdotci` |
|`cusparseZdotci` | | |`cusparseZdotci` |`hipsparseZdotci` |
|`cusparseSgthr` |`hipsparseSgthr` | |`cusparseSgthr` |`hipsparseSgthr` |
|`cusparseDgthr` |`hipsparseDgthr` | |`cusparseDgthr` |`hipsparseDgthr` |
|`cusparseCgthr` | | |`cusparseCgthr` |`hipsparseCgthr` |
|`cusparseZgthr` | | |`cusparseZgthr` |`hipsparseZgthr` |
|`cusparseSgthrz` |`hipsparseSgthrz` | |`cusparseSgthrz` |`hipsparseSgthrz` |
|`cusparseDgthrz` |`hipsparseDgthrz` | |`cusparseDgthrz` |`hipsparseDgthrz` |
|`cusparseCgthrz` | | |`cusparseCgthrz` |`hipsparseCgthrz` |
|`cusparseZgthrz` | | |`cusparseZgthrz` |`hipsparseZgthrz` |
|`cusparseSroti` |`hipsparseSroti` | |`cusparseSroti` |`hipsparseSroti` |
|`cusparseDroti` |`hipsparseDroti` | |`cusparseDroti` |`hipsparseDroti` |
|`cusparseSsctr` |`hipsparseSsctr` | |`cusparseSsctr` |`hipsparseSsctr` |
|`cusparseDsctr` |`hipsparseDsctr` | |`cusparseDsctr` |`hipsparseDsctr` |
|`cusparseCsctr` | | |`cusparseCsctr` |`hipsparseCsctr` |
|`cusparseZsctr` | | |`cusparseZsctr` |`hipsparseZsctr` |
## **4. cuSPARSE Level 2 Function Reference** ## **4. cuSPARSE Level 2 Function Reference**
@@ -213,8 +213,8 @@
|`cusparseZbsrxmv` | | |`cusparseZbsrxmv` | |
|`cusparseScsrmv` |`hipsparseScsrmv` | |`cusparseScsrmv` |`hipsparseScsrmv` |
|`cusparseDcsrmv` |`hipsparseDcsrmv` | |`cusparseDcsrmv` |`hipsparseDcsrmv` |
|`cusparseCcsrmv` | | |`cusparseCcsrmv` |`hipsparseCcsrmv` |
|`cusparseZcsrmv` | | |`cusparseZcsrmv` |`hipsparseZcsrmv` |
|`cusparseCsrmvEx` | | 8.0 | |`cusparseCsrmvEx` | | 8.0 |
|`cusparseCsrmvEx_bufferSize` | | 8.0 | |`cusparseCsrmvEx_bufferSize` | | 8.0 |
|`cusparseScsrmv_mp` | | 8.0 | |`cusparseScsrmv_mp` | | 8.0 |
@@ -260,23 +260,23 @@
|`cusparseScsrsv2_bufferSizeExt` |`hipsparseScsrsv2_bufferSizeExt` | |`cusparseScsrsv2_bufferSizeExt` |`hipsparseScsrsv2_bufferSizeExt` |
|`cusparseDcsrsv2_bufferSize` |`hipsparseDcsrsv2_bufferSize` | |`cusparseDcsrsv2_bufferSize` |`hipsparseDcsrsv2_bufferSize` |
|`cusparseDcsrsv2_bufferSizeExt` |`hipsparseDcsrsv2_bufferSizeExt` | |`cusparseDcsrsv2_bufferSizeExt` |`hipsparseDcsrsv2_bufferSizeExt` |
|`cusparseCcsrsv2_bufferSize` | | |`cusparseCcsrsv2_bufferSize` |`hipsparseCcsrsv2_bufferSize` |
|`cusparseCcsrsv2_bufferSizeExt` | | |`cusparseCcsrsv2_bufferSizeExt` |`hipsparseCcsrsv2_bufferSizeExt` |
|`cusparseZcsrsv2_bufferSize` | | |`cusparseZcsrsv2_bufferSize` |`hipsparseZcsrsv2_bufferSize` |
|`cusparseZcsrsv2_bufferSizeExt` | | |`cusparseZcsrsv2_bufferSizeExt` |`hipsparseZcsrsv2_bufferSizeExt` |
|`cusparseScsrsv2_analysis` |`hipsparseScsrsv2_analysis` | |`cusparseScsrsv2_analysis` |`hipsparseScsrsv2_analysis` |
|`cusparseDcsrsv2_analysis` |`hipsparseDcsrsv2_analysis` | |`cusparseDcsrsv2_analysis` |`hipsparseDcsrsv2_analysis` |
|`cusparseCcsrsv2_analysis` | | |`cusparseCcsrsv2_analysis` |`hipsparseCcsrsv2_analysis` |
|`cusparseZcsrsv2_analysis` | | |`cusparseZcsrsv2_analysis` |`hipsparseZcsrsv2_analysis` |
|`cusparseScsrsv2_solve` |`hipsparseScsrsv2_solve` | |`cusparseScsrsv2_solve` |`hipsparseScsrsv2_solve` |
|`cusparseDcsrsv2_solve` |`hipsparseDcsrsv2_solve` | |`cusparseDcsrsv2_solve` |`hipsparseDcsrsv2_solve` |
|`cusparseCcsrsv2_solve` | | |`cusparseCcsrsv2_solve` |`hipsparseCcsrsv2_solve` |
|`cusparseZcsrsv2_solve` | | |`cusparseZcsrsv2_solve` |`hipsparseZcsrsv2_solve` |
|`cusparseXcsrsv2_zeroPivot` |`hipsparseXcsrsv2_zeroPivot` | |`cusparseXcsrsv2_zeroPivot` |`hipsparseXcsrsv2_zeroPivot` |
|`cusparseShybmv` |`hipsparseShybmv` | |`cusparseShybmv` |`hipsparseShybmv` |
|`cusparseDhybmv` |`hipsparseDhybmv` | |`cusparseDhybmv` |`hipsparseDhybmv` |
|`cusparseChybmv` | | |`cusparseChybmv` |`hipsparseChybmv` |
|`cusparseZhybmv` | | |`cusparseZhybmv` |`hipsparseZhybmv` |
|`cusparseShybsv_analysis` | | |`cusparseShybsv_analysis` | |
|`cusparseDhybsv_analysis` | | |`cusparseDhybsv_analysis` | |
|`cusparseChybsv_analysis` | | |`cusparseChybsv_analysis` | |
@@ -292,12 +292,12 @@
|-----------------------------------------------------------|-------------------------------------------------|:----------------:| |-----------------------------------------------------------|-------------------------------------------------|:----------------:|
|`cusparseScsrmm` |`hipsparseScsrmm` | |`cusparseScsrmm` |`hipsparseScsrmm` |
|`cusparseDcsrmm` |`hipsparseDcsrmm` | |`cusparseDcsrmm` |`hipsparseDcsrmm` |
|`cusparseCcsrmm` | | |`cusparseCcsrmm` |`hipsparseCcsrmm` |
|`cusparseZcsrmm` | | |`cusparseZcsrmm` |`hipsparseZcsrmm` |
|`cusparseScsrmm2` |`hipsparseScsrmm2` | |`cusparseScsrmm2` |`hipsparseScsrmm2` |
|`cusparseDcsrmm2` |`hipsparseDcsrmm2` | |`cusparseDcsrmm2` |`hipsparseDcsrmm2` |
|`cusparseCcsrmm2` | | |`cusparseCcsrmm2` |`hipsparseCcsrmm2` |
|`cusparseZcsrmm2` | | |`cusparseZcsrmm2` |`hipsparseZcsrmm2` |
|`cusparseScsrsm_analysis` | | |`cusparseScsrsm_analysis` | |
|`cusparseDcsrsm_analysis` | | |`cusparseDcsrsm_analysis` | |
|`cusparseCcsrsm_analysis` | | |`cusparseCcsrsm_analysis` | |
@@ -366,17 +366,17 @@
|`cusparseXcsrgemmNnz` |`hipsparseXcsrgemmNnz` | |`cusparseXcsrgemmNnz` |`hipsparseXcsrgemmNnz` |
|`cusparseScsrgemm` |`hipsparseScsrgemm` | |`cusparseScsrgemm` |`hipsparseScsrgemm` |
|`cusparseDcsrgemm` |`hipsparseDcsrgemm` | |`cusparseDcsrgemm` |`hipsparseDcsrgemm` |
|`cusparseCcsrgemm` | | |`cusparseCcsrgemm` |`hipsparseCcsrgemm` |
|`cusparseZcsrgemm` | | |`cusparseZcsrgemm` |`hipsparseZcsrgemm` |
|`cusparseXcsrgemm2Nnz` |`hipsparseXcsrgemm2Nnz` | |`cusparseXcsrgemm2Nnz` |`hipsparseXcsrgemm2Nnz` |
|`cusparseScsrgemm2` |`hipsparseScsrgemm2` | |`cusparseScsrgemm2` |`hipsparseScsrgemm2` |
|`cusparseDcsrgemm2` |`hipsparseDcsrgemm2` | |`cusparseDcsrgemm2` |`hipsparseDcsrgemm2` |
|`cusparseCcsrgemm2` | | |`cusparseCcsrgemm2` |`hipsparseCcsrgemm2` |
|`cusparseZcsrgemm2` | | |`cusparseZcsrgemm2` |`hipsparseZcsrgemm2` |
|`cusparseScsrgemm2_bufferSizeExt` |`hipsparseScsrgemm2_bufferSizeExt` | |`cusparseScsrgemm2_bufferSizeExt` |`hipsparseScsrgemm2_bufferSizeExt` |
|`cusparseDcsrgemm2_bufferSizeExt` |`hipsparseDcsrgemm2_bufferSizeExt` | |`cusparseDcsrgemm2_bufferSizeExt` |`hipsparseDcsrgemm2_bufferSizeExt` |
|`cusparseCcsrgemm2_bufferSizeExt` | | |`cusparseCcsrgemm2_bufferSizeExt` |`hipsparseCcsrgemm2_bufferSizeExt` |
|`cusparseZcsrgemm2_bufferSizeExt` | | |`cusparseZcsrgemm2_bufferSizeExt` |`hipsparseZcsrgemm2_bufferSizeExt` |
## **7. cuSPARSE Preconditioners Reference** ## **7. cuSPARSE Preconditioners Reference**
@@ -442,18 +442,18 @@
|`cusparseScsrilu02_bufferSizeExt` |`hipsparseScsrilu02_bufferSizeExt` | |`cusparseScsrilu02_bufferSizeExt` |`hipsparseScsrilu02_bufferSizeExt` |
|`cusparseDcsrilu02_bufferSize` |`hipsparseDcsrilu02_bufferSize` | |`cusparseDcsrilu02_bufferSize` |`hipsparseDcsrilu02_bufferSize` |
|`cusparseDcsrilu02_bufferSizeExt` |`hipsparseDcsrilu02_bufferSizeExt` | |`cusparseDcsrilu02_bufferSizeExt` |`hipsparseDcsrilu02_bufferSizeExt` |
|`cusparseCcsrilu02_bufferSize` | | |`cusparseCcsrilu02_bufferSize` |`hipsparseCcsrilu02_bufferSize` |
|`cusparseCcsrilu02_bufferSizeExt` | | |`cusparseCcsrilu02_bufferSizeExt` |`hipsparseCcsrilu02_bufferSizeExt` |
|`cusparseZcsrilu02_bufferSize` | | |`cusparseZcsrilu02_bufferSize` |`hipsparseZcsrilu02_bufferSize` |
|`cusparseZcsrilu02_bufferSizeExt` | | |`cusparseZcsrilu02_bufferSizeExt` |`hipsparseZcsrilu02_bufferSizeExt` |
|`cusparseScsrilu02_analysis` |`hipsparseScsrilu02_analysis` | |`cusparseScsrilu02_analysis` |`hipsparseScsrilu02_analysis` |
|`cusparseDcsrilu02_analysis` |`hipsparseDcsrilu02_analysis` | |`cusparseDcsrilu02_analysis` |`hipsparseDcsrilu02_analysis` |
|`cusparseCcsrilu02_analysis` | | |`cusparseCcsrilu02_analysis` |`hipsparseCcsrilu02_analysis` |
|`cusparseZcsrilu02_analysis` | | |`cusparseZcsrilu02_analysis` |`hipsparseZcsrilu02_analysis` |
|`cusparseScsrilu02` |`hipsparseScsrilu02` | |`cusparseScsrilu02` |`hipsparseScsrilu02` |
|`cusparseDcsrilu02` |`hipsparseDcsrilu02` | |`cusparseDcsrilu02` |`hipsparseDcsrilu02` |
|`cusparseCcsrilu02` | | |`cusparseCcsrilu02` |`hipsparseCcsrilu02` |
|`cusparseZcsrilu02` | | |`cusparseZcsrilu02` |`hipsparseZcsrilu02` |
|`cusparseXbsric02_zeroPivot` |`hipsparseXcsrilu02_zeroPivot` | |`cusparseXbsric02_zeroPivot` |`hipsparseXcsrilu02_zeroPivot` |
|`cusparseSbsrilu02_numericBoost` | | |`cusparseSbsrilu02_numericBoost` | |
|`cusparseDbsrilu02_numericBoost` | | |`cusparseDbsrilu02_numericBoost` | |
@@ -621,8 +621,8 @@
|`cusparseXcsr2coo` |`hipsparseXcsr2coo` | |`cusparseXcsr2coo` |`hipsparseXcsr2coo` |
|`cusparseScsr2csc` |`hipsparseScsr2csc` | |`cusparseScsr2csc` |`hipsparseScsr2csc` |
|`cusparseDcsr2csc` |`hipsparseDcsr2csc` | |`cusparseDcsr2csc` |`hipsparseDcsr2csc` |
|`cusparseCcsr2csc` | | |`cusparseCcsr2csc` |`hipsparseCcsr2csc` |
|`cusparseZcsr2csc` | | |`cusparseZcsr2csc` |`hipsparseZcsr2csc` |
|`cusparseCsr2cscEx` | | 8.0 | |`cusparseCsr2cscEx` | | 8.0 |
|`cusparseCsr2cscEx2` | | 10.1 | |`cusparseCsr2cscEx2` | | 10.1 |
|`cusparseCsr2cscEx2_bufferSize` | | 10.1 | |`cusparseCsr2cscEx2_bufferSize` | | 10.1 |
@@ -636,8 +636,8 @@
|`cusparseZcsr2csr_compress` | | 8.0 | |`cusparseZcsr2csr_compress` | | 8.0 |
|`cusparseScsr2hyb` |`hipsparseScsr2hyb` | |`cusparseScsr2hyb` |`hipsparseScsr2hyb` |
|`cusparseDcsr2hyb` |`hipsparseDcsr2hyb` | |`cusparseDcsr2hyb` |`hipsparseDcsr2hyb` |
|`cusparseCcsr2hyb` | | |`cusparseCcsr2hyb` |`hipsparseCcsr2hyb` |
|`cusparseZcsr2hyb` | | |`cusparseZcsr2hyb` |`hipsparseZcsr2hyb` |
|`cusparseSdense2csc` | | |`cusparseSdense2csc` | |
|`cusparseDdense2csc` | | |`cusparseDdense2csc` | |
|`cusparseCdense2csc` | | |`cusparseCdense2csc` | |
projects/clr/hipamd/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp
+46 -46
View File
@@ -71,34 +71,34 @@ const std::map<llvm::StringRef, hipCounter> CUDA_SPARSE_FUNCTION_MAP{
// 6. cuSPARSE Level 1 Function Reference // 6. cuSPARSE Level 1 Function Reference
{"cusparseSaxpyi", {"hipsparseSaxpyi", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseSaxpyi", {"hipsparseSaxpyi", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseDaxpyi", {"hipsparseDaxpyi", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDaxpyi", {"hipsparseDaxpyi", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseCaxpyi", {"hipsparseCaxpyi", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseCaxpyi", {"hipsparseCaxpyi", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseZaxpyi", {"hipsparseZaxpyi", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZaxpyi", {"hipsparseZaxpyi", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseSdoti", {"hipsparseSdoti", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseSdoti", {"hipsparseSdoti", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseDdoti", {"hipsparseDdoti", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDdoti", {"hipsparseDdoti", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseCdoti", {"hipsparseCdoti", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseCdoti", {"hipsparseCdoti", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseZdoti", {"hipsparseZdoti", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZdoti", {"hipsparseZdoti", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseCdotci", {"hipsparseCdotci", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseCdotci", {"hipsparseCdotci", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseZdotci", {"hipsparseZdotci", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZdotci", {"hipsparseZdotci", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseSgthr", {"hipsparseSgthr", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseSgthr", {"hipsparseSgthr", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseDgthr", {"hipsparseDgthr", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDgthr", {"hipsparseDgthr", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseCgthr", {"hipsparseCgthr", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseCgthr", {"hipsparseCgthr", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseZgthr", {"hipsparseZgthr", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZgthr", {"hipsparseZgthr", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseSgthrz", {"hipsparseSgthrz", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseSgthrz", {"hipsparseSgthrz", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseDgthrz", {"hipsparseDgthrz", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDgthrz", {"hipsparseDgthrz", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseCgthrz", {"hipsparseCgthrz", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseCgthrz", {"hipsparseCgthrz", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseZgthrz", {"hipsparseZgthrz", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZgthrz", {"hipsparseZgthrz", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseSroti", {"hipsparseSroti", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseSroti", {"hipsparseSroti", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseDroti", {"hipsparseDroti", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDroti", {"hipsparseDroti", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseSsctr", {"hipsparseSsctr", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseSsctr", {"hipsparseSsctr", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseDsctr", {"hipsparseDsctr", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDsctr", {"hipsparseDsctr", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseCsctr", {"hipsparseCsctr", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseCsctr", {"hipsparseCsctr", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseZsctr", {"hipsparseZsctr", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZsctr", {"hipsparseZsctr", "", CONV_LIB_FUNC, API_SPARSE}},
// 7. cuSPARSE Level 2 Function Reference // 7. cuSPARSE Level 2 Function Reference
{"cusparseSbsrmv", {"hipsparseSbsrmv", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseSbsrmv", {"hipsparseSbsrmv", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}},
@@ -113,8 +113,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_SPARSE_FUNCTION_MAP{
{"cusparseScsrmv", {"hipsparseScsrmv", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseScsrmv", {"hipsparseScsrmv", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseDcsrmv", {"hipsparseDcsrmv", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDcsrmv", {"hipsparseDcsrmv", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseCcsrmv", {"hipsparseCcsrmv", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseCcsrmv", {"hipsparseCcsrmv", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseZcsrmv", {"hipsparseZcsrmv", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZcsrmv", {"hipsparseZcsrmv", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseCsrmvEx", {"hipsparseCsrmvEx", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseCsrmvEx", {"hipsparseCsrmvEx", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}},
{"cusparseCsrmvEx_bufferSize", {"hipsparseCsrmvEx_bufferSize", "", 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<llvm::StringRef, hipCounter> CUDA_SPARSE_FUNCTION_MAP{
{"cusparseScsrsv2_bufferSizeExt", {"hipsparseScsrsv2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseScsrsv2_bufferSizeExt", {"hipsparseScsrsv2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseDcsrsv2_bufferSize", {"hipsparseDcsrsv2_bufferSize", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDcsrsv2_bufferSize", {"hipsparseDcsrsv2_bufferSize", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseDcsrsv2_bufferSizeExt", {"hipsparseDcsrsv2_bufferSizeExt", "", 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_bufferSize", {"hipsparseCcsrsv2_bufferSize", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseCcsrsv2_bufferSizeExt", {"hipsparseCcsrsv2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseCcsrsv2_bufferSizeExt", {"hipsparseCcsrsv2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseZcsrsv2_bufferSize", {"hipsparseZcsrsv2_bufferSize", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZcsrsv2_bufferSize", {"hipsparseZcsrsv2_bufferSize", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseZcsrsv2_bufferSizeExt", {"hipsparseZcsrsv2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZcsrsv2_bufferSizeExt", {"hipsparseZcsrsv2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseScsrsv2_analysis", {"hipsparseScsrsv2_analysis", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseScsrsv2_analysis", {"hipsparseScsrsv2_analysis", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseDcsrsv2_analysis", {"hipsparseDcsrsv2_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}}, {"cusparseCcsrsv2_analysis", {"hipsparseCcsrsv2_analysis", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseZcsrsv2_analysis", {"hipsparseZcsrsv2_analysis", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZcsrsv2_analysis", {"hipsparseZcsrsv2_analysis", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseScsrsv2_solve", {"hipsparseScsrsv2_solve", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseScsrsv2_solve", {"hipsparseScsrsv2_solve", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseDcsrsv2_solve", {"hipsparseDcsrsv2_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}}, {"cusparseCcsrsv2_solve", {"hipsparseCcsrsv2_solve", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseZcsrsv2_solve", {"hipsparseZcsrsv2_solve", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZcsrsv2_solve", {"hipsparseZcsrsv2_solve", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseXcsrsv2_zeroPivot", {"hipsparseXcsrsv2_zeroPivot", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseXcsrsv2_zeroPivot", {"hipsparseXcsrsv2_zeroPivot", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseShybmv", {"hipsparseShybmv", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseShybmv", {"hipsparseShybmv", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseDhybmv", {"hipsparseDhybmv", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDhybmv", {"hipsparseDhybmv", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseChybmv", {"hipsparseChybmv", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseChybmv", {"hipsparseChybmv", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseZhybmv", {"hipsparseZhybmv", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZhybmv", {"hipsparseZhybmv", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseShybsv_analysis", {"hipsparseShybsv_analysis", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseShybsv_analysis", {"hipsparseShybsv_analysis", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}},
{"cusparseDhybsv_analysis", {"hipsparseDhybsv_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<llvm::StringRef, hipCounter> CUDA_SPARSE_FUNCTION_MAP{
// 8. cuSPARSE Level 3 Function Reference // 8. cuSPARSE Level 3 Function Reference
{"cusparseScsrmm", {"hipsparseScsrmm", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseScsrmm", {"hipsparseScsrmm", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseDcsrmm", {"hipsparseDcsrmm", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDcsrmm", {"hipsparseDcsrmm", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseCcsrmm", {"hipsparseCcsrmm", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseCcsrmm", {"hipsparseCcsrmm", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseZcsrmm", {"hipsparseZcsrmm", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZcsrmm", {"hipsparseZcsrmm", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseScsrmm2", {"hipsparseScsrmm2", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseScsrmm2", {"hipsparseScsrmm2", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseDcsrmm2", {"hipsparseDcsrmm2", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDcsrmm2", {"hipsparseDcsrmm2", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseCcsrmm2", {"hipsparseCcsrmm2", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseCcsrmm2", {"hipsparseCcsrmm2", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseZcsrmm2", {"hipsparseZcsrmm2", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZcsrmm2", {"hipsparseZcsrmm2", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseScsrsm_analysis", {"hipsparseScsrsm_analysis", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseScsrsm_analysis", {"hipsparseScsrsm_analysis", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}},
{"cusparseDcsrsm_analysis", {"hipsparseDcsrsm_analysis", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseDcsrsm_analysis", {"hipsparseDcsrsm_analysis", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}},
@@ -292,19 +292,19 @@ const std::map<llvm::StringRef, hipCounter> CUDA_SPARSE_FUNCTION_MAP{
{"cusparseXcsrgemmNnz", {"hipsparseXcsrgemmNnz", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseXcsrgemmNnz", {"hipsparseXcsrgemmNnz", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseScsrgemm", {"hipsparseScsrgemm", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseScsrgemm", {"hipsparseScsrgemm", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseDcsrgemm", {"hipsparseDcsrgemm", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDcsrgemm", {"hipsparseDcsrgemm", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseCcsrgemm", {"hipsparseCcsrgemm", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseCcsrgemm", {"hipsparseCcsrgemm", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseZcsrgemm", {"hipsparseZcsrgemm", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZcsrgemm", {"hipsparseZcsrgemm", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseXcsrgemm2Nnz", {"hipsparseXcsrgemm2Nnz", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseXcsrgemm2Nnz", {"hipsparseXcsrgemm2Nnz", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseScsrgemm2", {"hipsparseScsrgemm2", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseScsrgemm2", {"hipsparseScsrgemm2", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseDcsrgemm2", {"hipsparseDcsrgemm2", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDcsrgemm2", {"hipsparseDcsrgemm2", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseCcsrgemm2", {"hipsparseCcsrgemm2", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseCcsrgemm2", {"hipsparseCcsrgemm2", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseZcsrgemm2", {"hipsparseZcsrgemm2", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZcsrgemm2", {"hipsparseZcsrgemm2", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseScsrgemm2_bufferSizeExt", {"hipsparseScsrgemm2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseScsrgemm2_bufferSizeExt", {"hipsparseScsrgemm2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseDcsrgemm2_bufferSizeExt", {"hipsparseDcsrgemm2_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}}, {"cusparseCcsrgemm2_bufferSizeExt", {"hipsparseCcsrgemm2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseZcsrgemm2_bufferSizeExt", {"hipsparseZcsrgemm2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZcsrgemm2_bufferSizeExt", {"hipsparseZcsrgemm2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}},
// 10. cuSPARSE Preconditioners Reference // 10. cuSPARSE Preconditioners Reference
// 10.1. Incomplete Cholesky Factorization : level 0 // 10.1. Incomplete Cholesky Factorization : level 0
@@ -372,20 +372,20 @@ const std::map<llvm::StringRef, hipCounter> CUDA_SPARSE_FUNCTION_MAP{
{"cusparseScsrilu02_bufferSizeExt", {"hipsparseScsrilu02_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseScsrilu02_bufferSizeExt", {"hipsparseScsrilu02_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseDcsrilu02_bufferSize", {"hipsparseDcsrilu02_bufferSize", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDcsrilu02_bufferSize", {"hipsparseDcsrilu02_bufferSize", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseDcsrilu02_bufferSizeExt", {"hipsparseDcsrilu02_bufferSizeExt", "", 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_bufferSize", {"hipsparseCcsrilu02_bufferSize", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseCcsrilu02_bufferSizeExt", {"hipsparseCcsrilu02_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseCcsrilu02_bufferSizeExt", {"hipsparseCcsrilu02_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseZcsrilu02_bufferSize", {"hipsparseZcsrilu02_bufferSize", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZcsrilu02_bufferSize", {"hipsparseZcsrilu02_bufferSize", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseZcsrilu02_bufferSizeExt", {"hipsparseZcsrilu02_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZcsrilu02_bufferSizeExt", {"hipsparseZcsrilu02_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseScsrilu02_analysis", {"hipsparseScsrilu02_analysis", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseScsrilu02_analysis", {"hipsparseScsrilu02_analysis", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseDcsrilu02_analysis", {"hipsparseDcsrilu02_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}}, {"cusparseCcsrilu02_analysis", {"hipsparseCcsrilu02_analysis", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseZcsrilu02_analysis", {"hipsparseZcsrilu02_analysis", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZcsrilu02_analysis", {"hipsparseZcsrilu02_analysis", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseScsrilu02", {"hipsparseScsrilu02", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseScsrilu02", {"hipsparseScsrilu02", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseDcsrilu02", {"hipsparseDcsrilu02", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDcsrilu02", {"hipsparseDcsrilu02", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseCcsrilu02", {"hipsparseCcsrilu02", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseCcsrilu02", {"hipsparseCcsrilu02", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseZcsrilu02", {"hipsparseZcsrilu02", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZcsrilu02", {"hipsparseZcsrilu02", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseXbsric02_zeroPivot", {"hipsparseXcsrilu02_zeroPivot", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseXbsric02_zeroPivot", {"hipsparseXcsrilu02_zeroPivot", "", CONV_LIB_FUNC, API_SPARSE}},
@@ -567,8 +567,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_SPARSE_FUNCTION_MAP{
{"cusparseScsr2csc", {"hipsparseScsr2csc", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseScsr2csc", {"hipsparseScsr2csc", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseDcsr2csc", {"hipsparseDcsr2csc", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDcsr2csc", {"hipsparseDcsr2csc", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseCcsr2csc", {"hipsparseCcsr2csc", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseCcsr2csc", {"hipsparseCcsr2csc", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseZcsr2csc", {"hipsparseZcsr2csc", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZcsr2csc", {"hipsparseZcsr2csc", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseCsr2cscEx", {"hipsparseCsr2cscEx", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseCsr2cscEx", {"hipsparseCsr2cscEx", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}},
{"cusparseCsr2cscEx2", {"hipsparseCsr2cscEx2", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseCsr2cscEx2", {"hipsparseCsr2cscEx2", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}},
@@ -586,8 +586,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_SPARSE_FUNCTION_MAP{
{"cusparseScsr2hyb", {"hipsparseScsr2hyb", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseScsr2hyb", {"hipsparseScsr2hyb", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseDcsr2hyb", {"hipsparseDcsr2hyb", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseDcsr2hyb", {"hipsparseDcsr2hyb", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseCcsr2hyb", {"hipsparseCcsr2hyb", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseCcsr2hyb", {"hipsparseCcsr2hyb", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseZcsr2hyb", {"hipsparseZcsr2hyb", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZcsr2hyb", {"hipsparseZcsr2hyb", "", CONV_LIB_FUNC, API_SPARSE}},
{"cusparseSdense2csc", {"hipsparseSdense2csc", "", 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}}, {"cusparseDdense2csc", {"hipsparseDdense2csc", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}},
projects/clr/hipamd/hipify-clang/src/main.cpp
+7 -11
View File
@@ -158,16 +158,6 @@ void appendArgumentsAdjusters(ct::RefactoringTool &Tool, const std::string &sSou
Tool.appendArgumentsAdjuster(ct::getClangSyntaxOnlyAdjuster()); Tool.appendArgumentsAdjuster(ct::getClangSyntaxOnlyAdjuster());
} }
void addSeparatorIfNeeded(int& argc, const char** argv) {
std::vector<const char*> 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 generatePython() {
bool bToRoc = TranslateToRoc; bool bToRoc = TranslateToRoc;
TranslateToRoc = true; TranslateToRoc = true;
@@ -177,7 +167,13 @@ bool generatePython() {
} }
int main(int argc, const char **argv) { int main(int argc, const char **argv) {
addSeparatorIfNeeded(argc, argv); std::vector<const char*> 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(); llcompat::PrintStackTraceOnErrorSignal();
ct::CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::ZeroOrMore); ct::CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::ZeroOrMore);
if (!llcompat::CheckCompatibility()) { if (!llcompat::CheckCompatibility()) {
projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h
+22
View File
@@ -294,7 +294,13 @@ THE SOFTWARE.
union { union {
Native_vec_ data; Native_vec_ data;
#if __HIP_CLANG_ONLY__
struct {
T x;
};
#else
hip_impl::Scalar_accessor<T, Native_vec_, 0> x; hip_impl::Scalar_accessor<T, Native_vec_, 0> x;
#endif
}; };
using value_type = T; using value_type = T;
@@ -317,8 +323,15 @@ THE SOFTWARE.
union { union {
Native_vec_ data; Native_vec_ data;
#if __HIP_CLANG_ONLY__
struct {
T x;
T y;
};
#else
hip_impl::Scalar_accessor<T, Native_vec_, 0> x; hip_impl::Scalar_accessor<T, Native_vec_, 0> x;
hip_impl::Scalar_accessor<T, Native_vec_, 1> y; hip_impl::Scalar_accessor<T, Native_vec_, 1> y;
#endif
}; };
using value_type = T; using value_type = T;
@@ -498,10 +511,19 @@ THE SOFTWARE.
union { union {
Native_vec_ data; Native_vec_ data;
#if __HIP_CLANG_ONLY__
struct {
T x;
T y;
T z;
T w;
};
#else
hip_impl::Scalar_accessor<T, Native_vec_, 0> x; hip_impl::Scalar_accessor<T, Native_vec_, 0> x;
hip_impl::Scalar_accessor<T, Native_vec_, 1> y; hip_impl::Scalar_accessor<T, Native_vec_, 1> y;
hip_impl::Scalar_accessor<T, Native_vec_, 2> z; hip_impl::Scalar_accessor<T, Native_vec_, 2> z;
hip_impl::Scalar_accessor<T, Native_vec_, 3> w; hip_impl::Scalar_accessor<T, Native_vec_, 3> w;
#endif
}; };
using value_type = T; using value_type = T;
projects/clr/hipamd/src/hip_hcc_internal.h
+51
View File
@@ -91,6 +91,57 @@ extern int HIP_DUMP_CODE_OBJECT;
// TODO - remove when this is standard behavior. // TODO - remove when this is standard behavior.
extern int HCC_OPT_FLUSH; extern int HCC_OPT_FLUSH;
#define IMAGE_PITCH_ALIGNMENT 256
template <typename T> inline T alignDown(T value, size_t alignment) {
return (T)(value & ~(alignment - 1));
}
template <typename T> inline T* alignDown(T* value, size_t alignment) {
return (T*)alignDown((intptr_t)value, alignment);
}
template <typename T> inline T alignUp(T value, size_t alignment) {
return alignDown((T)(value + alignment - 1), alignment);
}
template <typename T> 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 to assign a short TID to each new thread, for HIP debugging purposes.
class TidInfo { class TidInfo {
public: public:
projects/clr/hipamd/src/hip_memory.cpp
+22 -53
View File
@@ -37,6 +37,7 @@ __device__ uint32_t __hip_device_page_flag[__HIP_NUM_PAGES];
namespace hip_internal { namespace hip_internal {
namespace { namespace {
inline inline
const char* hsa_to_string(hsa_status_t err) noexcept 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<hsa_amd_memory_pool_t*>(acc.get_hsa_am_region()); hsa_amd_memory_pool_t* allocRegion = static_cast<hsa_amd_memory_pool_t*>(acc.get_hsa_am_region());
hsa_amd_memory_pool_get_info(*allocRegion, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &allocGranularity); hsa_amd_memory_pool_get_info(*allocRegion, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &allocGranularity);
hsa_ext_image_descriptor_t imageDescriptor; size_t rowPitch = getElementSize(channelOrder, channelType) * alignUp(width, IMAGE_PITCH_ALIGNMENT);
imageDescriptor.geometry = geometry; if(HSA_EXT_IMAGE_GEOMETRY_2DA == geometry)
imageDescriptor.width = width; imageInfo.size = rowPitch * (height == 0 ? 1 : height) * (array_size == 0 ? 1 : array_size) ;
imageDescriptor.height = height; else
imageDescriptor.depth = depth; imageInfo.size = rowPitch * (height == 0 ? 1 : height) * (depth == 0 ? 1 : depth) ;
imageDescriptor.array_size = array_size;
imageDescriptor.format.channel_order = channelOrder;
imageDescriptor.format.channel_type = channelType;
hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW; imageInfo.alignment = IMAGE_PITCH_ALIGNMENT;
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;
}
size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment; size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment;
const unsigned am_flags = 0; const unsigned am_flags = 0;
*ptr = hip_internal::allocAndSharePtr("device_array", imageInfo.size, ctx, *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)) 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_order_t channelOrder;
hsa_ext_image_channel_type_t channelType; hsa_ext_image_channel_type_t channelType;
getChannelOrderAndType(desc, hipReadModeElementType, &channelOrder, &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<hsa_agent_t*>(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; 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)) 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. size_t rowPitch = getElementSize(channelOrder, channelType) * alignUp(width, IMAGE_PITCH_ALIGNMENT);
hc::accelerator acc;
hsa_agent_t* agent = static_cast<hsa_agent_t*>(acc.get_hsa_agent()); if(HSA_EXT_IMAGE_GEOMETRY_2DA == geometry)
if (!agent) imageInfo.size = rowPitch * (height == 0 ? 1 : height) * (array_size == 0 ? 1 : array_size);
return hipErrorInvalidResourceHandle; else
hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW; imageInfo.size = rowPitch * (height == 0 ? 1 : height) * (depth == 0 ? 1 : depth);
hsa_status_t status =
hsa_ext_image_data_get_info_with_layout(*agent, &imageDescriptor, permission, HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, 0, 0, &imageInfo); imageInfo.alignment = IMAGE_PITCH_ALIGNMENT;
if(HSA_STATUS_SUCCESS != status){
return hipErrorRuntimeOther;
}
return hipSuccess; return hipSuccess;
} }
projects/clr/hipamd/src/hip_texture.cpp
+13 -2
View File
@@ -295,6 +295,10 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou
pTexDesc->normalizedCoords); pTexDesc->normalizedCoords);
hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW; 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( if (HSA_STATUS_SUCCESS != hsa_ext_image_create_with_layout(
*agent, &imageDescriptor, devPtr, permission, *agent, &imageDescriptor, devPtr, permission,
HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, pitch, 0, &(pTexture->image)) || 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; 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( if (HSA_STATUS_SUCCESS != hsa_ext_image_create_with_layout(
*agent, &imageDescriptor, devPtr, permission, *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_STATUS_SUCCESS !=
hsa_ext_sampler_create(*agent, &samplerDescriptor, &(pTexture->sampler))) { hsa_ext_sampler_create(*agent, &samplerDescriptor, &(pTexture->sampler))) {
return hipErrorRuntimeOther; return hipErrorRuntimeOther;
@@ -503,6 +509,9 @@ hipError_t ihipBindTexture2DImpl(TlsData *tls, int dim, enum hipTextureReadMode
hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW; 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( if (HSA_STATUS_SUCCESS != hsa_ext_image_create_with_layout(
*agent, &imageDescriptor, devPtr, permission, *agent, &imageDescriptor, devPtr, permission,
HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, pitch, 0, &(pTexture->image)) || 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; 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( if (HSA_STATUS_SUCCESS != hsa_ext_image_create_with_layout(
*agent, &imageDescriptor, array->data, permission, *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_STATUS_SUCCESS !=
hsa_ext_sampler_create(*agent, &samplerDescriptor, &(pTexture->sampler))) { hsa_ext_sampler_create(*agent, &samplerDescriptor, &(pTexture->sampler))) {
return hipErrorRuntimeOther; return hipErrorRuntimeOther;