diff --git a/Jenkinsfile b/Jenkinsfile
index 6d37f10e3c..e6b60f398e 100644
--- a/Jenkinsfile
+++ b/Jenkinsfile
@@ -167,8 +167,6 @@ def docker_build_inside_image( def build_image, String inside_args, String platf
}
// Cap the maximum amount of testing, in case of hangs
- // Excluding hipVectorTypes test from automation; due to regression from HCC commit 2367133
- // Excluding hipFloatMath test from automation; due to regression from ROCDL commit 2fc04e1
timeout(time: 1, unit: 'HOURS')
{
stage("${platform} unit testing")
@@ -178,7 +176,7 @@ def docker_build_inside_image( def build_image, String inside_args, String platf
cd ${build_dir_rel}
make install -j\$(nproc)
make build_tests -i -j\$(nproc)
- ctest -E "(hipVectorTypes.tst|hipVectorTypesDevice.tst|hipFloatMath.tst)"
+ ctest
"""
// If unit tests output a junit or xunit file in the future, jenkins can parse that file
// to display test results on the dashboard
diff --git a/bin/hipify-perl b/bin/hipify-perl
index 3ea201d14f..792bf86d94 100755
--- a/bin/hipify-perl
+++ b/bin/hipify-perl
@@ -287,6 +287,9 @@ while (@ARGV) {
$ft{'mem'} += s/\bcudaMemcpy2DToArray\b/hipMemcpy2DToArray/g;
$ft{'mem'} += s/\bcudaMemcpyToArray\b/hipMemcpyToArray/g;
+ $ft{'mem'} += s/\bcudaGetSymbolAddress\b/hipGetSymbolAddress/g;
+ $ft{'mem'} += s/\bcudaGetSymbolSize\b/hipGetSymbolSize/g;
+
#--------
# Memory management:
$ft{'mem'} += s/\bcudaMalloc\b/hipMalloc/g;
diff --git a/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md b/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md
index 6190f6565a..087b49b977 100644
--- a/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md
+++ b/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md
@@ -132,8 +132,8 @@
| `cudaFreeHost` | `hipHostFree` |
| `cudaFreeMipmappedArray` | |
| `cudaGetMipmappedArrayLevel` | |
-| `cudaGetSymbolAddress` | |
-| `cudaGetSymbolSize` | |
+| `cudaGetSymbolAddress` | `hipGetSymbolAddress` |
+| `cudaGetSymbolSize` | `hipGetSymbolSize` |
| `cudaHostAlloc` | `hipHostMalloc` |
| `cudaHostGetDevicePointer` | `hipHostGetDevicePointer` |
| `cudaHostGetFlags` | `hipHostGetFlags` |
@@ -383,8 +383,8 @@
| `cudaCreateChannelDesc` | `hipCreateChannelDesc` |
| `cudaFuncGetAttributes` | |
| `cudaFuncSetCacheConfig` | |
-| `cudaGetSymbolAddress` | |
-| `cudaGetSymbolSize` | |
+| `cudaGetSymbolAddress` | `hipGetSymbolAddress` |
+| `cudaGetSymbolSize` | `hipGetSymbolSize` |
| `cudaGetTextureAlignmentOffset` | |
| `cudaLaunch` | |
| `cudaLaunchKernel` | |
diff --git a/docs/markdown/CUSPARSE_API_supported_by_HIP.md b/docs/markdown/CUSPARSE_API_supported_by_HIP.md
new file mode 100644
index 0000000000..cbc4074daf
--- /dev/null
+++ b/docs/markdown/CUSPARSE_API_supported_by_HIP.md
@@ -0,0 +1,89 @@
+# CUSPARSE API supported by HIP
+
+## **1. CUSPARSE Data types**
+
+| **type** | **CUDA** | **HIP** |
+|-------------:|---------------------------------------------------------------|------------------------------------------------------------|
+| enum |***`cusparseAction_t`*** |***`hipsparseAction_t`*** |
+| 0 |*`CUSPARSE_ACTION_SYMBOLIC`* |*`HIPSPARSE_ACTION_SYMBOLIC`* |
+| 1 |*`CUSPARSE_ACTION_NUMERIC`* |*`HIPSPARSE_ACTION_NUMERIC`* |
+| enum |***`cusparseDirection_t`*** | |
+| 0 |*`CUSPARSE_DIRECTION_ROW`* | |
+| 1 |*`CUSPARSE_DIRECTION_COLUMN`* | |
+| enum |***`cusparseHybPartition_t`*** |***`hipsparseHybPartition_t`*** |
+| 0 |*`CUSPARSE_HYB_PARTITION_AUTO`* |*`HIPSPARSE_HYB_PARTITION_AUTO`* |
+| 1 |*`CUSPARSE_HYB_PARTITION_USER`* |*`HIPSPARSE_HYB_PARTITION_USER`* |
+| 2 |*`CUSPARSE_HYB_PARTITION_MAX`* |*`HIPSPARSE_HYB_PARTITION_MAX`* |
+| enum |***`cusparseDiagType_t`*** |***`hipsparseDiagType_t`*** |
+| 0 |*`CUSPARSE_DIAG_TYPE_NON_UNIT`* |*`HIPSPARSE_DIAG_TYPE_NON_UNIT`* |
+| 1 |*`CUSPARSE_DIAG_TYPE_UNIT`* |*`HIPSPARSE_DIAG_TYPE_UNIT`* |
+| enum |***`cusparseFillMode_t`*** |***`hipsparseFillMode_t`*** |
+| 0 |*`CUSPARSE_FILL_MODE_LOWER`* |*`HIPSPARSE_FILL_MODE_LOWER`* |
+| 1 |*`CUSPARSE_FILL_MODE_UPPER`* |*`HIPSPARSE_FILL_MODE_UPPER`* |
+| enum |***`cusparseIndexBase_t`*** |***`hipsparseIndexBase_t`*** |
+| 0 |*`CUSPARSE_INDEX_BASE_ZERO`* |*`HIPSPARSE_INDEX_BASE_ZERO`* |
+| 1 |*`CUSPARSE_INDEX_BASE_ONE`* |*`HIPSPARSE_INDEX_BASE_ONE`* |
+| enum |***`cusparseMatrixType_t`*** |***`hipsparseMatrixType_t`*** |
+| 0 |*`CUSPARSE_MATRIX_TYPE_GENERAL`* |*`HIPSPARSE_MATRIX_TYPE_GENERAL`* |
+| 1 |*`CUSPARSE_MATRIX_TYPE_SYMMETRIC`* |*`HIPSPARSE_MATRIX_TYPE_SYMMETRIC`* |
+| 2 |*`CUSPARSE_MATRIX_TYPE_HERMITIAN`* |*`HIPSPARSE_MATRIX_TYPE_HERMITIAN`* |
+| 3 |*`CUSPARSE_MATRIX_TYPE_TRIANGULAR`* |*`HIPSPARSE_MATRIX_TYPE_TRIANGULAR`* |
+| enum |***`cusparseOperation_t`*** |***`hipsparseOperation_t`*** |
+| 0 |*`CUSPARSE_OPERATION_NON_TRANSPOSE`* |*`HIPSPARSE_OPERATION_NON_TRANSPOSE`* |
+| 1 |*`CUSPARSE_OPERATION_TRANSPOSE`* |*`HIPSPARSE_OPERATION_TRANSPOSE`* |
+| 2 |*`CUSPARSE_OPERATION_CONJUGATE_TRANSPOSE`* |*`HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE`* |
+| enum |***`cusparsePointerMode_t`*** |***`hipsparsePointerMode_t`*** |
+| 0 |*`CUSPARSE_POINTER_MODE_HOST`* |*`HIPSPARSE_POINTER_MODE_HOST`* |
+| 1 |*`CUSPARSE_POINTER_MODE_DEVICE`* |*`HIPSPARSE_POINTER_MODE_DEVICE`* |
+| enum |***`cusparseAlgMode_t`*** | |
+| 0 |*`CUSPARSE_ALG0`* | |
+| 1 |*`CUSPARSE_ALG1`* | |
+| 0 |*`CUSPARSE_ALG_NAIVE`* | |
+| 1 |*`CUSPARSE_ALG_MERGE_PATH`* | |
+| enum |***`cusparseSolvePolicy_t`*** |***`hipsparseSolvePolicy_t`*** |
+| 0 |*`CUSPARSE_SOLVE_POLICY_NO_LEVEL`* |*`HIPSPARSE_SOLVE_POLICY_NO_LEVEL`* |
+| 1 |*`CUSPARSE_SOLVE_POLICY_USE_LEVEL`* |*`HIPSPARSE_SOLVE_POLICY_USE_LEVEL`* |
+| enum |***`cusparseStatus_t`*** |***`hipsparseMatrixType_t`*** |
+| 0 |*`CUSPARSE_STATUS_SUCCESS`* |*`HIPSPARSE_STATUS_SUCCESS`* |
+| 1 |*`CUSPARSE_STATUS_NOT_INITIALIZED`* |*`HIPSPARSE_STATUS_NOT_INITIALIZED`* |
+| 2 |*`CUSPARSE_STATUS_ALLOC_FAILED`* |*`HIPSPARSE_STATUS_ALLOC_FAILED`* |
+| 3 |*`CUSPARSE_STATUS_INVALID_VALUE`* |*`HIPSPARSE_STATUS_INVALID_VALUE`* |
+| 4 |*`CUSPARSE_STATUS_ARCH_MISMATCH`* |*`HIPSPARSE_STATUS_ARCH_MISMATCH`* |
+| 5 |*`CUSPARSE_STATUS_MAPPING_ERROR`* |*`HIPSPARSE_STATUS_MAPPING_ERROR`* |
+| 6 |*`CUSPARSE_STATUS_EXECUTION_FAILED`* |*`HIPSPARSE_STATUS_EXECUTION_FAILED`* |
+| 7 |*`CUSPARSE_STATUS_INTERNAL_ERROR`* |*`HIPSPARSE_STATUS_INTERNAL_ERROR`* |
+| 8 |*`CUSPARSE_STATUS_MATRIX_TYPE_NOT_SUPPORTED`* |*`HIPSPARSE_STATUS_MATRIX_TYPE_NOT_SUPPORTED`* |
+| 9 |*`CUSPARSE_STATUS_ZERO_PIVOT`* |*`HIPSPARSE_STATUS_ZERO_PIVOT`* |
+| struct |`cusparseContext` | |
+| typedef |`cusparseHandle_t` |`hipsparseHandle_t` |
+| struct |`cusparseHybMat` | |
+| typedef |`cusparseHybMat_t` |`hipsparseHybMat_t` |
+| struct |`cusparseMatDescr` | |
+| typedef |`cusparseMatDescr_t` |`hipsparseMatDescr_t` |
+| struct |`cusparseSolveAnalysisInfo` | |
+| typedef |`cusparseSolveAnalysisInfo_t` | |
+| struct |`csrsv2Info` | |
+| typedef |`csrsv2Info_t` |`csrsv2Info_t` |
+| struct |`csrsm2Info` | |
+| typedef |`csrsm2Info_t` | |
+| struct |`bsrsv2Info` | |
+| typedef |`bsrsv2Info_t` | |
+| struct |`bsrsm2Info` | |
+| typedef |`bsrsm2Info_t` | |
+| struct |`bsric02Info` | |
+| typedef |`bsric02Info_t` | |
+| struct |`csrilu02Info` | |
+| typedef |`csrilu02Info_t` |`csrilu02Info_t` |
+| struct |`bsrilu02Info` | |
+| typedef |`bsrilu02Info_t` | |
+| struct |`csru2csrInfo` | |
+| typedef |`csru2csrInfo_t` | |
+| struct |`cusparseColorInfo` | |
+| typedef |`cusparseColorInfo_t` | |
+| struct |`pruneInfo` | |
+| typedef |`pruneInfo_t` | |
+
+## **2. CUSPARSE API functions**
+
+| **CUDA** | **HIP** |
+|-----------------------------------------------------------|-------------------------------------------------|
diff --git a/docs/markdown/hip_kernel_language.md b/docs/markdown/hip_kernel_language.md
index 7c209acadf..d69f5a04a8 100644
--- a/docs/markdown/hip_kernel_language.md
+++ b/docs/markdown/hip_kernel_language.md
@@ -159,7 +159,7 @@ void callMyKernel()
## Variable-Type Qualifiers
### `__constant__`
-The `__constant__` keyword is supported. The host writes constant memory before launching the kernel; from the GPU, this memory is read-only during kernel execution. The functions for accessing constant memory (hipGetSymbolAddress(), hipGetSymbolSize(), hipMemcpyToSymbol(), hipMemcpyToSymbolAsync, hipMemcpyFromSymbol, hipMemcpyFromSymbolAsync) are under development.
+The `__constant__` keyword is supported. The host writes constant memory before launching the kernel; from the GPU, this memory is read-only during kernel execution. The functions for accessing constant memory (hipGetSymbolAddress(), hipGetSymbolSize(), hipMemcpyToSymbol(), hipMemcpyToSymbolAsync(), hipMemcpyFromSymbol(), hipMemcpyFromSymbolAsync()) are available.
### `__shared__`
The `__shared__` keyword is supported.
diff --git a/hipify-clang/README.md b/hipify-clang/README.md
index 431c0a844a..2df3a73bff 100644
--- a/hipify-clang/README.md
+++ b/hipify-clang/README.md
@@ -27,13 +27,14 @@
- [cuRAND](../docs/markdown/CURAND_API_supported_by_HIP.md)
- [cuDNN](../docs/markdown/CUDNN_API_supported_by_HIP.md)
- [cuFFT](../docs/markdown/CUFFT_API_supported_by_HIP.md)
+- [cuSPARSE](../docs/markdown/CUSPARSE_API_supported_by_HIP.md)
## Dependencies
`hipify-clang` requires:
1. LLVM+CLANG of at least version 3.8.0, latest stable and recommended release: 6.0.1 (linux and windows).
-2. CUDA at least version 7.5, latest supported release is 9.2.
+2. CUDA at least version 7.5, latest supported release is 9.0.
| **LLVM release version** | **CUDA latest supported version** | **Comments** |
|:------------------------:|:---------------------------------:|:------------:|
diff --git a/hipify-clang/src/CUDA2HIP.cpp b/hipify-clang/src/CUDA2HIP.cpp
index c0879c1f98..01b23bf007 100644
--- a/hipify-clang/src/CUDA2HIP.cpp
+++ b/hipify-clang/src/CUDA2HIP.cpp
@@ -12,6 +12,8 @@ const std::map CUDA_INCLUDE_MAP{
{"cuda_fp16.h", {"hip/hip_fp16.h", CONV_INCLUDE, API_RUNTIME}},
{"cuda_texture_types.h", {"hip/hip_texture_types.h", CONV_INCLUDE, API_RUNTIME}},
{"vector_types.h", {"hip/hip_vector_types.h", CONV_INCLUDE, API_RUNTIME}},
+ // cuComplex includes
+ {"cuComplex.h", {"hip/hip_complex.h", CONV_INCLUDE_CUDA_MAIN_H, API_COMPLEX}},
// cuBLAS includes
{"cublas.h", {"hipblas.h", CONV_INCLUDE_CUDA_MAIN_H, API_BLAS}},
{"cublas_v2.h", {"hipblas.h", CONV_INCLUDE_CUDA_MAIN_H, API_BLAS}},
@@ -37,8 +39,9 @@ const std::map CUDA_INCLUDE_MAP{
{"cudnn.h", {"hipDNN.h", CONV_INCLUDE_CUDA_MAIN_H, API_DNN}},
// cuFFT includes
{"cufft.h", {"hipfft.h", CONV_INCLUDE_CUDA_MAIN_H, API_FFT}},
- // cuComplex includes
- {"cuComplex.h", {"hip/hip_complex.h", CONV_INCLUDE_CUDA_MAIN_H, API_COMPLEX}},
+ // cuBLAS includes
+ {"cusparse.h", {"hipsparse.h", CONV_INCLUDE_CUDA_MAIN_H, API_SPARSE}},
+ {"cusparse_v2.h", {"hipsparse.h", CONV_INCLUDE_CUDA_MAIN_H, API_SPARSE}},
};
const std::map& CUDA_RENAMES_MAP() {
@@ -61,5 +64,7 @@ const std::map& CUDA_RENAMES_MAP() {
ret.insert(CUDA_DNN_FUNCTION_MAP.begin(), CUDA_DNN_FUNCTION_MAP.end());
ret.insert(CUDA_FFT_TYPE_NAME_MAP.begin(), CUDA_FFT_TYPE_NAME_MAP.end());
ret.insert(CUDA_FFT_FUNCTION_MAP.begin(), CUDA_FFT_FUNCTION_MAP.end());
+ ret.insert(CUDA_SPARSE_TYPE_NAME_MAP.begin(), CUDA_SPARSE_TYPE_NAME_MAP.end());
+ ret.insert(CUDA_SPARSE_FUNCTION_MAP.begin(), CUDA_SPARSE_FUNCTION_MAP.end());
return ret;
};
diff --git a/hipify-clang/src/CUDA2HIP.h b/hipify-clang/src/CUDA2HIP.h
index 5c3a6fa246..637dbf6998 100644
--- a/hipify-clang/src/CUDA2HIP.h
+++ b/hipify-clang/src/CUDA2HIP.h
@@ -37,6 +37,10 @@ extern const std::map CUDA_DNN_FUNCTION_MAP;
extern const std::map CUDA_FFT_TYPE_NAME_MAP;
// Maps the names of CUDA FFT API functions to the corresponding HIP functions
extern const std::map CUDA_FFT_FUNCTION_MAP;
+// Maps the names of CUDA SPARSE API types to the corresponding HIP types
+extern const std::map CUDA_SPARSE_TYPE_NAME_MAP;
+// Maps the names of CUDA SPARSE API functions to the corresponding HIP functions
+extern const std::map CUDA_SPARSE_FUNCTION_MAP;
/**
* The union of all the above maps, except includes.
diff --git a/hipify-clang/src/CUDA2HIP_Complex_API_functions.cpp b/hipify-clang/src/CUDA2HIP_Complex_API_functions.cpp
index 3bc7c4f0a0..1adabc4f5d 100644
--- a/hipify-clang/src/CUDA2HIP_Complex_API_functions.cpp
+++ b/hipify-clang/src/CUDA2HIP_Complex_API_functions.cpp
@@ -1,6 +1,6 @@
#include "CUDA2HIP.h"
-// Maps the names of CUDA DRIVER API types to the corresponding HIP types
+// Maps the names of CUDA Complex API types to the corresponding HIP types
const std::map CUDA_COMPLEX_FUNCTION_MAP{
{"cuCrealf", {"hipCrealf", CONV_COMPLEX, API_COMPLEX}},
{"cuCimagf", {"hipCimagf", CONV_COMPLEX, API_COMPLEX}},
diff --git a/hipify-clang/src/CUDA2HIP_Complex_API_types.cpp b/hipify-clang/src/CUDA2HIP_Complex_API_types.cpp
index f371cf3b9a..6b7eec96fe 100644
--- a/hipify-clang/src/CUDA2HIP_Complex_API_types.cpp
+++ b/hipify-clang/src/CUDA2HIP_Complex_API_types.cpp
@@ -1,6 +1,6 @@
#include "CUDA2HIP.h"
-// Maps the names of CUDA DRIVER API types to the corresponding HIP types
+// Maps the names of CUDA Complex API types to the corresponding HIP types
const std::map CUDA_COMPLEX_TYPE_NAME_MAP{
{"cuFloatComplex", {"hipFloatComplex", CONV_TYPE, API_COMPLEX}},
{"cuDoubleComplex", {"hipDoubleComplex", CONV_TYPE, API_COMPLEX}},
diff --git a/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp b/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp
index 1eaa0903e1..3031cadf64 100644
--- a/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp
+++ b/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp
@@ -69,8 +69,8 @@ const std::map CUDA_RUNTIME_FUNCTION_MAP{
// NOTE: Not equal to cuMipmappedArrayDestroy due to different signatures
{"cudaFreeMipmappedArray", {"hipFreeMipmappedArray", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaGetMipmappedArrayLevel", {"hipGetMipmappedArrayLevel", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
- {"cudaGetSymbolAddress", {"hipGetSymbolAddress", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
- {"cudaGetSymbolSize", {"hipGetSymbolSize", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
+ {"cudaGetSymbolAddress", {"hipGetSymbolAddress", CONV_MEMORY, API_RUNTIME}},
+ {"cudaGetSymbolSize", {"hipGetSymbolSize", CONV_MEMORY, API_RUNTIME}},
// TODO: double check cuMemPrefetchAsync
{"cudaMemPrefetchAsync", {"hipMemPrefetchAsync", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
diff --git a/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp b/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp
new file mode 100644
index 0000000000..9331e297d7
--- /dev/null
+++ b/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp
@@ -0,0 +1,5 @@
+#include "CUDA2HIP.h"
+
+// Maps the names of CUDA SPARSE API types to the corresponding HIP types
+const std::map CUDA_SPARSE_FUNCTION_MAP{
+};
diff --git a/hipify-clang/src/CUDA2HIP_SPARSE_API_types.cpp b/hipify-clang/src/CUDA2HIP_SPARSE_API_types.cpp
new file mode 100644
index 0000000000..ad7ce43be5
--- /dev/null
+++ b/hipify-clang/src/CUDA2HIP_SPARSE_API_types.cpp
@@ -0,0 +1,111 @@
+#include "CUDA2HIP.h"
+
+// Maps the names of CUDA SPARSE API types to the corresponding HIP types
+const std::map CUDA_SPARSE_TYPE_NAME_MAP{
+
+ // 1. Structs
+ {"cusparseContext", {"hipsparseContext", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}},
+ {"cusparseHandle_t", {"hipsparseHandle_t", CONV_TYPE, API_SPARSE}},
+
+ {"cusparseHybMat", {"hipsparseHybMat", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}},
+ {"cusparseHybMat_t", {"hipsparseHybMat_t", CONV_TYPE, API_SPARSE}},
+
+ {"cusparseMatDescr", {"hipsparseMatDescr", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}},
+ {"cusparseMatDescr_t", {"hipsparseMatDescr_t", CONV_TYPE, API_SPARSE}},
+
+ {"cusparseSolveAnalysisInfo", {"hipsparseSolveAnalysisInfo", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}},
+ {"cusparseSolveAnalysisInfo_t", {"hipsparseSolveAnalysisInfo_t", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}},
+
+ {"csrsv2Info", {"csrsv2Info", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}},
+ {"csrsv2Info_t", {"csrsv2Info_t", CONV_TYPE, API_SPARSE}},
+
+ {"csrsm2Info", {"csrsm2Info", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}},
+ {"csrsm2Info_t", {"csrsm2Info_t", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}},
+
+ {"bsrsv2Info", {"bsrsv2Info", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}},
+ {"bsrsv2Info_t", {"bsrsv2Info_t", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}},
+
+ {"bsrsm2Info", {"bsrsm2Info", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}},
+ {"bsrsm2Info_t", {"bsrsm2Info_t", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}},
+
+ {"bsric02Info", {"bsric02Info", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}},
+ {"bsric02Info_t", {"bsric02Info_t", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}},
+
+ {"csrilu02Info", {"csrilu02Info", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}},
+ {"csrilu02Info_t", {"csrilu02Info_t", CONV_TYPE, API_SPARSE}},
+
+ {"bsrilu02Info", {"bsrilu02Info", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}},
+ {"bsrilu02Info_t", {"bsrilu02Info_t", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}},
+
+ {"csru2csrInfo", {"csru2csrInfo", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}},
+ {"csru2csrInfo_t", {"csru2csrInfo_t", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}},
+
+ {"cusparseColorInfo", {"hipsparseColorInfo", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}},
+ {"cusparseColorInfo_t", {"hipsparseColorInfo_t", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}},
+
+ {"pruneInfo", {"pruneInfo", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}},
+ {"pruneInfo_t", {"pruneInfo_t", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}},
+
+ // 2. Enums
+ {"cusparseAction_t", {"hipsparseAction_t", CONV_TYPE, API_SPARSE}},
+ {"CUSPARSE_ACTION_SYMBOLIC", {"HIPSPARSE_ACTION_SYMBOLIC", CONV_NUMERIC_LITERAL, API_SPARSE}},
+ {"CUSPARSE_ACTION_NUMERIC", {"HIPSPARSE_ACTION_NUMERIC", CONV_NUMERIC_LITERAL, API_SPARSE}},
+
+ {"cusparseDirection_t", {"hipsparseDirection_t", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}},
+ {"CUSPARSE_DIRECTION_ROW", {"HIPSPARSE_DIRECTION_ROW", CONV_NUMERIC_LITERAL, API_SPARSE, HIP_UNSUPPORTED}},
+ {"CUSPARSE_DIRECTION_COLUMN", {"HIPSPARSE_DIRECTION_COLUMN", CONV_NUMERIC_LITERAL, API_SPARSE, HIP_UNSUPPORTED}},
+
+ {"cusparseHybPartition_t", {"hipsparseHybPartition_t", CONV_TYPE, API_SPARSE}},
+ {"CUSPARSE_HYB_PARTITION_AUTO", {"HIPSPARSE_HYB_PARTITION_AUTO", CONV_NUMERIC_LITERAL, API_SPARSE}},
+ {"CUSPARSE_HYB_PARTITION_USER", {"HIPSPARSE_HYB_PARTITION_USER", CONV_NUMERIC_LITERAL, API_SPARSE}},
+ {"CUSPARSE_HYB_PARTITION_MAX", {"HIPSPARSE_HYB_PARTITION_MAX", CONV_NUMERIC_LITERAL, API_SPARSE}},
+
+ {"cusparseDiagType_t", {"hipsparseDiagType_t", CONV_TYPE, API_SPARSE}},
+ {"CUSPARSE_DIAG_TYPE_NON_UNIT", {"HIPSPARSE_DIAG_TYPE_NON_UNIT", CONV_NUMERIC_LITERAL, API_SPARSE}},
+ {"CUSPARSE_DIAG_TYPE_UNIT", {"HIPSPARSE_DIAG_TYPE_UNIT", CONV_NUMERIC_LITERAL, API_SPARSE}},
+
+ {"cusparseFillMode_t", {"hipsparseFillMode_t", CONV_TYPE, API_SPARSE}},
+ {"CUSPARSE_FILL_MODE_LOWER", {"HIPSPARSE_FILL_MODE_LOWER", CONV_NUMERIC_LITERAL, API_SPARSE}},
+ {"CUSPARSE_FILL_MODE_UPPER", {"HIPSPARSE_FILL_MODE_UPPER", CONV_NUMERIC_LITERAL, API_SPARSE}},
+
+ {"cusparseIndexBase_t", {"hipsparseIndexBase_t", CONV_TYPE, API_SPARSE}},
+ {"CUSPARSE_INDEX_BASE_ZERO", {"HIPSPARSE_INDEX_BASE_ZERO", CONV_NUMERIC_LITERAL, API_SPARSE}},
+ {"CUSPARSE_INDEX_BASE_ONE", {"HIPSPARSE_INDEX_BASE_ONE", CONV_NUMERIC_LITERAL, API_SPARSE}},
+
+ {"cusparseMatrixType_t", {"hipsparseMatrixType_t", CONV_TYPE, API_SPARSE}},
+ {"CUSPARSE_MATRIX_TYPE_GENERAL", {"HIPSPARSE_MATRIX_TYPE_GENERAL", CONV_NUMERIC_LITERAL, API_SPARSE}},
+ {"CUSPARSE_MATRIX_TYPE_SYMMETRIC", {"HIPSPARSE_MATRIX_TYPE_SYMMETRIC", CONV_NUMERIC_LITERAL, API_SPARSE}},
+ {"CUSPARSE_MATRIX_TYPE_HERMITIAN", {"HIPSPARSE_MATRIX_TYPE_HERMITIAN", CONV_NUMERIC_LITERAL, API_SPARSE}},
+ {"CUSPARSE_MATRIX_TYPE_TRIANGULAR", {"HIPSPARSE_MATRIX_TYPE_TRIANGULAR", CONV_NUMERIC_LITERAL, API_SPARSE}},
+
+ {"cusparseOperation_t", {"hipsparseOperation_t", CONV_TYPE, API_SPARSE}},
+ {"CUSPARSE_OPERATION_NON_TRANSPOSE", {"HIPSPARSE_OPERATION_NON_TRANSPOSE", CONV_NUMERIC_LITERAL, API_SPARSE}},
+ {"CUSPARSE_OPERATION_TRANSPOSE", {"HIPSPARSE_OPERATION_TRANSPOSE", CONV_NUMERIC_LITERAL, API_SPARSE}},
+ {"CUSPARSE_OPERATION_CONJUGATE_TRANSPOSE", {"HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE", CONV_NUMERIC_LITERAL, API_SPARSE}},
+
+ {"cusparsePointerMode_t", {"hipsparsePointerMode_t", CONV_TYPE, API_SPARSE}},
+ {"CUSPARSE_POINTER_MODE_HOST", {"HIPSPARSE_POINTER_MODE_HOST", CONV_NUMERIC_LITERAL, API_SPARSE}},
+ {"CUSPARSE_POINTER_MODE_DEVICE", {"HIPSPARSE_POINTER_MODE_DEVICE", CONV_NUMERIC_LITERAL, API_SPARSE}},
+
+ {"cusparseAlgMode_t", {"hipsparseAlgMode_t", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}},
+ {"CUSPARSE_ALG0", {"CUSPARSE_ALG0", CONV_NUMERIC_LITERAL, API_SPARSE, HIP_UNSUPPORTED}},
+ {"CUSPARSE_ALG1", {"CUSPARSE_ALG1", CONV_NUMERIC_LITERAL, API_SPARSE, HIP_UNSUPPORTED}},
+ {"CUSPARSE_ALG_NAIVE", {"CUSPARSE_ALG_NAIVE", CONV_NUMERIC_LITERAL, API_SPARSE, HIP_UNSUPPORTED}},
+ {"CUSPARSE_ALG_MERGE_PATH", {"CUSPARSE_ALG_MERGE_PATH", CONV_NUMERIC_LITERAL, API_SPARSE, HIP_UNSUPPORTED}},
+
+ {"cusparseSolvePolicy_t", {"hipsparseSolvePolicy_t", CONV_TYPE, API_SPARSE}},
+ {"CUSPARSE_SOLVE_POLICY_NO_LEVEL", {"HIPSPARSE_SOLVE_POLICY_NO_LEVEL", CONV_NUMERIC_LITERAL, API_SPARSE}},
+ {"CUSPARSE_SOLVE_POLICY_USE_LEVEL", {"HIPSPARSE_SOLVE_POLICY_USE_LEVEL", CONV_NUMERIC_LITERAL, API_SPARSE}},
+
+ {"cusparseStatus_t", {"hipsparseMatrixType_t", CONV_TYPE, API_SPARSE}},
+ {"CUSPARSE_STATUS_SUCCESS", {"HIPSPARSE_STATUS_SUCCESS", CONV_NUMERIC_LITERAL, API_SPARSE}},
+ {"CUSPARSE_STATUS_NOT_INITIALIZED", {"HIPSPARSE_STATUS_NOT_INITIALIZED", CONV_NUMERIC_LITERAL, API_SPARSE}},
+ {"CUSPARSE_STATUS_ALLOC_FAILED", {"HIPSPARSE_STATUS_ALLOC_FAILED", CONV_NUMERIC_LITERAL, API_SPARSE}},
+ {"CUSPARSE_STATUS_INVALID_VALUE", {"HIPSPARSE_STATUS_INVALID_VALUE", CONV_NUMERIC_LITERAL, API_SPARSE}},
+ {"CUSPARSE_STATUS_ARCH_MISMATCH", {"HIPSPARSE_STATUS_ARCH_MISMATCH", CONV_NUMERIC_LITERAL, API_SPARSE}},
+ {"CUSPARSE_STATUS_MAPPING_ERROR", {"HIPSPARSE_STATUS_MAPPING_ERROR", CONV_NUMERIC_LITERAL, API_SPARSE}},
+ {"CUSPARSE_STATUS_EXECUTION_FAILED", {"HIPSPARSE_STATUS_EXECUTION_FAILED", CONV_NUMERIC_LITERAL, API_SPARSE}},
+ {"CUSPARSE_STATUS_INTERNAL_ERROR", {"HIPSPARSE_STATUS_INTERNAL_ERROR", CONV_NUMERIC_LITERAL, API_SPARSE}},
+ {"CUSPARSE_STATUS_MATRIX_TYPE_NOT_SUPPORTED", {"HIPSPARSE_STATUS_MATRIX_TYPE_NOT_SUPPORTED", CONV_NUMERIC_LITERAL, API_SPARSE}},
+ {"CUSPARSE_STATUS_ZERO_PIVOT", {"HIPSPARSE_STATUS_ZERO_PIVOT", CONV_NUMERIC_LITERAL, API_SPARSE}},
+};
diff --git a/hipify-clang/src/HipifyAction.cpp b/hipify-clang/src/HipifyAction.cpp
index b370df794e..71e7b759ac 100644
--- a/hipify-clang/src/HipifyAction.cpp
+++ b/hipify-clang/src/HipifyAction.cpp
@@ -155,6 +155,10 @@ bool HipifyAction::Exclude(const hipCounter & hipToken) {
if (insertedComplexHeader) { return true; }
insertedComplexHeader = true;
return false;
+ case API_SPARSE:
+ if (insertedSPARSEHeader) { return true; }
+ insertedSPARSEHeader = true;
+ return false;
default:
return false;
}
diff --git a/hipify-clang/src/HipifyAction.h b/hipify-clang/src/HipifyAction.h
index 9d30a72592..99cb52e74c 100644
--- a/hipify-clang/src/HipifyAction.h
+++ b/hipify-clang/src/HipifyAction.h
@@ -29,6 +29,7 @@ private:
bool insertedRAND_kernelHeader = false;
bool insertedDNNHeader = false;
bool insertedFFTHeader = false;
+ bool insertedSPARSEHeader = false;
bool insertedComplexHeader = false;
bool firstHeader = false;
bool pragmaOnce = false;
diff --git a/hipify-clang/src/Statistics.cpp b/hipify-clang/src/Statistics.cpp
index 2d37c3156c..95aa15c5ce 100644
--- a/hipify-clang/src/Statistics.cpp
+++ b/hipify-clang/src/Statistics.cpp
@@ -43,7 +43,14 @@ const char *counterNames[NUM_CONV_TYPES] = {
};
const char *apiNames[NUM_API_TYPES] = {
- "CUDA Driver API", "CUDA RT API", "CUBLAS API", "CURAND API", "CUDNN API", "CUFFT API", "cuComplex API"
+ "CUDA Driver API",
+ "CUDA RT API",
+ "cuComplex API",
+ "cuBLAS API",
+ "cuRAND API",
+ "cuDNN API",
+ "cuFFT API",
+ "cuSPARSE"
};
namespace {
diff --git a/hipify-clang/src/Statistics.h b/hipify-clang/src/Statistics.h
index 0ce8e0de67..7c81a15e4d 100644
--- a/hipify-clang/src/Statistics.h
+++ b/hipify-clang/src/Statistics.h
@@ -103,11 +103,12 @@ constexpr int NUM_CONV_TYPES = (int) ConvTypes::CONV_LAST;
enum ApiTypes {
API_DRIVER = 0,
API_RUNTIME,
+ API_COMPLEX,
API_BLAS,
API_RAND,
API_DNN,
API_FFT,
- API_COMPLEX,
+ API_SPARSE,
API_LAST
};
constexpr int NUM_API_TYPES = (int) ApiTypes::API_LAST;
diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h
index c9ff32d197..3567a67854 100644
--- a/include/hip/hcc_detail/hip_runtime_api.h
+++ b/include/hip/hcc_detail/hip_runtime_api.h
@@ -1385,6 +1385,32 @@ hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, size_t siz
size_t offset __dparm(0), hipMemcpyKind kind __dparm(hipMemcpyHostToDevice));
+/**
+ * @brief Copies the memory address of symbol @p symbolName to @p devPtr
+ *
+ * @param[in] symbolName - Symbol on device
+ * @param[out] devPtr - Pointer to a pointer to the memory referred to by the symbol
+ * @return #hipSuccess, #hipErrorNotInitialized, #hipErrorNotFound
+ *
+ * @see hipGetSymbolSize, hipMemcpyToSymbol, hipMemcpyFromSymbol, hipMemcpyToSymbolAsync,
+ * hipMemcpyFromSymbolAsync
+ */
+hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName);
+
+
+/**
+ * @brief Copies the size of symbol @p symbolName to @p size
+ *
+ * @param[in] symbolName - Symbol on device
+ * @param[out] size - Pointer to the size of the symbol
+ * @return #hipSuccess, #hipErrorNotInitialized, #hipErrorNotFound
+ *
+ * @see hipGetSymbolSize, hipMemcpyToSymbol, hipMemcpyFromSymbol, hipMemcpyToSymbolAsync,
+ * hipMemcpyFromSymbolAsync
+ */
+hipError_t hipGetSymbolSize(size_t* size, const void* symbolName);
+
+
/**
* @brief Copies @p sizeBytes bytes from the memory area pointed to by @p src to the memory area
* pointed to by @p offset bytes from the start of symbol @p symbol
@@ -2334,6 +2360,9 @@ hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func);
hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t hmod,
const char* name);
+hipError_t ihipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t hmod,
+ const char* name);
+
hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const char* name);
/**
* @brief builds module from code object which resides in host memory. Image is pointer to that
diff --git a/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h
index e68e8ac328..02c4b7ee61 100644
--- a/include/hip/nvcc_detail/hip_runtime_api.h
+++ b/include/hip/nvcc_detail/hip_runtime_api.h
@@ -551,6 +551,14 @@ inline static hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolN
dst, symbolName, sizeBytes, offset, hipMemcpyKindToCudaMemcpyKind(kind), stream));
}
+inline static hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName) {
+ return hipCUDAErrorTohipError(cudaGetSymbolAddress(devPtr, symbolName));
+}
+
+inline static hipError_t hipGetSymbolSize(size_t* size, const void* symbolName) {
+ return hipCUDAErrorTohipError(cudaGetSymbolSize(size, symbolName));
+}
+
inline static hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch,
size_t width, size_t height, hipMemcpyKind kind) {
return hipCUDAErrorTohipError(
diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp
index 7c25b714f8..6a6bed395f 100644
--- a/src/hip_memory.cpp
+++ b/src/hip_memory.cpp
@@ -909,14 +909,22 @@ hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags)
}
am_status = hc::am_memory_host_lock(device->_acc, hostPtr, sizeBytes, &vecAcc[0],
vecAcc.size());
-#if USE_APP_PTR_FOR_CTX
- hc::am_memtracker_update(hostPtr, device->_deviceId, flags, ctx);
-#else
- hc::am_memtracker_update(hostPtr, device->_deviceId, flags);
-#endif
+ if ( am_status == AM_SUCCESS ) {
+ am_status = hc::am_memtracker_getinfo(&amPointerInfo, hostPtr);
- tprintf(DB_MEM, " %s registered ptr=%p and allowed access to %zu peers\n", __func__,
- hostPtr, vecAcc.size());
+ if ( am_status == AM_SUCCESS ) {
+ void *devPtr = amPointerInfo._devicePointer;
+ #if USE_APP_PTR_FOR_CTX
+ hc::am_memtracker_update(hostPtr, device->_deviceId, flags, ctx);
+ hc::am_memtracker_update(devPtr, device->_deviceId, flags, ctx);
+ #else
+ hc::am_memtracker_update(hostPtr, device->_deviceId, flags);
+ hc::am_memtracker_update(devPtr, device->_deviceId, flags);
+ #endif
+ tprintf(DB_MEM, " %s registered ptr=%p and allowed access to %zu peers\n", __func__,
+ hostPtr, vecAcc.size());
+ };
+ };
if (am_status == AM_SUCCESS) {
hip_status = hipSuccess;
} else {
@@ -953,7 +961,7 @@ inline hipDeviceptr_t agent_address_for_symbol(const char* symbolName) {
#if __hcc_workweek__ >= 17481
size_t byte_cnt = 0u;
- hipModuleGetGlobal(&r, &byte_cnt, 0, symbolName);
+ ihipModuleGetGlobal(&r, &byte_cnt, 0, symbolName);
#else
auto ctx = ihipGetTlsDefaultCtx();
auto acc = ctx->getDevice()->_acc;
@@ -1098,6 +1106,23 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t co
return ihipLogStatus(e);
}
+
+hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName) {
+ HIP_INIT_API(devPtr, symbolName);
+
+ size_t size = 0;
+ return ihipModuleGetGlobal(devPtr, &size, 0, static_cast(symbolName));
+}
+
+
+hipError_t hipGetSymbolSize(size_t* size, const void* symbolName) {
+ HIP_INIT_API(size, symbolName);
+
+ void* devPtr = nullptr;
+ return ihipModuleGetGlobal(&devPtr, size, 0, static_cast(symbolName));
+}
+
+
//---
hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, kind);
diff --git a/src/hip_module.cpp b/src/hip_module.cpp
index 019bafbe43..786d1e8d5c 100644
--- a/src/hip_module.cpp
+++ b/src/hip_module.cpp
@@ -472,14 +472,19 @@ hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t h
const char* name) {
HIP_INIT_API(dptr, bytes, hmod, name);
- if (!dptr || !bytes) return ihipLogStatus(hipErrorInvalidValue);
+ return ihipLogStatus(ihipModuleGetGlobal(dptr, bytes, hmod, name));
+}
- if (!name) return ihipLogStatus(hipErrorNotInitialized);
+hipError_t ihipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t hmod,
+ const char* name) {
+ if (!dptr || !bytes) return hipErrorInvalidValue;
+
+ if (!name) return hipErrorNotInitialized;
const auto r = hmod ? read_agent_global_from_module(dptr, bytes, hmod, name)
: read_agent_global_from_process(dptr, bytes, name);
- return ihipLogStatus(r);
+ return r;
}
namespace
diff --git a/tests/hipify-clang/headers_test_09.cu b/tests/hipify-clang/headers_test_09.cu
index 5dbfaffc2d..13d2b688f9 100644
--- a/tests/hipify-clang/headers_test_09.cu
+++ b/tests/hipify-clang/headers_test_09.cu
@@ -48,6 +48,7 @@
// CHECK: #include
// CHECK: #include "hipfft.h"
+// CHECK: #include "hipsparse.h"
#include
@@ -95,3 +96,5 @@
#include
#include "cufft.h"
+
+#include "cusparse.h"
diff --git a/tests/src/deviceLib/hipTestDeviceSymbol.cpp b/tests/src/deviceLib/hipTestDeviceSymbol.cpp
index 4bac9a902b..a17dd75a31 100644
--- a/tests/src/deviceLib/hipTestDeviceSymbol.cpp
+++ b/tests/src/deviceLib/hipTestDeviceSymbol.cpp
@@ -40,6 +40,12 @@ __global__ void Assign(int* Out) {
globalOut[tid] = globalIn[tid];
}
+__device__ __constant__ int globalConst[NUM];
+
+__global__ void checkAddress(int* addr, bool* out) {
+ *out = (globalConst == addr);
+}
+
int main() {
int *A, *Am, *B, *Ad, *C, *Cm;
A = new int[NUM];
@@ -101,6 +107,20 @@ int main() {
assert(A[i] == B[i]);
assert(A[i] == C[i]);
}
+
+ bool *checkOkD;
+ bool checkOk = false;
+ size_t symbolSize = 0;
+ int *symbolAddress;
+ hipGetSymbolSize(&symbolSize, HIP_SYMBOL(globalConst));
+ hipGetSymbolAddress((void**) &symbolAddress, HIP_SYMBOL(globalConst));
+ hipMalloc((void**)&checkOkD, sizeof(bool));
+ hipLaunchKernelGGL(checkAddress, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, symbolAddress, checkOkD);
+ hipMemcpy(&checkOk, checkOkD, sizeof(bool), hipMemcpyDeviceToHost);
+ hipFree(checkOkD);
+ assert(checkOk);
+ assert(symbolSize == SIZE);
+
hipHostFree(Am);
hipHostFree(Cm);
hipFree(Ad);