AI/rocm-systems
2
0
Derivar 0
Código Questões Pedidos de integração Operações Pacotes Planeamentos Lançamentos Wiki Trabalho

Merge branch 'master' of https://github.com/ROCm-Developer-Tools/HIP

Ver a proveniência
Este cometimento está contido em:
Qianfeng Zhang
2018-11-22 19:11:49 +08:00
ascendente b88c8769ea 40d3184dd1
cometimento 4f85e98384
22 ficheiros modificados com 346 adições e 25 eliminações
Descarregar ficheiro patch Descarregar ficheiro diff Expandir todos os ficheiros Colapsar todos os ficheiros
hipamd/bin/hipify-perl
+3
Ver ficheiro
@@ -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;
hipamd/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md
+4 -4
Ver ficheiro
@@ -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` | |
hipamd/docs/markdown/CUSPARSE_API_supported_by_HIP.md
+89
Ver ficheiro
@@ -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** |
|-----------------------------------------------------------|-------------------------------------------------|
hipamd/docs/markdown/hip_kernel_language.md
+1 -1
Ver ficheiro
@@ -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.
hipamd/hipify-clang/README.md
+2 -1
Ver ficheiro
@@ -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)
## <a name="dependencies"></a> 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** |
|:------------------------:|:---------------------------------:|:------------:|
hipamd/hipify-clang/src/CUDA2HIP.cpp
+7 -2
Ver ficheiro
@@ -12,6 +12,8 @@ const std::map <llvm::StringRef, hipCounter> 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 <llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter>& CUDA_RENAMES_MAP() {
@@ -61,5 +64,7 @@ const std::map<llvm::StringRef, hipCounter>& 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;
};
hipamd/hipify-clang/src/CUDA2HIP.h
+4
Ver ficheiro
@@ -37,6 +37,10 @@ extern const std::map<llvm::StringRef, hipCounter> CUDA_DNN_FUNCTION_MAP;
extern const std::map<llvm::StringRef, hipCounter> CUDA_FFT_TYPE_NAME_MAP;
// Maps the names of CUDA FFT API functions to the corresponding HIP functions
extern const std::map<llvm::StringRef, hipCounter> CUDA_FFT_FUNCTION_MAP;
// Maps the names of CUDA SPARSE API types to the corresponding HIP types
extern const std::map<llvm::StringRef, hipCounter> CUDA_SPARSE_TYPE_NAME_MAP;
// Maps the names of CUDA SPARSE API functions to the corresponding HIP functions
extern const std::map<llvm::StringRef, hipCounter> CUDA_SPARSE_FUNCTION_MAP;
/**
* The union of all the above maps, except includes.
hipamd/hipify-clang/src/CUDA2HIP_Complex_API_functions.cpp
+1 -1
Ver ficheiro
@@ -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<llvm::StringRef, hipCounter> CUDA_COMPLEX_FUNCTION_MAP{
{"cuCrealf", {"hipCrealf", CONV_COMPLEX, API_COMPLEX}},
{"cuCimagf", {"hipCimagf", CONV_COMPLEX, API_COMPLEX}},
hipamd/hipify-clang/src/CUDA2HIP_Complex_API_types.cpp
+1 -1
Ver ficheiro
@@ -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<llvm::StringRef, hipCounter> CUDA_COMPLEX_TYPE_NAME_MAP{
{"cuFloatComplex", {"hipFloatComplex", CONV_TYPE, API_COMPLEX}},
{"cuDoubleComplex", {"hipDoubleComplex", CONV_TYPE, API_COMPLEX}},
hipamd/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp
+2 -2
Ver ficheiro
@@ -69,8 +69,8 @@ const std::map<llvm::StringRef, hipCounter> 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}},
hipamd/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp
+5
Ver ficheiro
@@ -0,0 +1,5 @@
#include "CUDA2HIP.h"
// Maps the names of CUDA SPARSE API types to the corresponding HIP types
const std::map<llvm::StringRef, hipCounter> CUDA_SPARSE_FUNCTION_MAP{
};
hipamd/hipify-clang/src/CUDA2HIP_SPARSE_API_types.cpp
+111
Ver ficheiro
@@ -0,0 +1,111 @@
#include "CUDA2HIP.h"
// Maps the names of CUDA SPARSE API types to the corresponding HIP types
const std::map<llvm::StringRef, hipCounter> 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}},
};
hipamd/hipify-clang/src/HipifyAction.cpp
+4
Ver ficheiro
@@ -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;
}
hipamd/hipify-clang/src/HipifyAction.h
+1
Ver ficheiro
@@ -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;
hipamd/hipify-clang/src/Statistics.cpp
+8 -1
Ver ficheiro
@@ -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 {
hipamd/hipify-clang/src/Statistics.h
+2 -1
Ver ficheiro
@@ -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;
hipamd/include/hip/hcc_detail/hip_runtime_api.h
+29
Ver ficheiro
@@ -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
hipamd/include/hip/nvcc_detail/hip_runtime_api.h
+8
Ver ficheiro
@@ -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(
hipamd/src/hip_memory.cpp
+33 -8
Ver ficheiro
@@ -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<const char*>(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<const char*>(symbolName));
}
//---
hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, kind);
hipamd/src/hip_module.cpp
+8 -3
Ver ficheiro
@@ -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
hipamd/tests/hipify-clang/headers_test_09.cu
+3
Ver ficheiro
@@ -48,6 +48,7 @@
// CHECK: #include <string>
// CHECK: #include "hipfft.h"
// CHECK: #include "hipsparse.h"
#include <cuda.h>
@@ -95,3 +96,5 @@
#include <string>
#include "cufft.h"
#include "cusparse.h"
hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp
+20
Ver ficheiro
@@ -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);