AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

Merge 'master' into 'amd-master'

Просмотр исходного кода 
Change-Id: I26e2e1777867aab4e5f8d8a27c8bb006a30f867f
Этот коммит содержится в:
Jenkins
2018-11-21 04:09:43 -06:00
родитель 09e5eb003e bcea027bf1
Коммит 9def89b4fe
23 изменённых файлов: 347 добавлений и 28 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
Jenkinsfile
поставляемый
+1 -3
Просмотреть файл
@@ -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
bin/hipify-perl
+3
Просмотреть файл
@@ -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;
docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md
+4 -4
Просмотреть файл
@@ -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` | |
docs/markdown/CUSPARSE_API_supported_by_HIP.md
+89
Просмотреть файл
@@ -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** |
|-----------------------------------------------------------|-------------------------------------------------|
docs/markdown/hip_kernel_language.md
+1 -1
Просмотреть файл
@@ -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.
hipify-clang/README.md
+2 -1
Просмотреть файл
@@ -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** |
|:------------------------:|:---------------------------------:|:------------:|
hipify-clang/src/CUDA2HIP.cpp
+7 -2
Просмотреть файл
@@ -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;
};
hipify-clang/src/CUDA2HIP.h
+4
Просмотреть файл
@@ -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.
hipify-clang/src/CUDA2HIP_Complex_API_functions.cpp
+1 -1
Просмотреть файл
@@ -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}},
hipify-clang/src/CUDA2HIP_Complex_API_types.cpp
+1 -1
Просмотреть файл
@@ -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}},
hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp
+2 -2
Просмотреть файл
@@ -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}},
hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp
+5
Просмотреть файл
@@ -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{
};
hipify-clang/src/CUDA2HIP_SPARSE_API_types.cpp
+111
Просмотреть файл
@@ -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}},
};
hipify-clang/src/HipifyAction.cpp
+4
Просмотреть файл
@@ -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;
}
hipify-clang/src/HipifyAction.h
+1
Просмотреть файл
@@ -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;
hipify-clang/src/Statistics.cpp
+8 -1
Просмотреть файл
@@ -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 {
hipify-clang/src/Statistics.h
+2 -1
Просмотреть файл
@@ -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;
include/hip/hcc_detail/hip_runtime_api.h
+29
Просмотреть файл
@@ -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
include/hip/nvcc_detail/hip_runtime_api.h
+8
Просмотреть файл
@@ -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(
src/hip_memory.cpp
+33 -8
Просмотреть файл
@@ -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);
src/hip_module.cpp
+8 -3
Просмотреть файл
@@ -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
tests/hipify-clang/headers_test_09.cu
+3
Просмотреть файл
@@ -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"
tests/src/deviceLib/hipTestDeviceSymbol.cpp
+20
Просмотреть файл
@@ -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);