From 3da00458ac3a515da1709a307f04cd846408e182 Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Sun, 17 Dec 2023 21:27:08 -0800 Subject: [PATCH 1/6] attempt to fix found issue with hiprtc --- .../include/utility/config.hpp | 13 +++++ .../include/utility/type.hpp | 52 ------------------- 2 files changed, 13 insertions(+), 52 deletions(-) diff --git a/src/composable_kernel/composable_kernel/include/utility/config.hpp b/src/composable_kernel/composable_kernel/include/utility/config.hpp index 92307214f4..e87b8a7b48 100644 --- a/src/composable_kernel/composable_kernel/include/utility/config.hpp +++ b/src/composable_kernel/composable_kernel/include/utility/config.hpp @@ -7,6 +7,19 @@ #endif #include "bfloat16_dev.hpp" +#ifdef __HIPCC_RTC__ +#ifdef WORKAROUND_ISSUE_HIPRTC_TRUE_TYPE +/// Definitions from , conflict with +/// /opt/rocm/include/hip/amd_detail/amd_hip_vector_types.h. + +typedef int int32_t; +typedef unsigned int uint32_t; + +#else +#include // int8_t, int16_t +#endif +#endif // __HIPCC_RTC__ + // "Constant" address space for kernel parameter #define CONSTANT __attribute__((address_space(4))) diff --git a/src/composable_kernel/composable_kernel/include/utility/type.hpp b/src/composable_kernel/composable_kernel/include/utility/type.hpp index 4e5d4e5134..10ec197d07 100644 --- a/src/composable_kernel/composable_kernel/include/utility/type.hpp +++ b/src/composable_kernel/composable_kernel/include/utility/type.hpp @@ -14,53 +14,6 @@ namespace std { -template -struct remove_reference -{ - typedef T type; -}; -template -struct remove_reference -{ - typedef T type; -}; -template -struct remove_reference -{ - typedef T type; -}; - -template -using remove_reference_t = typename remove_reference::type; - -template -struct remove_const -{ - typedef T type; -}; -template -struct remove_const -{ - typedef T type; -}; - -template -struct remove_volatile -{ - typedef T type; -}; -template -struct remove_volatile -{ - typedef T type; -}; - -template -struct remove_cv -{ - typedef typename remove_volatile::type>::type type; -}; - template struct is_pointer_helper : std::false_type { @@ -71,11 +24,6 @@ struct is_pointer_helper : std::true_type { }; -template -struct is_pointer : is_pointer_helper::type> -{ -}; - } // namespace std #else #include // std::remove_reference, std::remove_cv, is_pointer From 12237d21609064f9a9767f0a00489d8afe11bffd Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Sun, 17 Dec 2023 21:33:42 -0800 Subject: [PATCH 2/6] Revert "attempt to fix found issue with hiprtc" This reverts commit 3da00458ac3a515da1709a307f04cd846408e182. --- .../include/utility/config.hpp | 13 ----- .../include/utility/type.hpp | 52 +++++++++++++++++++ 2 files changed, 52 insertions(+), 13 deletions(-) diff --git a/src/composable_kernel/composable_kernel/include/utility/config.hpp b/src/composable_kernel/composable_kernel/include/utility/config.hpp index e87b8a7b48..92307214f4 100644 --- a/src/composable_kernel/composable_kernel/include/utility/config.hpp +++ b/src/composable_kernel/composable_kernel/include/utility/config.hpp @@ -7,19 +7,6 @@ #endif #include "bfloat16_dev.hpp" -#ifdef __HIPCC_RTC__ -#ifdef WORKAROUND_ISSUE_HIPRTC_TRUE_TYPE -/// Definitions from , conflict with -/// /opt/rocm/include/hip/amd_detail/amd_hip_vector_types.h. - -typedef int int32_t; -typedef unsigned int uint32_t; - -#else -#include // int8_t, int16_t -#endif -#endif // __HIPCC_RTC__ - // "Constant" address space for kernel parameter #define CONSTANT __attribute__((address_space(4))) diff --git a/src/composable_kernel/composable_kernel/include/utility/type.hpp b/src/composable_kernel/composable_kernel/include/utility/type.hpp index 10ec197d07..4e5d4e5134 100644 --- a/src/composable_kernel/composable_kernel/include/utility/type.hpp +++ b/src/composable_kernel/composable_kernel/include/utility/type.hpp @@ -14,6 +14,53 @@ namespace std { +template +struct remove_reference +{ + typedef T type; +}; +template +struct remove_reference +{ + typedef T type; +}; +template +struct remove_reference +{ + typedef T type; +}; + +template +using remove_reference_t = typename remove_reference::type; + +template +struct remove_const +{ + typedef T type; +}; +template +struct remove_const +{ + typedef T type; +}; + +template +struct remove_volatile +{ + typedef T type; +}; +template +struct remove_volatile +{ + typedef T type; +}; + +template +struct remove_cv +{ + typedef typename remove_volatile::type>::type type; +}; + template struct is_pointer_helper : std::false_type { @@ -24,6 +71,11 @@ struct is_pointer_helper : std::true_type { }; +template +struct is_pointer : is_pointer_helper::type> +{ +}; + } // namespace std #else #include // std::remove_reference, std::remove_cv, is_pointer From df0b906bb545ca8f380d1dbc01c144464eff6435 Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Sun, 17 Dec 2023 21:40:11 -0800 Subject: [PATCH 3/6] Use HIP version instead of removing code --- .../composable_kernel/include/utility/config.hpp | 13 +++++++++++++ .../composable_kernel/include/utility/type.hpp | 9 +++++---- 2 files changed, 18 insertions(+), 4 deletions(-) diff --git a/src/composable_kernel/composable_kernel/include/utility/config.hpp b/src/composable_kernel/composable_kernel/include/utility/config.hpp index 92307214f4..b9b6ff168f 100644 --- a/src/composable_kernel/composable_kernel/include/utility/config.hpp +++ b/src/composable_kernel/composable_kernel/include/utility/config.hpp @@ -7,6 +7,19 @@ #endif #include "bfloat16_dev.hpp" +#ifdef __HIPCC_RTC__ && HIP_PACKAGE_VERSION_FLAT >= 6000023494ULL +#ifdef WORKAROUND_ISSUE_HIPRTC_TRUE_TYPE +/// Definitions from , conflict with +/// /opt/rocm/include/hip/amd_detail/amd_hip_vector_types.h. + +typedef int int32_t; +typedef unsigned int uint32_t; + +#else +#include // int8_t, int16_t +#endif +#endif // __HIPCC_RTC__ + // "Constant" address space for kernel parameter #define CONSTANT __attribute__((address_space(4))) diff --git a/src/composable_kernel/composable_kernel/include/utility/type.hpp b/src/composable_kernel/composable_kernel/include/utility/type.hpp index 4e5d4e5134..0e601c48b1 100644 --- a/src/composable_kernel/composable_kernel/include/utility/type.hpp +++ b/src/composable_kernel/composable_kernel/include/utility/type.hpp @@ -13,7 +13,7 @@ /// which defines std::true_type as well (which is wrong). namespace std { - +if HIP_PACKAGE_VERSION_FLAT < 6000023494ULL template struct remove_reference { @@ -62,17 +62,18 @@ struct remove_cv }; template -struct is_pointer_helper : std::false_type +struct is_pointer : is_pointer_helper::type> { }; +#endif template -struct is_pointer_helper : std::true_type +struct is_pointer_helper : std::false_type { }; template -struct is_pointer : is_pointer_helper::type> +struct is_pointer_helper : std::true_type { }; From df7ebce9647464e4fbd9e431cd19a64b29b44987 Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Sun, 17 Dec 2023 21:52:15 -0800 Subject: [PATCH 4/6] Remove prolematic GitHub action --- .github/workflows/update_develop_nightly.yml | 20 -------------------- 1 file changed, 20 deletions(-) delete mode 100644 .github/workflows/update_develop_nightly.yml diff --git a/.github/workflows/update_develop_nightly.yml b/.github/workflows/update_develop_nightly.yml deleted file mode 100644 index a8b5aabdb9..0000000000 --- a/.github/workflows/update_develop_nightly.yml +++ /dev/null @@ -1,20 +0,0 @@ -name: Sync branch - -on: - pull_request: - workflow_dispatch: - push: - branches: - - develop -jobs: - sync-branch: - name: Update nightly branch - runs-on: ubuntu-latest - steps: - - name: Checkout repository - uses: actions/checkout@main - - uses: connor-baer/action-sync-branch@main - with: - branch: develop_nightly - token: ${{ secrets.GITHUB_TOKEN }} - force: false \ No newline at end of file From dbbcd837ba29adc85dcdf75fad6e57b274fc68b7 Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Mon, 18 Dec 2023 00:00:23 -0800 Subject: [PATCH 5/6] fix clang tidy issues --- .../composable_kernel/include/utility/config.hpp | 4 ++-- .../composable_kernel/include/utility/type.hpp | 12 +++++++----- 2 files changed, 9 insertions(+), 7 deletions(-) diff --git a/src/composable_kernel/composable_kernel/include/utility/config.hpp b/src/composable_kernel/composable_kernel/include/utility/config.hpp index b9b6ff168f..198d282de5 100644 --- a/src/composable_kernel/composable_kernel/include/utility/config.hpp +++ b/src/composable_kernel/composable_kernel/include/utility/config.hpp @@ -7,8 +7,8 @@ #endif #include "bfloat16_dev.hpp" -#ifdef __HIPCC_RTC__ && HIP_PACKAGE_VERSION_FLAT >= 6000023494ULL -#ifdef WORKAROUND_ISSUE_HIPRTC_TRUE_TYPE +#ifdef __HIPCC_RTC__ +#ifdef WORKAROUND_ISSUE_HIPRTC_TRUE_TYPE && HIP_PACKAGE_VERSION_FLAT >= 6000023494ULL /// Definitions from , conflict with /// /opt/rocm/include/hip/amd_detail/amd_hip_vector_types.h. diff --git a/src/composable_kernel/composable_kernel/include/utility/type.hpp b/src/composable_kernel/composable_kernel/include/utility/type.hpp index 0e601c48b1..3b43117410 100644 --- a/src/composable_kernel/composable_kernel/include/utility/type.hpp +++ b/src/composable_kernel/composable_kernel/include/utility/type.hpp @@ -13,12 +13,13 @@ /// which defines std::true_type as well (which is wrong). namespace std { -if HIP_PACKAGE_VERSION_FLAT < 6000023494ULL +if HIP_PACKAGE_VERSION_FLAT < 6000023494ULL + template -struct remove_reference -{ - typedef T type; -}; + struct remove_reference + { + typedef T type; + }; template struct remove_reference { @@ -65,6 +66,7 @@ template struct is_pointer : is_pointer_helper::type> { }; + #endif template From af8bdbfa94d1d118e7692504cd02946c90f946c5 Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Mon, 18 Dec 2023 00:04:16 -0800 Subject: [PATCH 6/6] fix clang format issues --- .../composable_kernel/include/utility/config.hpp | 4 +++- .../composable_kernel/include/utility/type.hpp | 10 +++++----- 2 files changed, 8 insertions(+), 6 deletions(-) diff --git a/src/composable_kernel/composable_kernel/include/utility/config.hpp b/src/composable_kernel/composable_kernel/include/utility/config.hpp index 198d282de5..f481498d56 100644 --- a/src/composable_kernel/composable_kernel/include/utility/config.hpp +++ b/src/composable_kernel/composable_kernel/include/utility/config.hpp @@ -8,13 +8,15 @@ #include "bfloat16_dev.hpp" #ifdef __HIPCC_RTC__ -#ifdef WORKAROUND_ISSUE_HIPRTC_TRUE_TYPE && HIP_PACKAGE_VERSION_FLAT >= 6000023494ULL +#ifdef WORKAROUND_ISSUE_HIPRTC_TRUE_TYPE +#if HIP_PACKAGE_VERSION_FLAT >= 6000023494ULL /// Definitions from , conflict with /// /opt/rocm/include/hip/amd_detail/amd_hip_vector_types.h. typedef int int32_t; typedef unsigned int uint32_t; +#endif #else #include // int8_t, int16_t #endif diff --git a/src/composable_kernel/composable_kernel/include/utility/type.hpp b/src/composable_kernel/composable_kernel/include/utility/type.hpp index 3b43117410..069761809d 100644 --- a/src/composable_kernel/composable_kernel/include/utility/type.hpp +++ b/src/composable_kernel/composable_kernel/include/utility/type.hpp @@ -13,13 +13,13 @@ /// which defines std::true_type as well (which is wrong). namespace std { -if HIP_PACKAGE_VERSION_FLAT < 6000023494ULL +#if(HIP_PACKAGE_VERSION_FLAT < 6000023494ULL) template - struct remove_reference - { - typedef T type; - }; +struct remove_reference +{ + typedef T type; +}; template struct remove_reference {