From e126f7254c8efddcaff96e8392685ccc8c74e797 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Tue, 2 Apr 2019 17:28:06 -0400 Subject: [PATCH 01/11] minor workaround for lazy binding --- include/hip/hcc_detail/code_object_bundle.hpp | 9 +++++---- src/code_object_bundle.cpp | 8 +------- 2 files changed, 6 insertions(+), 11 deletions(-) diff --git a/include/hip/hcc_detail/code_object_bundle.hpp b/include/hip/hcc_detail/code_object_bundle.hpp index 89b341afa7..8b1ae8aa49 100644 --- a/include/hip/hcc_detail/code_object_bundle.hpp +++ b/include/hip/hcc_detail/code_object_bundle.hpp @@ -89,9 +89,10 @@ struct Bundled_code { std::vector blob; }; +#define magic_string_ "__CLANG_OFFLOAD_BUNDLE__" + class Bundled_code_header { // DATA - STATICS - static constexpr const char magic_string_[] = "__CLANG_OFFLOAD_BUNDLE__"; static constexpr auto magic_string_sz_ = sizeof(magic_string_) - 1; // DATA @@ -143,12 +144,12 @@ class Bundled_code_header { std::vector{std::istreambuf_iterator{is}, std::istreambuf_iterator{}}, x); } - // FRIENDS - ACCESSORS friend inline bool valid(const Bundled_code_header& x) { - return std::equal(magic_string_, magic_string_ + magic_string_sz_, - x.header_.bundler_magic_string_); + const std::string ms = {magic_string_}; + return std::equal(ms.begin(), ms.end(), x.header_.bundler_magic_string_); } + friend inline const std::vector& bundles(const Bundled_code_header& x) { return x.bundles_; } diff --git a/src/code_object_bundle.cpp b/src/code_object_bundle.cpp index 6285cb3828..feef90a61a 100644 --- a/src/code_object_bundle.cpp +++ b/src/code_object_bundle.cpp @@ -10,9 +10,6 @@ using namespace std; -// DATA - STATICS -constexpr const char hip_impl::Bundled_code_header::magic_string_[]; - // CREATORS hip_impl::Bundled_code_header::Bundled_code_header(const vector& x) : Bundled_code_header{x.cbegin(), x.cend()} {} @@ -22,12 +19,9 @@ hip_impl::Bundled_code_header::Bundled_code_header( // hipLoadModuleData is so poorly specified (for no fault of its own). if (!p) return; + if (!valid(*static_cast(p))) return; auto ph = static_cast(p); - if (!equal(magic_string_, magic_string_ + magic_string_sz_, ph->bundler_magic_string_)) { - return; - } - size_t sz = sizeof(Header_) + ph->bundle_cnt_ * sizeof(Bundled_code::Header); auto pb = static_cast(p) + sizeof(Header_); auto n = ph->bundle_cnt_; From c61f265657b30e3460f3107bc6cc41ddcdaa61ad Mon Sep 17 00:00:00 2001 From: David Salinas Date: Fri, 29 Mar 2019 17:57:17 -0400 Subject: [PATCH 02/11] append the ELF flags for sram-ecc and xnack to the target triple per code object --- include/hip/hcc_detail/code_object_bundle.hpp | 14 ++++++++++++++ 1 file changed, 14 insertions(+) diff --git a/include/hip/hcc_detail/code_object_bundle.hpp b/include/hip/hcc_detail/code_object_bundle.hpp index 89b341afa7..173a915409 100644 --- a/include/hip/hcc_detail/code_object_bundle.hpp +++ b/include/hip/hcc_detail/code_object_bundle.hpp @@ -31,6 +31,8 @@ THE SOFTWARE. #include #include #include +#include +#include "elfio/elfio.hpp" namespace hip_impl { @@ -124,6 +126,18 @@ class Bundled_code_header { std::copy_n(f + y.header.offset, y.header.bundle_sz, std::back_inserter(y.blob)); + // we need to read the blob (ELF) for this bundle to get the feature flags + std::stringstream tmp{std::string{y.blob.cbegin(), y.blob.cend()}}; + ELFIO::elfio reader; + if (reader.load(tmp)) { + unsigned int xnack_enabled = reader.get_flags() & 0x100; + unsigned int sram_ecc_enabled = reader.get_flags() & 0x200; + if (xnack_enabled) + y.triple.append("+xnack"); + if (sram_ecc_enabled) + y.triple.append("+sram-ecc"); + } + it += y.header.triple_sz; x.bundled_code_size = std::max(x.bundled_code_size, From 5a0e14fbb1dde157d4fa5f6e50c22b9626a77e44 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 15 Apr 2019 12:38:01 +0530 Subject: [PATCH 03/11] [ci] Enable tests on ROCm 2.3 Change-Id: Id344ef600b0868f36f2e7ac08d5664234d88835b --- Jenkinsfile | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/Jenkinsfile b/Jenkinsfile index 5e6ee065fd..e94585ceb5 100644 --- a/Jenkinsfile +++ b/Jenkinsfile @@ -296,13 +296,13 @@ def docker_upload_dockerhub( String local_org, String image_name, String remote_ String build_config = 'Release' String job_name = env.JOB_NAME.toLowerCase( ) -// The following launches 3 builds in parallel: rocm-head, rocm-2.1.x and cuda-10.x -parallel rocm_2_1: +// The following launches 3 builds in parallel: rocm-head, rocm-2.3.x and cuda-10.x +parallel rocm_2_3: { node('hip-rocm') { - String hcc_ver = 'rocm-2.1.x' - String from_image = 'ci_test_nodes/rocm-2.1.x/ubuntu-16.04:latest' + String hcc_ver = 'rocm-2.3.x' + String from_image = 'ci_test_nodes/rocm-2.3.x/ubuntu-16.04:latest' String inside_args = '--device=/dev/kfd --device=/dev/dri --group-add=video' // Checkout source code, dependencies and version files From d3a810ab11c76a2f93088af1304d793bdd9835d2 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 15 Apr 2019 15:46:46 +0300 Subject: [PATCH 04/11] [HIPIFY] cuDNN 7.5.0.56 support --- docs/markdown/CUDNN_API_supported_by_HIP.md | 52 +++++++++++++++++++ .../src/CUDA2HIP_DNN_API_functions.cpp | 24 +++++++++ hipify-clang/src/CUDA2HIP_DNN_API_types.cpp | 32 ++++++++++++ 3 files changed, 108 insertions(+) diff --git a/docs/markdown/CUDNN_API_supported_by_HIP.md b/docs/markdown/CUDNN_API_supported_by_HIP.md index 05c49e2bff..be0fd442fc 100644 --- a/docs/markdown/CUDNN_API_supported_by_HIP.md +++ b/docs/markdown/CUDNN_API_supported_by_HIP.md @@ -47,6 +47,8 @@ | struct* |`cudnnReduceTensorDescriptor_t` |`hipdnnReduceTensorDescriptor_t` | | struct |`cudnnCTCLossStruct` | | | struct* |`cudnnCTCLossDescriptor_t` | | +| struct |`cudnnTensorTransformStruct` | | +| struct* |`cudnnTensorTransformDescriptor_t` | | | enum |***`cudnnDataType_t`*** |***`hipdnnDataType_t`*** | | 0 |*`CUDNN_DATA_FLOAT`* |*`HIPDNN_DATA_FLOAT`* | | 1 |*`CUDNN_DATA_DOUBLE`* |*`HIPDNN_DATA_DOUBLE`* | @@ -71,6 +73,9 @@ | 0 |*`CUDNN_TENSOR_NCHW`* |*`HIPDNN_TENSOR_NCHW`* | | 1 |*`CUDNN_TENSOR_NHWC`* |*`HIPDNN_TENSOR_NHWC`* | | 2 |*`CUDNN_TENSOR_NCHW_VECT_C`* |*`HIPDNN_TENSOR_NCHW_VECT_C`* | +| enum |***`cudnnFoldingDirection_t`*** | | +| 0 |*`CUDNN_TRANSFORM_FOLD`* | | +| 1 |*`CUDNN_TRANSFORM_UNFOLD`* | | | enum |***`cudnnOpTensorOp_t`*** |***`hipdnnOpTensorOp_t`*** | | 0 |*`CUDNN_OP_TENSOR_ADD`* |*`HIPDNN_OP_TENSOR_ADD`* | | 1 |*`CUDNN_OP_TENSOR_MUL`* |*`HIPDNN_OP_TENSOR_MUL`* | @@ -169,6 +174,11 @@ | 1 |*`CUDNN_RNN_TANH`* |*`HIPDNN_RNN_TANH`* | | 2 |*`CUDNN_LSTM`* |*`HIPDNN_LSTM`* | | 3 |*`CUDNN_GRU`* |*`HIPDNN_GRU`* | +| enum |***`cudnnRNNBiasMode_t`*** | | +| 0 |*`CUDNN_RNN_NO_BIAS`* | | +| 1 |*`CUDNN_RNN_SINGLE_INP_BIAS`* | | +| 2 |*`CUDNN_RNN_DOUBLE_BIAS`* | | +| 3 |*`CUDNN_RNN_SINGLE_REC_BIAS`* | | | enum |***`cudnnDirectionMode_t`*** |***`hipdnnDirectionMode_t`*** | | 0 |*`CUDNN_UNIDIRECTIONAL`* |*`HIPDNN_UNIDIRECTIONAL`* | | 1 |*`CUDNN_BIDIRECTIONAL`* |*`HIPDNN_BIDIRECTIONAL`* | @@ -215,6 +225,27 @@ | enum |***`cudnnRNNPaddingMode_t`*** | | | 0 |*`CUDNN_RNN_PADDED_IO_DISABLED`* | | | 1 |*`CUDNN_RNN_PADDED_IO_ENABLED`* | | +| enum |***`cudnnSeqDataAxis_t`*** | | +| 0 |*`CUDNN_SEQDATA_TIME_DIM`* | | +| 1 |*`CUDNN_SEQDATA_BATCH_DIM`* | | +| 2 |*`CUDNN_SEQDATA_BEAM_DIM`* | | +| 3 |*`CUDNN_SEQDATA_VECT_DIM`* | | +| define |`CUDNN_SEQDATA_DIM_COUNT` | | +| struct |`cudnnSeqDataStruct` | | +| struct* |`cudnnSeqDataDescriptor_t` | | +| enum |***`cudnnAttnQueryMap_t`*** | | +| 0 |*`CUDNN_ATTN_QUERYMAP_ALL_TO_ONE`* | | +| 1 |*`CUDNN_ATTN_QUERYMAP_ONE_TO_ONE`* | | +| struct |`cudnnAttnStruct` | | +| struct* |`cudnnAttnDescriptor_t` | | +| enum |***`cudnnMultiHeadAttnWeightKind_t`*** | | +| 0 |*`CUDNN_MH_ATTN_Q_WEIGHTS`* | | +| 1 |*`CUDNN_MH_ATTN_K_WEIGHTS`* | | +| 2 |*`CUDNN_MH_ATTN_V_WEIGHTS`* | | +| 3 |*`CUDNN_MH_ATTN_O_WEIGHTS`* | | +| enum |***`cudnnWgradMode_t`*** | | +| 0 |*`CUDNN_WGRAD_MODE_ADD`* | | +| 1 |*`CUDNN_WGRAD_MODE_SET`* | | ## **2. CUDNN API functions** @@ -239,6 +270,12 @@ |`cudnnGetTensorSizeInBytes` | | |`cudnnDestroyTensorDescriptor` |`hipdnnDestroyTensorDescriptor` | |`cudnnTransformTensor` | | +|`cudnnTransformTensorEx` | | +|`cudnnInitTransformDest` | | +|`cudnnCreateTensorTransformDescriptor` | | +|`cudnnSetTensorTransformDescriptor` | | +|`cudnnGetTensorTransformDescriptor` | | +|`cudnnDestroyTensorTransformDescriptor` | | |`cudnnAddTensor` |`hipdnnAddTensor` | |`cudnnCreateOpTensorDescriptor` |`hipdnnCreateOpTensorDescriptor` | |`cudnnSetOpTensorDescriptor` |`hipdnnSetOpTensorDescriptor` | @@ -383,6 +420,8 @@ |`cudnnDestroyRNNDataDescriptor` | | |`cudnnSetRNNDataDescriptor` | | |`cudnnGetRNNDataDescriptor` | | +|`cudnnSetRNNBiasMode` | | +|`cudnnGetRNNBiasMode` | | |`cudnnCreateCTCLossDescriptor` | | |`cudnnSetCTCLossDescriptor` | | |`cudnnGetCTCLossDescriptor` | | @@ -410,3 +449,16 @@ |`cudnnGetBatchNormalizationTrainingExReserveSpaceSize` | | |`cudnnRNNSetClip` | | |`cudnnRNNGetClip` | | +|`cudnnCreateSeqDataDescriptor` | | +|`cudnnDestroySeqDataDescriptor` | | +|`cudnnSetSeqDataDescriptor` | | +|`cudnnGetSeqDataDescriptor` | | +|`cudnnCreateAttnDescriptor` | | +|`cudnnDestroyAttnDescriptor` | | +|`cudnnSetAttnDescriptor` | | +|`cudnnGetAttnDescriptor` | | +|`cudnnGetMultiHeadAttnBuffers` | | +|`cudnnGetMultiHeadAttnWeights` | | +|`cudnnMultiHeadAttnForward` | | +|`cudnnMultiHeadAttnBackwardData` | | +|`cudnnMultiHeadAttnBackwardWeights` | | diff --git a/hipify-clang/src/CUDA2HIP_DNN_API_functions.cpp b/hipify-clang/src/CUDA2HIP_DNN_API_functions.cpp index 880e03be79..d99a201b8e 100644 --- a/hipify-clang/src/CUDA2HIP_DNN_API_functions.cpp +++ b/hipify-clang/src/CUDA2HIP_DNN_API_functions.cpp @@ -49,6 +49,12 @@ const std::map CUDA_DNN_FUNCTION_MAP{ {"cudnnGetTensorSizeInBytes", {"hipdnnGetTensorSizeInBytes", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, {"cudnnDestroyTensorDescriptor", {"hipdnnDestroyTensorDescriptor", "", CONV_LIB_FUNC, API_DNN}}, {"cudnnTransformTensor", {"hipdnnTransformTensor", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnTransformTensorEx", {"hipdnnTransformTensorEx", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnInitTransformDest", {"hipdnnInitTransformDest", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnCreateTensorTransformDescriptor", {"hipdnnCreateTensorTransformDescriptor", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnSetTensorTransformDescriptor", {"hipdnnSetTensorTransformDescriptor", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnGetTensorTransformDescriptor", {"hipdnnGetTensorTransformDescriptor", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnDestroyTensorTransformDescriptor", {"hipdnnDestroyTensorTransformDescriptor", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, {"cudnnAddTensor", {"hipdnnAddTensor", "", CONV_LIB_FUNC, API_DNN}}, {"cudnnCreateOpTensorDescriptor", {"hipdnnCreateOpTensorDescriptor", "", CONV_LIB_FUNC, API_DNN}}, {"cudnnSetOpTensorDescriptor", {"hipdnnSetOpTensorDescriptor", "", CONV_LIB_FUNC, API_DNN}}, @@ -223,6 +229,8 @@ const std::map CUDA_DNN_FUNCTION_MAP{ {"cudnnDestroyRNNDataDescriptor", {"hipdnnDestroyRNNDataDescriptor", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, {"cudnnSetRNNDataDescriptor", {"hipdnnSetRNNDataDescriptor", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, {"cudnnGetRNNDataDescriptor", {"hipdnnGetRNNDataDescriptor", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnSetRNNBiasMode", {"hipdnnSetRNNBiasMode", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnGetRNNBiasMode", {"hipdnnGetRNNBiasMode", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, // cuDNN Connectionist Temporal Classification loss functions {"cudnnCreateCTCLossDescriptor", {"hipdnnCreateCTCLossDescriptor", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, @@ -250,4 +258,20 @@ const std::map CUDA_DNN_FUNCTION_MAP{ {"cudnnRNNSetClip", {"hipdnnRNNSetClip", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, {"cudnnRNNGetClip", {"hipdnnRNNGetClip", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, + // cuDNN Sequence functions + {"cudnnCreateSeqDataDescriptor", {"hipdnnCreateSeqDataDescriptor", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnDestroySeqDataDescriptor", {"hipdnnDestroySeqDataDescriptor", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnSetSeqDataDescriptor", {"hipdnnSetSeqDataDescriptor", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnGetSeqDataDescriptor", {"hipdnnGetSeqDataDescriptor", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, + + // cuDNN Multihead Attention functions + {"cudnnCreateAttnDescriptor", {"hipdnnCreateAttnDescriptor", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnDestroyAttnDescriptor", {"hipdnnDestroyAttnDescriptor", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnSetAttnDescriptor", {"hipdnnSetAttnDescriptor", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnGetAttnDescriptor", {"hipdnnGetAttnDescriptor", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnGetMultiHeadAttnBuffers", {"hipdnnGetMultiHeadAttnBuffers", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnGetMultiHeadAttnWeights", {"hipdnnGetMultiHeadAttnWeights", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnMultiHeadAttnForward", {"hipdnnMultiHeadAttnForward", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnMultiHeadAttnBackwardData", {"hipdnnMultiHeadAttnBackwardData", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnMultiHeadAttnBackwardWeights", {"hipdnnMultiHeadAttnBackwardWeights", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, }; diff --git a/hipify-clang/src/CUDA2HIP_DNN_API_types.cpp b/hipify-clang/src/CUDA2HIP_DNN_API_types.cpp index 508726e20e..431eaab72c 100644 --- a/hipify-clang/src/CUDA2HIP_DNN_API_types.cpp +++ b/hipify-clang/src/CUDA2HIP_DNN_API_types.cpp @@ -35,6 +35,7 @@ const std::map CUDA_DNN_TYPE_NAME_MAP{ {"CUDNN_SEV_ERROR_EN", {"HIPDNN_SEV_ERROR_EN", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, {"CUDNN_SEV_WARNING_EN", {"HIPDNN_SEV_WARNING_EN", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, {"CUDNN_SEV_INFO_EN", {"HIPDNN_SEV_INFO_EN", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, + {"CUDNN_SEQDATA_DIM_COUNT", {"HIPDNN_SEQDATA_DIM_COUNT", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 4 // cuDNN enums {"cudnnStatus_t", {"hipdnnStatus_t", "", CONV_TYPE, API_DNN}}, @@ -60,6 +61,9 @@ const std::map CUDA_DNN_TYPE_NAME_MAP{ {"CUDNN_TENSOR_NCHW", {"HIPDNN_TENSOR_NCHW", "", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 {"CUDNN_TENSOR_NHWC", {"HIPDNN_TENSOR_NHWC", "", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 {"CUDNN_TENSOR_NCHW_VECT_C", {"HIPDNN_TENSOR_NCHW_VECT_C", "", CONV_NUMERIC_LITERAL, API_DNN}}, // 2 + {"cudnnFoldingDirection_t", {"hipdnnFoldingDirection_t", "", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + {"CUDNN_TRANSFORM_FOLD", {"HIPDNN_TRANSFORM_FOLD", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 0U + {"CUDNN_TRANSFORM_UNFOLD", {"HIPDNN_TRANSFORM_UNFOLD", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1U {"cudnnDataType_t", {"hipdnnDataType_t", "", CONV_TYPE, API_DNN}}, {"CUDNN_DATA_FLOAT", {"HIPDNN_DATA_FLOAT", "", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 {"CUDNN_DATA_DOUBLE", {"HIPDNN_DATA_DOUBLE", "", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 @@ -147,6 +151,11 @@ const std::map CUDA_DNN_TYPE_NAME_MAP{ {"CUDNN_RNN_TANH", {"HIPDNN_RNN_TANH", "", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 {"CUDNN_LSTM", {"HIPDNN_LSTM", "", CONV_NUMERIC_LITERAL, API_DNN}}, // 2 {"CUDNN_GRU", {"HIPDNN_GRU", "", CONV_NUMERIC_LITERAL, API_DNN}}, // 3 + {"cudnnRNNBiasMode_t", {"hipdnnRNNBiasMode_t", "", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + {"CUDNN_RNN_NO_BIAS", {"HIPDNN_RNN_NO_BIAS", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 0 + {"CUDNN_RNN_SINGLE_INP_BIAS", {"HIPDNN_RNN_SINGLE_INP_BIAS", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1 + {"CUDNN_RNN_DOUBLE_BIAS", {"HIPDNN_RNN_DOUBLE_BIAS", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 2 + {"CUDNN_RNN_SINGLE_REC_BIAS", {"HIPDNN_RNN_SINGLE_REC_BIAS", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 3 {"cudnnOpTensorOp_t", {"hipdnnOpTensorOp_t", "", CONV_TYPE, API_DNN}}, {"CUDNN_OP_TENSOR_ADD", {"HIPDNN_OP_TENSOR_ADD", "", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 {"CUDNN_OP_TENSOR_MUL", {"HIPDNN_OP_TENSOR_MUL", "", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 @@ -215,6 +224,22 @@ const std::map CUDA_DNN_TYPE_NAME_MAP{ {"cudnnRNNPaddingMode_t", {"hipdnnRNNPaddingMode_t", "", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, {"CUDNN_RNN_PADDED_IO_DISABLED", {"HIPDNN_RNN_PADDED_IO_DISABLED", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 0 {"CUDNN_RNN_PADDED_IO_ENABLED", {"HIPDNN_RNN_PADDED_IO_ENABLED", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1 + {"cudnnSeqDataAxis_t", {"hipdnnSeqDataAxis_t", "", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + {"CUDNN_SEQDATA_TIME_DIM", {"HIPDNN_SEQDATA_TIME_DIM", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 0 + {"CUDNN_SEQDATA_BATCH_DIM", {"HIPDNN_SEQDATA_BATCH_DIM", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1 + {"CUDNN_SEQDATA_BEAM_DIM", {"HIPDNN_SEQDATA_BEAM_DIM", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 2 + {"CUDNN_SEQDATA_VECT_DIM", {"HIPDNN_SEQDATA_VECT_DIM", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 3 + {"cudnnAttnQueryMap_t", {"hipdnnAttnQueryMap_t", "", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + {"CUDNN_ATTN_QUERYMAP_ALL_TO_ONE", {"HIPDNN_ATTN_QUERYMAP_ALL_TO_ONE", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 0 + {"CUDNN_ATTN_QUERYMAP_ONE_TO_ONE", {"HIPDNN_ATTN_QUERYMAP_ONE_TO_ONE", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1 + {"cudnnMultiHeadAttnWeightKind_t", {"hipdnnMultiHeadAttnWeightKind_t", "", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + {"CUDNN_MH_ATTN_Q_WEIGHTS", {"HIPDNN_MH_ATTN_Q_WEIGHTS", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 0 + {"CUDNN_MH_ATTN_K_WEIGHTS", {"HIPDNN_MH_ATTN_K_WEIGHTS", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1 + {"CUDNN_MH_ATTN_V_WEIGHTS", {"HIPDNN_MH_ATTN_V_WEIGHTS", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 2 + {"CUDNN_MH_ATTN_O_WEIGHTS", {"HIPDNN_MH_ATTN_O_WEIGHTS", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 3 + {"cudnnWgradMode_t", {"hipdnnWgradMode_t", "", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + {"CUDNN_WGRAD_MODE_ADD", {"HIPDNN_WGRAD_MODE_ADD", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 0 + {"CUDNN_WGRAD_MODE_SET", {"HIPDNN_WGRAD_MODE_SET", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1 // cuDNN types {"cudnnContext", {"hipdnnContext", "", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, @@ -239,6 +264,8 @@ const std::map CUDA_DNN_TYPE_NAME_MAP{ {"cudnnReduceTensorDescriptor_t", {"hipdnnReduceTensorDescriptor_t", "", CONV_TYPE, API_DNN}}, {"cudnnCTCLossStruct", {"hipdnnCTCLossStruct", "", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, {"cudnnCTCLossDescriptor_t", {"hipdnnCTCLossDescriptor_t", "", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnTensorTransformStruct", {"hipdnnTensorTransformStruct", "", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnTensorTransformDescriptor_t", {"hipdnnTensorTransformDescriptor_t", "", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, {"cudnnConvolutionFwdAlgoPerf_t", {"hipdnnConvolutionFwdAlgoPerf_t", "", CONV_TYPE, API_DNN}}, {"cudnnConvolutionBwdFilterAlgoPerf_t", {"hipdnnConvolutionBwdFilterAlgoPerf_t", "", CONV_TYPE, API_DNN}}, {"cudnnConvolutionBwdDataAlgoPerf_t", {"hipdnnConvolutionBwdDataAlgoPerf_t", "", CONV_TYPE, API_DNN}}, @@ -257,4 +284,9 @@ const std::map CUDA_DNN_TYPE_NAME_MAP{ {"cudnnCallback_t", {"hipdnnCallback_t", "", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, {"cudnnRNNDataStruct", {"hipdnnRNNDataStruct", "", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, {"cudnnRNNDataDescriptor_t", {"hipdnnRNNDataDescriptor_t", "", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnSeqDataStruct", {"hipdnnSeqDataStruct", "", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnSeqDataDescriptor_t", {"hipdnnSeqDataDescriptor_t", "", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnAttnStruct", {"hipdnnAttnStruct", "", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnAttnDescriptor_t", {"hipdnnAttnDescriptor_t", "", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + }; From 8a548bf40b45cc4f3c90642859b9b2f68437379f Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Mon, 15 Apr 2019 23:04:39 -0400 Subject: [PATCH 05/11] Use COMgr to read Kernel Args Metadata (#1006) * Add CMAKE dep to amd_comgr * Use COMGR for read_kernarg_metadata in COV2 * Do not assume kernargs exist * Add proper metadata destroy cleanup * Use a process function for easier destroy * Remove old read_kernarg_metadata * Clean up HCC, prints, names * Use COMGR in CMAKE by default * Move metadata lookup for keyword values into helper * Remove C string usage for lookup_keyword_value * Guard COMGR for non-NVCC path * Add hip_hcc dependency on comgr package * Add lifetime to metadata nodes * Find COMGR config file for amd_comgr target * Move set_active data earlier --- CMakeLists.txt | 15 ++ bin/hipcc | 2 + include/hip/hcc_detail/program_state.hpp | 244 ++++++++++++++++------- packaging/hip_hcc.txt | 8 +- 4 files changed, 188 insertions(+), 81 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 78d2bc7a30..a71d847af7 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -193,6 +193,18 @@ if(USE_PROF_API EQUAL 1) endif() endif() +################ +# Detect Code Object Manager +################ + +if(HIP_PLATFORM STREQUAL "hcc") + find_package(amd_comgr REQUIRED CONFIG + PATHS + /opt/rocm/lib/cmake/amd_comgr + ) + MESSAGE(STATUS "Code Object Manager found at ${amd_comgr_DIR}.") +endif() + ############################# # Build steps ############################# @@ -264,6 +276,9 @@ if(HIP_PLATFORM STREQUAL "hcc") target_link_libraries(hip_hcc_static PRIVATE hc_am) endif() + target_link_libraries(hip_hcc PRIVATE amd_comgr) + target_link_libraries(hip_hcc_static PRIVATE amd_comgr) + string(REPLACE " " ";" HCC_CXX_FLAGS_LIST ${HCC_CXX_FLAGS}) foreach(TARGET hip_hcc hip_hcc_static) target_include_directories(${TARGET} SYSTEM INTERFACE $/include>;${HSA_PATH}/include) diff --git a/bin/hipcc b/bin/hipcc index 6f020662de..2648e45ef9 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -265,9 +265,11 @@ if ($HIP_PLATFORM eq "clang") { $HIPCXXFLAGS .= " -isystem $HIP_PATH/include/hip/hcc_detail/cuda"; $HIPCXXFLAGS .= " -isystem $HSA_PATH/include"; + $HIPCXXFLAGS .= " -isystem /opt/rocm/include"; $HIPCXXFLAGS .= " -Wno-deprecated-register"; $HIPLDFLAGS .= " -L$HSA_PATH/lib -L$ROCM_PATH/lib -lhsa-runtime64 -lhc_am "; + $HIPLDFLAGS .= " -L/opt/rocm/lib -Wl,--rpath=/opt/rocm/lib -lamd_comgr "; # $HIPLDFLAGS .= " -L$HCC_HOME/compiler/lib -lLLVMAMDGPUDesc -lLLVMAMDGPUUtils -lLLVMMC -lLLVMCore -lLLVMSupport "; # Add trace marker library: diff --git a/include/hip/hcc_detail/program_state.hpp b/include/hip/hcc_detail/program_state.hpp index f05f41d3f5..76ccf33b87 100644 --- a/include/hip/hcc_detail/program_state.hpp +++ b/include/hip/hcc_detail/program_state.hpp @@ -39,6 +39,7 @@ THE SOFTWARE. #include #include #include +#include #include @@ -55,6 +56,8 @@ THE SOFTWARE. #include #include #include +#include +#include struct ihipModuleSymbol_t; using hipFunction_t = ihipModuleSymbol_t*; @@ -558,92 +561,185 @@ const std::unordered_map< } inline -std::size_t parse_args( - const std::string& metadata, - std::size_t f, - std::size_t l, - std::vector>& size_align) { - if (f == l) return f; - if (!size_align.empty()) return l; +void checkError( + amd_comgr_status_t status, + char const *str) { + if (status != AMD_COMGR_STATUS_SUCCESS) { + const char *status_str; + status = amd_comgr_status_string(status, &status_str); + if (status == AMD_COMGR_STATUS_SUCCESS) + std::cerr << "FAILED: " << str << "\n REASON: " << status_str << std::endl; + hip_throw(std::runtime_error{"Metadata parsing failed."}); + } +} - do { - static constexpr size_t size_sz{5}; - f = metadata.find("Size:", f) + size_sz; +class comgr_metadata_node { + public: + amd_comgr_metadata_node_t node; + bool active; + comgr_metadata_node() : active(false) {} + ~comgr_metadata_node() { + if(active) + checkError(amd_comgr_destroy_metadata(node), "amd_comgr_destroy_metadata"); + } + bool is_active() { return active; } + void set_active(bool value) { active = value; } + comgr_metadata_node(const comgr_metadata_node&) = delete; + comgr_metadata_node(comgr_metadata_node&&) = delete; + comgr_metadata_node& operator=(const comgr_metadata_node&) = delete; + comgr_metadata_node& operator=(comgr_metadata_node&&) = delete; - if (l <= f) return f; +}; - auto size = std::strtoul(&metadata[f], nullptr, 10); +class comgr_data { + public: + amd_comgr_data_t data; + bool active; + comgr_data() : active(false) {} + ~comgr_data() { + if(active) + checkError(amd_comgr_release_data(data), "amd_comgr_release_data"); + } + bool is_active() { return active; } + void set_active(bool value) { active = value; } + comgr_data(const comgr_data&) = delete; + comgr_data(comgr_data&&) = delete; + comgr_data& operator=(const comgr_data&) = delete; + comgr_data& operator=(comgr_data&&) = delete; - static constexpr size_t align_sz{6}; - f = metadata.find("Align:", f) + align_sz; +}; - char* l{}; - auto align = std::strtoul(&metadata[f], &l, 10); +inline +std::string lookup_keyword_value( + amd_comgr_metadata_node_t& in_node, + std::string keyword) { + amd_comgr_status_t status; + size_t value_size; + comgr_metadata_node value_meta; - f += (l - &metadata[f]) + 1; + status = amd_comgr_metadata_lookup(in_node, keyword.c_str(), &value_meta.node); + checkError(status, "amd_comgr_metadata_lookup"); + value_meta.set_active(true); + status = amd_comgr_get_metadata_string(value_meta.node, &value_size, NULL); + checkError(status, "amd_comgr_get_metadata_string"); + // Since value_size returns size with null terminator, we don't include for C++ string size + value_size--; + std::string value(value_size, '\0'); + status = amd_comgr_get_metadata_string(value_meta.node, &value_size, &value[0]); + checkError(status, "amd_comgr_get_metadata_string"); - size_align.emplace_back(size, align); - } while (true); + return value; } inline -void read_kernarg_metadata( - ELFIO::elfio& reader, +void process_kernarg_metadata( + amd_comgr_metadata_node_t blob_meta, std::unordered_map< std::string, std::vector>>& kernargs) { - // TODO: this is inefficient. - auto it = find_section_if(reader, [](const ELFIO::section* x) { - return x->get_type() == SHT_NOTE; - }); + amd_comgr_status_t status; + amd_comgr_metadata_kind_t mkindLookup; + comgr_metadata_node kernelList; + std::string kernel_name; + size_t num_kernels = 0; + size_t num_kern_args = 0; - if (!it) return; + // Kernels is a list of MAPS!! + status = amd_comgr_metadata_lookup(blob_meta, "Kernels", &kernelList.node); + // FIXME - few hip memset kernels are missing Kernels node + if(status != AMD_COMGR_STATUS_SUCCESS) + return; + kernelList.set_active(true); - const ELFIO::note_section_accessor acc{reader, it}; - for (decltype(acc.get_notes_num()) i = 0; i != acc.get_notes_num(); ++i) { - ELFIO::Elf_Word type{}; - std::string name{}; - void* desc{}; - ELFIO::Elf_Word desc_size{}; - - acc.get_note(i, type, name, desc, desc_size); - - if (name != "AMD") continue; // TODO: switch to using NT_AMD_AMDGPU_HSA_METADATA. - - std::string tmp{ - static_cast(desc), static_cast(desc) + desc_size}; - - auto dx = tmp.find("Kernels:"); - - if (dx == std::string::npos) continue; - - static constexpr decltype(tmp.size()) kernels_sz{8}; - dx += kernels_sz; - - do { - dx = tmp.find("Name:", dx); - - if (dx == std::string::npos) break; - - static constexpr decltype(tmp.size()) name_sz{5}; - dx = tmp.find_first_not_of(" '", dx + name_sz); - - auto fn = tmp.substr(dx, tmp.find_first_of("'\n", dx) - dx); - dx += fn.size(); - - auto dx1 = tmp.find("CodeProps", dx); - dx = tmp.find("Args:", dx); - - if (dx1 < dx) { - dx = dx1; - continue; - } - if (dx == std::string::npos) break; - - static constexpr decltype(tmp.size()) args_sz{5}; - dx = parse_args(tmp, dx + args_sz, dx1, kernargs[fn]); - } while (true); + status = amd_comgr_get_metadata_kind(kernelList.node, &mkindLookup); + if (mkindLookup != AMD_COMGR_METADATA_KIND_LIST) { + hip_throw(std::runtime_error{"Lookup of Kernels didn't return a list\n"}); } + + status = amd_comgr_get_metadata_list_size(kernelList.node, &num_kernels); + checkError(status, "amd_comgr_get_metadata_list_size"); + for (int i = 0; i < num_kernels; i++) { + comgr_metadata_node kernelMap; + status = amd_comgr_index_list_metadata(kernelList.node, i, &kernelMap.node); + checkError(status, "amd_comgr_index_list_metadata"); + kernelMap.set_active(true); + + kernel_name = std::move(lookup_keyword_value(kernelMap.node, "Name")); + + // Check if this kernel was already processed + if(!kernargs[kernel_name].empty()) { + continue; + } + + comgr_metadata_node kernArgList; + status = amd_comgr_metadata_lookup(kernelMap.node, "Args", &kernArgList.node); + if (status == AMD_COMGR_STATUS_SUCCESS ) { + kernArgList.set_active(true); + status = amd_comgr_get_metadata_kind(kernArgList.node, &mkindLookup); + if (mkindLookup != AMD_COMGR_METADATA_KIND_LIST) { + hip_throw(std::runtime_error{"Lookup of Args didn't return a list\n"}); + } + + if (status == AMD_COMGR_STATUS_SUCCESS ) { + status = amd_comgr_get_metadata_list_size(kernArgList.node, &num_kern_args); + checkError(status, "amd_comgr_get_metadata_list_size"); + for (int k_ar = 0; k_ar < num_kern_args; k_ar++) { + comgr_metadata_node kernArgMap; + status = amd_comgr_index_list_metadata(kernArgList.node, k_ar, &kernArgMap.node); + checkError(status, "amd_comgr_index_list_metadata"); + kernArgMap.set_active(true); + size_t k_arg_size, k_arg_align; + + k_arg_size = std::stoul(lookup_keyword_value(kernArgMap.node, "Size")); + k_arg_align = std::stoul(lookup_keyword_value(kernArgMap.node, "Align")); + + // Save it into our kernargs + kernargs[kernel_name].emplace_back(k_arg_size, k_arg_align); + + } // end kernArgMap + } + } // end kernArgList + } // end kernelMap +} // end kernelList + +inline +void read_kernarg_metadata( + std::string blob, + std::unordered_map< + std::string, + std::vector>>& kernargs) { + + const char *blob_buf = blob.data(); + long blob_size = blob.size(); + + amd_comgr_status_t status; + comgr_data blob_data; + status = amd_comgr_create_data(AMD_COMGR_DATA_KIND_RELOCATABLE, &blob_data.data); + checkError(status, "amd_comgr_create_data"); + blob_data.set_active(true); + + status = amd_comgr_set_data(blob_data.data, blob_size, blob_buf); + if(status != AMD_COMGR_STATUS_SUCCESS) + return; + + // We have a valid code object now + status = amd_comgr_set_data_name(blob_data.data, "HIP Code Object"); + checkError(status, "amd_comgr_set_data_name"); + + comgr_metadata_node blob_meta; + status = amd_comgr_get_data_metadata(blob_data.data, &blob_meta.node); + checkError(status, "amd_comgr_get_data_metadata"); + blob_meta.set_active(true); + + // Root is a map + amd_comgr_metadata_kind_t blob_mkind; + status = amd_comgr_get_metadata_kind(blob_meta.node, &blob_mkind); + checkError(status, "amd_comgr_get_metadata_kind"); + if (blob_mkind != AMD_COMGR_METADATA_KIND_MAP) { + hip_throw(std::runtime_error{"Root is not map\n"}); + } + + process_kernarg_metadata(blob_meta.node, kernargs); } inline @@ -657,13 +753,7 @@ const std::unordered_map< std::call_once(f, []() { for (auto&& isa_blobs : code_object_blobs()) { for (auto&& blob : isa_blobs.second) { - std::stringstream tmp{std::string{blob.cbegin(), blob.cend()}}; - - ELFIO::elfio reader; - - if (!reader.load(tmp)) continue; - - read_kernarg_metadata(reader, r); + read_kernarg_metadata(std::string{blob.cbegin(), blob.cend()}, r); } } }); diff --git a/packaging/hip_hcc.txt b/packaging/hip_hcc.txt index fe866e47f9..c2da2fdfb8 100644 --- a/packaging/hip_hcc.txt +++ b/packaging/hip_hcc.txt @@ -30,9 +30,9 @@ set(CPACK_GENERATOR "TGZ;DEB;RPM") set(CPACK_BINARY_DEB "ON") set(CPACK_DEBIAN_PACKAGE_CONTROL_EXTRA "${PROJECT_BINARY_DIR}/postinst;${PROJECT_BINARY_DIR}/prerm") if(@COMPILE_HIP_ATP_MARKER@) - set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip_base (= ${CPACK_PACKAGE_VERSION}), ${HCC_PACKAGE_NAME} (= @HCC_PACKAGE_VERSION@), rocm-profiler") + set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip_base (= ${CPACK_PACKAGE_VERSION}), ${HCC_PACKAGE_NAME} (= @HCC_PACKAGE_VERSION@), rocm-profiler, comgr (>= 1.1)") else() - set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip_base (= ${CPACK_PACKAGE_VERSION}), ${HCC_PACKAGE_NAME} (= @HCC_PACKAGE_VERSION@)") + set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip_base (= ${CPACK_PACKAGE_VERSION}), ${HCC_PACKAGE_NAME} (= @HCC_PACKAGE_VERSION@), comgr (>= 1.1)") endif() set(CPACK_BINARY_RPM "ON") set(CPACK_RPM_PACKAGE_ARCHITECTURE "${CMAKE_SYSTEM_PROCESSOR}") @@ -40,9 +40,9 @@ set(CPACK_RPM_POST_INSTALL_SCRIPT_FILE "${PROJECT_BINARY_DIR}/postinst") set(CPACK_RPM_PRE_UNINSTALL_SCRIPT_FILE "${PROJECT_BINARY_DIR}/prerm") set(CPACK_RPM_PACKAGE_AUTOREQPROV " no") if(@COMPILE_HIP_ATP_MARKER@) - set(CPACK_RPM_PACKAGE_REQUIRES "hip_base = ${CPACK_PACKAGE_VERSION}, ${HCC_PACKAGE_NAME} = @HCC_PACKAGE_VERSION@, rocm-profiler") + set(CPACK_RPM_PACKAGE_REQUIRES "hip_base = ${CPACK_PACKAGE_VERSION}, ${HCC_PACKAGE_NAME} = @HCC_PACKAGE_VERSION@, rocm-profiler, comgr") else() - set(CPACK_RPM_PACKAGE_REQUIRES "hip_base = ${CPACK_PACKAGE_VERSION}, ${HCC_PACKAGE_NAME} = @HCC_PACKAGE_VERSION@") + set(CPACK_RPM_PACKAGE_REQUIRES "hip_base = ${CPACK_PACKAGE_VERSION}, ${HCC_PACKAGE_NAME} = @HCC_PACKAGE_VERSION@, comgr") endif() set(CPACK_RPM_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "/opt") set(CPACK_SOURCE_GENERATOR "TGZ") From dc07b1b06c618b182accec5093e60b0720534ffc Mon Sep 17 00:00:00 2001 From: Mr-LiuSw <33505548+Mr-LiuSw@users.noreply.github.com> Date: Tue, 16 Apr 2019 11:05:36 +0800 Subject: [PATCH 06/11] add little changes in hip_runtime_api.h to work with c language (#1017) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * Update hip_runtime_api.h when i try to use mpicc or gcc to compile a c language code which call some hip runtime api , error occured as > /path/to/hcc_detail/hip_runtime_api.h:2268:33: error: unknown type name ‘hipFuncAttributes’; > hipFuncGetAttributes(hipFuncAttributes* attr, const void* func); add ' struct ' for the first parameter of hipFuncGetAttributes will get ride of this problem. --- include/hip/hcc_detail/hip_runtime_api.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index cee4e81054..ccfe7aefc3 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -2549,7 +2549,7 @@ hipError_t hipModuleGetFunction(hipFunction_t* function, hipModule_t module, con * @returns hipSuccess, hipErrorInvalidDeviceFunction */ -hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func); +hipError_t hipFuncGetAttributes(struct hipFuncAttributes* attr, const void* func); struct Agent_global { From 2b3037a6ea18a143332fd91af4cdfefb35b9a437 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Mon, 15 Apr 2019 20:05:55 -0700 Subject: [PATCH 07/11] In hipFree, synchronize owner of memory (#1018) * In hipFree, if memory is associated with a device, synchronize that device's streams. This changes the behavior from synchronizing the currently set TLS device. * All devices sync in hipFree for _appId=-1 case. * Revert "All devices sync in hipFree for _appId=-1 case." This reverts commit 1efb34d6a8426661e45bc5f763422a1147aeac10. * add HIP_SYNC_FREE env var --- src/hip_hcc.cpp | 3 +++ src/hip_hcc_internal.h | 2 ++ src/hip_memory.cpp | 27 +++++++++++++++++++++++---- 3 files changed, 28 insertions(+), 4 deletions(-) diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index 2955db5e37..ac285e6a83 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -92,6 +92,7 @@ int HIP_EVENT_SYS_RELEASE = 0; int HIP_HOST_COHERENT = 1; int HIP_SYNC_HOST_ALLOC = 1; +int HIP_SYNC_FREE = 0; int HIP_INIT_ALLOC = -1; @@ -1279,6 +1280,8 @@ void HipReadEnv() { READ_ENV_I(release, HIP_SYNC_STREAM_WAIT, 0, "hipStreamWaitEvent will synchronize to host"); + READ_ENV_I(release, HIP_SYNC_FREE, 0, + "Force all calls to hipFree to sync all devices and all streams"); READ_ENV_I(release, HIP_HOST_COHERENT, 0, "If set, all host memory will be allocated as fine-grained system memory. This " diff --git a/src/hip_hcc_internal.h b/src/hip_hcc_internal.h index 0e0db0c628..b40fac93a5 100644 --- a/src/hip_hcc_internal.h +++ b/src/hip_hcc_internal.h @@ -83,6 +83,8 @@ extern int HIP_SYNC_NULL_STREAM; extern int HIP_INIT_ALLOC; extern int HIP_FORCE_NULL_STREAM; +extern int HIP_SYNC_FREE; + extern int HIP_DUMP_CODE_OBJECT; // TODO - remove when this is standard behavior. diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp index 8196be91ff..b6a1164a4a 100644 --- a/src/hip_memory.cpp +++ b/src/hip_memory.cpp @@ -1902,10 +1902,6 @@ hipError_t hipFree(void* ptr) { hipError_t hipStatus = hipErrorInvalidDevicePointer; - // Synchronize to ensure all work has finished. - ihipGetTlsDefaultCtx()->locked_waitAllStreams(); // ignores non-blocking streams, this waits - // for all activity to finish. - if (ptr) { hc::accelerator acc; #if (__hcc_workweek__ >= 17332) @@ -1916,6 +1912,29 @@ hipError_t hipFree(void* ptr) { am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, ptr); if (status == AM_SUCCESS) { if (amPointerInfo._hostPointer == NULL) { + if (HIP_SYNC_FREE) { + // Synchronize all devices, all streams + // to ensure all work has finished on all devices. + // This is disabled by default. + for (unsigned i = 0; i < g_deviceCnt; i++) { + ihipGetPrimaryCtx(i)->locked_waitAllStreams(); + } + } + else { + ihipCtx_t* ctx; + if (amPointerInfo._appId != -1) { +#if USE_APP_PTR_FOR_CTX + ctx = static_cast(amPointerInfo._appPtr); +#else + ctx = ihipGetPrimaryCtx(amPointerInfo._appId); +#endif + } else { + ctx = ihipGetTlsDefaultCtx(); + } + // Synchronize to ensure all work has finished on device owning the memory. + ctx->locked_waitAllStreams(); // ignores non-blocking streams, this waits + // for all activity to finish. + } hc::am_free(ptr); hipStatus = hipSuccess; } From 1b36987c5f683fc49d71b0758c09a45a8b54ffa0 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 16 Apr 2019 11:01:01 +0300 Subject: [PATCH 08/11] [HIPIFY][cuDNN] Add partial cudnnRNNBiasMode_t support --- docs/markdown/CUDNN_API_supported_by_HIP.md | 14 +++++++------- hipify-clang/src/CUDA2HIP_DNN_API_types.cpp | 10 +++++----- 2 files changed, 12 insertions(+), 12 deletions(-) diff --git a/docs/markdown/CUDNN_API_supported_by_HIP.md b/docs/markdown/CUDNN_API_supported_by_HIP.md index be0fd442fc..ba1eab75a9 100644 --- a/docs/markdown/CUDNN_API_supported_by_HIP.md +++ b/docs/markdown/CUDNN_API_supported_by_HIP.md @@ -2,8 +2,8 @@ ## **1. CUDNN Data types** -| **type** | **CUDA** | **HIP** | -|-------------:|---------------------------------------------------------------|------------------------------------------------------------| +| **type** | **CUDA** | **HIP** |**HIP value** (if differs) | +|-------------:|---------------------------------------------------------------|------------------------------------------------------------|---------------------------| | define |`CUDNN_VERSION` |`HIPDNN_VERSION` | | struct |`cudnnContext` | | | struct* |`cudnnHandle_t` |`hipdnnHandle_t` | @@ -174,11 +174,11 @@ | 1 |*`CUDNN_RNN_TANH`* |*`HIPDNN_RNN_TANH`* | | 2 |*`CUDNN_LSTM`* |*`HIPDNN_LSTM`* | | 3 |*`CUDNN_GRU`* |*`HIPDNN_GRU`* | -| enum |***`cudnnRNNBiasMode_t`*** | | -| 0 |*`CUDNN_RNN_NO_BIAS`* | | -| 1 |*`CUDNN_RNN_SINGLE_INP_BIAS`* | | -| 2 |*`CUDNN_RNN_DOUBLE_BIAS`* | | -| 3 |*`CUDNN_RNN_SINGLE_REC_BIAS`* | | +| enum |***`cudnnRNNBiasMode_t`*** |***`hipdnnRNNBiasMode_t`*** | +| 0 |*`CUDNN_RNN_NO_BIAS`* |*`HIPDNN_RNN_NO_BIAS`* | +| 1 |*`CUDNN_RNN_SINGLE_INP_BIAS`* |*`HIPDNN_RNN_WITH_BIAS`* | +| 2 |*`CUDNN_RNN_DOUBLE_BIAS`* |*`HIPDNN_RNN_WITH_BIAS`* | 1 | +| 3 |*`CUDNN_RNN_SINGLE_REC_BIAS`* |*`HIPDNN_RNN_WITH_BIAS`* | 1 | | enum |***`cudnnDirectionMode_t`*** |***`hipdnnDirectionMode_t`*** | | 0 |*`CUDNN_UNIDIRECTIONAL`* |*`HIPDNN_UNIDIRECTIONAL`* | | 1 |*`CUDNN_BIDIRECTIONAL`* |*`HIPDNN_BIDIRECTIONAL`* | diff --git a/hipify-clang/src/CUDA2HIP_DNN_API_types.cpp b/hipify-clang/src/CUDA2HIP_DNN_API_types.cpp index 431eaab72c..349e243082 100644 --- a/hipify-clang/src/CUDA2HIP_DNN_API_types.cpp +++ b/hipify-clang/src/CUDA2HIP_DNN_API_types.cpp @@ -151,11 +151,11 @@ const std::map CUDA_DNN_TYPE_NAME_MAP{ {"CUDNN_RNN_TANH", {"HIPDNN_RNN_TANH", "", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 {"CUDNN_LSTM", {"HIPDNN_LSTM", "", CONV_NUMERIC_LITERAL, API_DNN}}, // 2 {"CUDNN_GRU", {"HIPDNN_GRU", "", CONV_NUMERIC_LITERAL, API_DNN}}, // 3 - {"cudnnRNNBiasMode_t", {"hipdnnRNNBiasMode_t", "", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, - {"CUDNN_RNN_NO_BIAS", {"HIPDNN_RNN_NO_BIAS", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 0 - {"CUDNN_RNN_SINGLE_INP_BIAS", {"HIPDNN_RNN_SINGLE_INP_BIAS", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1 - {"CUDNN_RNN_DOUBLE_BIAS", {"HIPDNN_RNN_DOUBLE_BIAS", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 2 - {"CUDNN_RNN_SINGLE_REC_BIAS", {"HIPDNN_RNN_SINGLE_REC_BIAS", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 3 + {"cudnnRNNBiasMode_t", {"hipdnnRNNBiasMode_t", "", CONV_TYPE, API_DNN}}, + {"CUDNN_RNN_NO_BIAS", {"HIPDNN_RNN_NO_BIAS", "", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 + {"CUDNN_RNN_SINGLE_INP_BIAS", {"HIPDNN_RNN_WITH_BIAS", "", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 + {"CUDNN_RNN_DOUBLE_BIAS", {"HIPDNN_RNN_WITH_BIAS", "", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 + {"CUDNN_RNN_SINGLE_REC_BIAS", {"HIPDNN_RNN_WITH_BIAS", "", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 {"cudnnOpTensorOp_t", {"hipdnnOpTensorOp_t", "", CONV_TYPE, API_DNN}}, {"CUDNN_OP_TENSOR_ADD", {"HIPDNN_OP_TENSOR_ADD", "", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 {"CUDNN_OP_TENSOR_MUL", {"HIPDNN_OP_TENSOR_MUL", "", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 From 032c3bf5b8ca28e0f444a23bba533cc8f2999bef Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 16 Apr 2019 12:52:58 +0300 Subject: [PATCH 09/11] [HIPIFY][BLAS] cuBLAS 10.1 support --- docs/markdown/CUBLAS_API_supported_by_HIP.md | 12 ++++++++++++ hipify-clang/src/CUDA2HIP_BLAS_API_functions.cpp | 10 ++++++++++ hipify-clang/src/CUDA2HIP_BLAS_API_types.cpp | 3 +++ 3 files changed, 25 insertions(+) diff --git a/docs/markdown/CUBLAS_API_supported_by_HIP.md b/docs/markdown/CUBLAS_API_supported_by_HIP.md index 388b86d8ef..7f3cbc6702 100644 --- a/docs/markdown/CUBLAS_API_supported_by_HIP.md +++ b/docs/markdown/CUBLAS_API_supported_by_HIP.md @@ -20,9 +20,12 @@ | 0 |*`CUBLAS_OP_N`* |*`HIPBLAS_OP_N`* | 111 | | 1 |*`CUBLAS_OP_T`* |*`HIPBLAS_OP_T`* | 112 | | 2 |*`CUBLAS_OP_C`* |*`HIPBLAS_OP_C`* | 113 | +| 2 |*`CUBLAS_OP_HERMITAN`* |*`HIPBLAS_OP_C`* | 113 | +| 3 |*`CUBLAS_OP_CONJG`* | | | enum |***`cublasFillMode_t`*** |***`hipblasFillMode_t`*** | | 0 |*`CUBLAS_FILL_MODE_LOWER`* |*`HIPBLAS_FILL_MODE_LOWER`* | 121 | | 1 |*`CUBLAS_FILL_MODE_UPPER`* |*`HIPBLAS_FILL_MODE_UPPER`* | 122 | +| 1 |*`CUBLAS_FILL_MODE_FULL`* |*`HIPBLAS_FILL_MODE_FULL`* | 123 | | enum |***`cublasDiagType_t`*** |***`hipblasDiagType_t`*** | | 0 |*`CUBLAS_DIAG_NON_UNIT`* |*`HIPBLAS_DIAG_NON_UNIT`* | 131 | | 1 |*`CUBLAS_DIAG_UNIT`* |*`HIPBLAS_DIAG_UNIT`* | 132 | @@ -98,6 +101,7 @@ |`cublasGetVersion` | | |`cublasGetVersion_v2` | | |`cublasGetProperty` | | +|`cublasGetCudartVersion` | | |`cublasGetStream` |`hipblasGetStream` | |`cublasGetStream_v2` |`hipblasGetStream` | |`cublasSetStream` |`hipblasSetStream` | @@ -173,6 +177,7 @@ |`cublasDcopy` |`hipblasDcopy` | |`cublasDcopy_v2` |`hipblasDcopy` | |`cublasCcopy` | | +|`cublasCopyEx` | | |`cublasCcopy_v2` | | |`cublasZcopy` | | |`cublasZcopy_v2` | | @@ -184,6 +189,7 @@ |`cublasCswap_v2` | | |`cublasZswap` | | |`cublasZswap_v2` | | +|`cublasIamaxEx` | | |`cublasIsamax` |`hipblasIsamax` | |`cublasIsamax_v2` |`hipblasIsamax` | |`cublasIdamax` |`hipblasIdamax` | @@ -192,6 +198,7 @@ |`cublasIcamax_v2` | | |`cublasIzamax` | | |`cublasIzamax_v2` | | +|`cublasIaminEx` | | |`cublasIsamin` | | |`cublasIsamin_v2` | | |`cublasIdamin` | | @@ -200,6 +207,7 @@ |`cublasIcamin_v2` | | |`cublasIzamin` | | |`cublasIzamin_v2` | | +|`cublasAsumEx` | | |`cublasSasum` |`hipblasSasum` | |`cublasSasum_v2` |`hipblasSasum` | |`cublasDasum` |`hipblasDasum` | @@ -208,6 +216,7 @@ |`cublasScasum_v2` | | |`cublasDzasum` | | |`cublasDzasum_v2` | | +|`cublasRotEx` | | |`cublasSrot` | | |`cublasSrot_v2` | | |`cublasDrot` | | @@ -216,6 +225,7 @@ |`cublasCrot_v2` | | |`cublasZrot` | | |`cublasZrot_v2` | | +|`cublasRotgEx` | | |`cublasZdrot` | | |`cublasZdrot_v2` | | |`cublasSrotg` | | @@ -226,10 +236,12 @@ |`cublasCrotg_v2` | | |`cublasZrotg` | | |`cublasZrotg_v2` | | +|`cublasRotmEx` | | |`cublasSrotm` | | |`cublasSrotm_v2` | | |`cublasDrotm` | | |`cublasDrotm_v2` | | +|`cublasRotmgEx` | | |`cublasSrotmg` | | |`cublasSrotmg_v2` | | |`cublasDrotmg` | | diff --git a/hipify-clang/src/CUDA2HIP_BLAS_API_functions.cpp b/hipify-clang/src/CUDA2HIP_BLAS_API_functions.cpp index b6bdf4ebe9..9d7f559e1a 100644 --- a/hipify-clang/src/CUDA2HIP_BLAS_API_functions.cpp +++ b/hipify-clang/src/CUDA2HIP_BLAS_API_functions.cpp @@ -421,6 +421,7 @@ const std::map CUDA_BLAS_FUNCTION_MAP{ {"cublasGetStream_v2", {"hipblasGetStream", "rocblas_get_stream", CONV_LIB_FUNC, API_BLAS}}, {"cublasGetPointerMode_v2", {"hipblasGetPointerMode", "rocblas_set_pointer_mode", CONV_LIB_FUNC, API_BLAS}}, {"cublasSetPointerMode_v2", {"hipblasSetPointerMode", "rocblas_get_pointer_mode", CONV_LIB_FUNC, API_BLAS}}, + {"cublasGetCudartVersion", {"hipblasGetCudartVersion", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, // GEMV {"cublasSgemv_v2", {"hipblasSgemv", "rocblas_sgemv", CONV_LIB_FUNC, API_BLAS}}, @@ -631,36 +632,42 @@ const std::map CUDA_BLAS_FUNCTION_MAP{ {"cublasZaxpy_v2", {"hipblasZaxpy", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, // COPY + {"cublasCopyEx", {"hipblasCopyEx", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, {"cublasScopy_v2", {"hipblasScopy", "rocblas_scopy", CONV_LIB_FUNC, API_BLAS}}, {"cublasDcopy_v2", {"hipblasDcopy", "rocblas_dcopy", CONV_LIB_FUNC, API_BLAS}}, {"cublasCcopy_v2", {"hipblasCcopy", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, {"cublasZcopy_v2", {"hipblasZcopy", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, // SWAP + {"cublasSwapEx", {"hipblasSwapEx", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, {"cublasSswap_v2", {"hipblasSswap", "rocblas_sswap", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasDswap_v2", {"hipblasDswap", "rocblas_dswap", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasCswap_v2", {"hipblasCswap", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, {"cublasZswap_v2", {"hipblasZswap", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, // AMAX + {"cublasIamaxEx", {"hipblasIamaxEx", "", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasIsamax_v2", {"hipblasIsamax", "rocblas_isamax", CONV_LIB_FUNC, API_BLAS}}, {"cublasIdamax_v2", {"hipblasIdamax", "rocblas_idamax", CONV_LIB_FUNC, API_BLAS}}, {"cublasIcamax_v2", {"hipblasIcamax", "", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasIzamax_v2", {"hipblasIzamax", "", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}}, // AMIN + {"cublasIaminEx", {"hipblasIaminEx", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, {"cublasIsamin_v2", {"hipblasIsamin", "rocblas_isamin", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasIdamin_v2", {"hipblasIdamin", "rocblas_idamin", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasIcamin_v2", {"hipblasIcamin", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, {"cublasIzamin_v2", {"hipblasIzamin", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, // ASUM + {"cublasAsumEx", {"hipblasAsumEx", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, {"cublasSasum_v2", {"hipblasSasum", "rocblas_sasum", CONV_LIB_FUNC, API_BLAS}}, {"cublasDasum_v2", {"hipblasDasum", "rocblas_dasum", CONV_LIB_FUNC, API_BLAS}}, {"cublasScasum_v2", {"hipblasScasum", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, {"cublasDzasum_v2", {"hipblasDzasum", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, // ROT + {"cublasRotEx", {"hipblasRotEx", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, {"cublasSrot_v2", {"hipblasSrot", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, {"cublasDrot_v2", {"hipblasDrot", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, {"cublasCrot_v2", {"hipblasCrot", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, @@ -669,16 +676,19 @@ const std::map CUDA_BLAS_FUNCTION_MAP{ {"cublasZdrot_v2", {"hipblasZdrot", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, // ROTG + {"cublasRotgEx", {"hipblasRotgEx", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, {"cublasSrotg_v2", {"hipblasSrotg", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, {"cublasDrotg_v2", {"hipblasDrotg", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, {"cublasCrotg_v2", {"hipblasCrotg", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, {"cublasZrotg_v2", {"hipblasZrotg", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, // ROTM + {"cublasRotmEx", {"hipblasRotmEx", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, {"cublasSrotm_v2", {"hipblasSrotm", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, {"cublasDrotm_v2", {"hipblasDrotm", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, // ROTMG + {"cublasRotmgEx", {"hipblasRotmgEx", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, {"cublasSrotmg_v2", {"hipblasSrotmg", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, {"cublasDrotmg_v2", {"hipblasDrotmg", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, }; diff --git a/hipify-clang/src/CUDA2HIP_BLAS_API_types.cpp b/hipify-clang/src/CUDA2HIP_BLAS_API_types.cpp index 6725d6b5a0..2e2168efa0 100644 --- a/hipify-clang/src/CUDA2HIP_BLAS_API_types.cpp +++ b/hipify-clang/src/CUDA2HIP_BLAS_API_types.cpp @@ -29,6 +29,8 @@ const std::map CUDA_BLAS_TYPE_NAME_MAP{ {"CUBLAS_OP_N", {"HIPBLAS_OP_N", "rocblas_operation_none", CONV_NUMERIC_LITERAL, API_BLAS}}, {"CUBLAS_OP_T", {"HIPBLAS_OP_T", "rocblas_operation_transpose", CONV_NUMERIC_LITERAL, API_BLAS}}, {"CUBLAS_OP_C", {"HIPBLAS_OP_C", "rocblas_operation_conjugate_transpose", CONV_NUMERIC_LITERAL, API_BLAS}}, + {"CUBLAS_OP_HERMITAN", {"HIPBLAS_OP_C", "rocblas_operation_conjugate_transpose", CONV_NUMERIC_LITERAL, API_BLAS}}, + {"CUBLAS_OP_CONJG", {"HIPBLAS_OP_CONJG", "rocblas_operation_conjugate", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // Blas statuses {"cublasStatus", {"hipblasStatus_t", "rocblas_status", CONV_TYPE, API_BLAS}}, @@ -48,6 +50,7 @@ const std::map CUDA_BLAS_TYPE_NAME_MAP{ {"cublasFillMode_t", {"hipblasFillMode_t", "rocblas_fill", CONV_TYPE, API_BLAS}}, {"CUBLAS_FILL_MODE_LOWER", {"HIPBLAS_FILL_MODE_LOWER", "rocblas_fill_lower", CONV_NUMERIC_LITERAL, API_BLAS}}, {"CUBLAS_FILL_MODE_UPPER", {"HIPBLAS_FILL_MODE_UPPER", "rocblas_fill_upper", CONV_NUMERIC_LITERAL, API_BLAS}}, + {"CUBLAS_FILL_MODE_FULL", {"HIPBLAS_FILL_MODE_FULL", "rocblas_fill_full", CONV_NUMERIC_LITERAL, API_BLAS}}, // Blas Diag Types {"cublasDiagType_t", {"hipblasDiagType_t", "rocblas_diagonal", CONV_TYPE, API_BLAS}}, From e1c87d8caefa10eb0358a72a8009ff864344b30d Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 16 Apr 2019 14:59:44 +0300 Subject: [PATCH 10/11] [HIPIFY][SPARSE] cuSPARSE 10.1 support --- .../markdown/CUSPARSE_API_supported_by_HIP.md | 54 +++++++++++++++++++ .../src/CUDA2HIP_SPARSE_API_functions.cpp | 29 ++++++++-- .../src/CUDA2HIP_SPARSE_API_types.cpp | 29 ++++++++++ 3 files changed, 109 insertions(+), 3 deletions(-) diff --git a/docs/markdown/CUSPARSE_API_supported_by_HIP.md b/docs/markdown/CUSPARSE_API_supported_by_HIP.md index 4dadea2c93..a90873d6c7 100644 --- a/docs/markdown/CUSPARSE_API_supported_by_HIP.md +++ b/docs/markdown/CUSPARSE_API_supported_by_HIP.md @@ -82,6 +82,27 @@ | typedef |`cusparseColorInfo_t` | | | struct |`pruneInfo` | | | typedef |`pruneInfo_t` | | +| enum |***`cusparseCsr2CscAlg_t`*** | | +| 1 |*`CUSPARSE_CSR2CSC_ALG1`* | | +| 2 |*`CUSPARSE_CSR2CSC_ALG2`* | | +| enum |***`cusparseFormat_t`*** | | +| 1 |*`CUSPARSE_FORMAT_CSR`* | | +| 2 |*`CUSPARSE_FORMAT_CSC`* | | +| 3 |*`CUSPARSE_FORMAT_COO`* | | +| enum |***`cusparseOrder_t`*** | | +| 1 |*`CUSPARSE_ORDER_COL`* | | +| 2 |*`CUSPARSE_ORDER_ROW`* | | +| enum |***`cusparseSpMMAlg_t`*** | | +| 1 |*`CUSPARSE_COOMM_ALG1`* | | +| 2 |*`CUSPARSE_COOMM_ALG2`* | | +| 3 |*`CUSPARSE_COOMM_ALG3`* | | +| enum |***`cusparseIndexType_t`*** | | +| 1 |*`CUSPARSE_INDEX_16U`* | | +| 2 |*`CUSPARSE_INDEX_32I`* | | +| struct |`cusparseSpMatDescr` | | +| typedef |`cusparseSpMatDescr_t` | | +| struct |`cusparseDnMatDescr` | | +| typedef |`cusparseDnMatDescr_t` | | ## **2. cuSPARSE Helper Function Reference** @@ -530,6 +551,8 @@ |`cusparseCcsr2csc` | | |`cusparseZcsr2csc` | | |`cusparseCsr2cscEx` | | +|`cusparseCsr2cscEx2` | | +|`cusparseCsr2cscEx2_bufferSize` | | |`cusparseScsr2dense` | | |`cusparseDcsr2dense` | | |`cusparseCcsr2dense` | | @@ -616,3 +639,34 @@ |`cusparseDnnz_compress` | | |`cusparseCnnz_compress` | | |`cusparseZnnz_compress` | | + +## **10. cuSPARSE Generic API Reference** + +## ***10.1. Generic Sparse API helper functions*** + +| **CUDA** | **HIP** | +|-----------------------------------------------------------|-------------------------------------------------| +|`cusparseCreateCoo` | | +|`cusparseDestroySpMat` | | +|`cusparseCooGet` | | +|`cusparseSpMatGetFormat` | | +|`cusparseSpMatGetIndexBase` | | +|`cusparseSpMatSetNumBatches` | | +|`cusparseSpMatGetNumBatches` | | + +## ***10.2. Generic Dense API helper functions*** + +| **CUDA** | **HIP** | +|-----------------------------------------------------------|-------------------------------------------------| +|`cusparseCreateDnMat` | | +|`cusparseDestroyDnMat` | | +|`cusparseDnMatGet` | | +|`cusparseDnMatSetStridedBatch` | | +|`cusparseDnMatGetStridedBatch` | | + +## ***10.3. Generic SpMM API functions*** + +| **CUDA** | **HIP** | +|-----------------------------------------------------------|-------------------------------------------------| +|`cusparseSpMM` | | +|`cusparseSpMM_bufferSize` | | diff --git a/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp b/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp index 210ab1ee06..f5a969b9be 100644 --- a/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp +++ b/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp @@ -516,6 +516,8 @@ const std::map CUDA_SPARSE_FUNCTION_MAP{ {"cusparseZcsr2csc", {"hipsparseZcsr2csc", "", 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_bufferSize", {"hipsparseCsr2cscEx2_bufferSize", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseScsr2dense", {"hipsparseScsr2dense", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseDcsr2dense", {"hipsparseDcsr2dense", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, @@ -616,9 +618,9 @@ const std::map CUDA_SPARSE_FUNCTION_MAP{ {"cusparseSpruneDense2csrNnzByPercentage", {"hipsparseSpruneDense2csrNnzByPercentage", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseDpruneDense2csrNnzByPercentage", {"hipsparseDpruneDense2csrNnzByPercentage", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseHpruneCsr2csrByPercentage_bufferSizeExt", {"hipsparseHpruneCsr2csrByPercentage_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseSpruneCsr2csrByPercentage_bufferSizeExt", {"hipsparseSpruneCsr2csrByPercentage_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseDpruneCsr2csrByPercentage_bufferSizeExt", {"hipsparseDpruneCsr2csrByPercentage_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseHpruneCsr2csrByPercentage_bufferSizeExt", {"hipsparseHpruneCsr2csrByPercentage_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseSpruneCsr2csrByPercentage_bufferSizeExt", {"hipsparseSpruneCsr2csrByPercentage_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDpruneCsr2csrByPercentage_bufferSizeExt", {"hipsparseDpruneCsr2csrByPercentage_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseHpruneCsr2csrNnzByPercentage", {"hipsparseHpruneCsr2csrNnzByPercentage", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseSpruneCsr2csrNnzByPercentage", {"hipsparseSpruneCsr2csrNnzByPercentage", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, @@ -628,4 +630,25 @@ const std::map CUDA_SPARSE_FUNCTION_MAP{ {"cusparseDnnz_compress", {"hipsparseDnnz_compress", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseCnnz_compress", {"hipsparseCnnz_compress", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZnnz_compress", {"hipsparseZnnz_compress", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + // 13. cuSPARSE Generic API Reference + // Generic Sparse API helper functions + {"cusparseCreateCoo", {"hipsparseCreateCoo", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDestroySpMat", {"hipsparseDestroySpMat", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCooGet", {"hipsparseCooGet", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseSpMatGetFormat", {"hipsparseSpMatGetFormat", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseSpMatGetIndexBase", {"hipsparseSpMatGetIndexBase", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseSpMatSetNumBatches", {"hipsparseSpMatSetNumBatches", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseSpMatGetNumBatches", {"hipsparseSpMatGetNumBatches", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + // Generic Dense API helper functions + {"cusparseCreateDnMat", {"hipsparseCreateDnMat", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDestroyDnMat", {"hipsparseDestroyDnMat", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDnMatGet", {"hipsparseDnMatGet", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDnMatSetStridedBatch", {"hipsparseDnMatSetStridedBatch", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDnMatGetStridedBatch", {"hipsparseDnMatGetStridedBatch", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + + // Generic SpMM API functions + {"cusparseSpMM", {"hipsparseSpMM", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseSpMM_bufferSize", {"hipsparseSpMM_bufferSize", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, }; diff --git a/hipify-clang/src/CUDA2HIP_SPARSE_API_types.cpp b/hipify-clang/src/CUDA2HIP_SPARSE_API_types.cpp index 02e5785c4c..364c98f515 100644 --- a/hipify-clang/src/CUDA2HIP_SPARSE_API_types.cpp +++ b/hipify-clang/src/CUDA2HIP_SPARSE_API_types.cpp @@ -68,6 +68,12 @@ const std::map CUDA_SPARSE_TYPE_NAME_MAP{ {"pruneInfo", {"pruneInfo", "", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}}, {"pruneInfo_t", {"pruneInfo_t", "", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseSpMatDescr", {"hipsparseSpMatDescr", "", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseSpMatDescr_t", {"hipsparseSpMatDescr_t", "", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseDnMatDescr", {"hipsparseDnMatDescr", "", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseDnMatDescr_t", {"hipsparseDnMatDescr_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}}, @@ -130,4 +136,27 @@ const std::map CUDA_SPARSE_TYPE_NAME_MAP{ {"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}}, + + {"cusparseCsr2CscAlg_t", {"hipsparseCsr2CscAlg_t", "", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}}, + {"CUSPARSE_CSR2CSC_ALG1", {"HIPSPARSE_CSR2CSC_ALG1", "", CONV_NUMERIC_LITERAL, API_SPARSE, HIP_UNSUPPORTED}}, + {"CUSPARSE_CSR2CSC_ALG2", {"HIPSPARSE_CSR2CSC_ALG2", "", CONV_NUMERIC_LITERAL, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseFormat_t", {"hipsparseFormat_t", "", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}}, + {"CUSPARSE_FORMAT_CSR", {"HIPSPARSE_FORMAT_CSR", "", CONV_NUMERIC_LITERAL, API_SPARSE, HIP_UNSUPPORTED}}, + {"CUSPARSE_FORMAT_CSC", {"HIPSPARSE_FORMAT_CSC", "", CONV_NUMERIC_LITERAL, API_SPARSE, HIP_UNSUPPORTED}}, + {"CUSPARSE_FORMAT_COO", {"HIPSPARSE_FORMAT_COO", "", CONV_NUMERIC_LITERAL, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseOrder_t", {"hipsparseOrder_t", "", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}}, + {"CUSPARSE_ORDER_COL", {"HIPSPARSE_ORDER_COL", "", CONV_NUMERIC_LITERAL, API_SPARSE, HIP_UNSUPPORTED}}, + {"CUSPARSE_ORDER_ROW", {"HIPSPARSE_ORDER_ROW", "", CONV_NUMERIC_LITERAL, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseSpMMAlg_t", {"hipsparseSpMMAlg_t", "", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}}, + {"CUSPARSE_COOMM_ALG1", {"HIPSPARSE_COOMM_ALG1", "", CONV_NUMERIC_LITERAL, API_SPARSE, HIP_UNSUPPORTED}}, + {"CUSPARSE_COOMM_ALG2", {"HIPSPARSE_COOMM_ALG2", "", CONV_NUMERIC_LITERAL, API_SPARSE, HIP_UNSUPPORTED}}, + {"CUSPARSE_COOMM_ALG3", {"HIPSPARSE_COOMM_ALG3", "", CONV_NUMERIC_LITERAL, API_SPARSE, HIP_UNSUPPORTED}}, + + {"cusparseIndexType_t", {"hipsparseIndexType_t", "", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}}, + {"CUSPARSE_INDEX_16U", {"HIPSPARSE_INDEX_16U", "", CONV_NUMERIC_LITERAL, API_SPARSE, HIP_UNSUPPORTED}}, + {"CUSPARSE_INDEX_32I", {"HIPSPARSE_INDEX_32I", "", CONV_NUMERIC_LITERAL, API_SPARSE, HIP_UNSUPPORTED}}, + }; From 5843530a06e4bcbda66d3cefa1e7b32580f2f4c3 Mon Sep 17 00:00:00 2001 From: David Salinas Date: Tue, 16 Apr 2019 19:33:34 +0000 Subject: [PATCH 11/11] Revert "append the ELF flags for sram-ecc and xnack to the target triple per code object" This reverts commit c61f265657b30e3460f3107bc6cc41ddcdaa61ad. --- include/hip/hcc_detail/code_object_bundle.hpp | 14 -------------- 1 file changed, 14 deletions(-) diff --git a/include/hip/hcc_detail/code_object_bundle.hpp b/include/hip/hcc_detail/code_object_bundle.hpp index 1b24bf55fb..8b1ae8aa49 100644 --- a/include/hip/hcc_detail/code_object_bundle.hpp +++ b/include/hip/hcc_detail/code_object_bundle.hpp @@ -31,8 +31,6 @@ THE SOFTWARE. #include #include #include -#include -#include "elfio/elfio.hpp" namespace hip_impl { @@ -127,18 +125,6 @@ class Bundled_code_header { std::copy_n(f + y.header.offset, y.header.bundle_sz, std::back_inserter(y.blob)); - // we need to read the blob (ELF) for this bundle to get the feature flags - std::stringstream tmp{std::string{y.blob.cbegin(), y.blob.cend()}}; - ELFIO::elfio reader; - if (reader.load(tmp)) { - unsigned int xnack_enabled = reader.get_flags() & 0x100; - unsigned int sram_ecc_enabled = reader.get_flags() & 0x200; - if (xnack_enabled) - y.triple.append("+xnack"); - if (sram_ecc_enabled) - y.triple.append("+sram-ecc"); - } - it += y.header.triple_sz; x.bundled_code_size = std::max(x.bundled_code_size,