diff --git a/hipamd/hip-config-clang.cmake.in b/hipamd/hip-config-clang.cmake.in
index e4862f3d6c..9680b2ff8e 100644
--- a/hipamd/hip-config-clang.cmake.in
+++ b/hipamd/hip-config-clang.cmake.in
@@ -50,6 +50,7 @@ set_and_check( hip_BIN_INSTALL_DIR "@PACKAGE_BIN_INSTALL_DIR@" )
set_and_check(hip_HIPCC_EXECUTABLE "${hip_BIN_INSTALL_DIR}/hipcc")
set_and_check(hip_HIPCONFIG_EXECUTABLE "${hip_BIN_INSTALL_DIR}/hipconfig")
+find_dependency(amd_comgr)
include( "${CMAKE_CURRENT_LIST_DIR}/hip-targets.cmake" )
set( hip_LIBRARIES hip::host hip::device)
diff --git a/hipamd/hip-config-hcc.cmake.in b/hipamd/hip-config-hcc.cmake.in
index 7e236ba0b4..c0ffc6e2af 100644
--- a/hipamd/hip-config-hcc.cmake.in
+++ b/hipamd/hip-config-hcc.cmake.in
@@ -51,6 +51,7 @@ set_and_check(hip_HIPCC_EXECUTABLE "${hip_BIN_INSTALL_DIR}/hipcc")
set_and_check(hip_HIPCONFIG_EXECUTABLE "${hip_BIN_INSTALL_DIR}/hipconfig")
find_dependency(hcc)
+find_dependency(amd_comgr)
include( "${CMAKE_CURRENT_LIST_DIR}/hip-targets.cmake" )
set( hip_LIBRARIES hip::host hip::device)
diff --git a/hipamd/hipify-clang/README.md b/hipamd/hipify-clang/README.md
index 7744085f16..51fcc08bc1 100644
--- a/hipamd/hipify-clang/README.md
+++ b/hipamd/hipify-clang/README.md
@@ -53,9 +53,9 @@ If the target CUDA is [9.1](https://developer.nvidia.com/cuda-91-download-archiv
| [**6.0.1**](http://releases.llvm.org/download.html#6.0.1) | [**9.0**](https://developer.nvidia.com/cuda-90-download-archive) | +
**LATEST STABLE RELEASE** | + |
| [7.0.0](http://releases.llvm.org/download.html#7.0.0) | [9.2](https://developer.nvidia.com/cuda-92-download-archive) | -
not working due to
the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811)
+
[patch](patches/patch_for_clang_7.0.0_bug_38811.zip)* | -
not working due to
the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
| [7.0.1](http://releases.llvm.org/download.html#7.0.1) | [9.2](https://developer.nvidia.com/cuda-92-download-archive) | -
not working due to
the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811)
+
[patch](patches/patch_for_clang_7.0.1_bug_38811.zip)* | -
not working due to
the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
-| 7.1.0 | 9.2 (?) | -
LLVM 7.1.0
is not yet released | -
LLVM 7.1.0
is not yet released |
+| [7.1.0](http://releases.llvm.org/download.html#7.1.0) | [9.2](https://developer.nvidia.com/cuda-92-download-archive) | -
not working due to
the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811)
+
[patch](patches/patch_for_clang_7.1.0_bug_38811.zip)* | -
not working due to
the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
| [**8.0.0**](http://releases.llvm.org/download.html#8.0.0) | [**10.0**](https://developer.nvidia.com/cuda-10.0-download-archive) | -
not working due to
the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811)
+
[patch](patches/patch_for_clang_8.0.0_bug_38811.zip)* | +
**LATEST STABLE RELEASE** |
-| 8.0.1 | 10.1 (?) | -
LLVM 8.0.1
is not yet released | -
LLVM 8.0.1
is not yet released |
+| 8.0.1 | 10.2 | -
LLVM 8.0.1
is not yet released | -
LLVM 8.0.1
is not yet released |
`*` Download the patch and unpack it into your LLVM distributive directory; a few header files will be overwritten; rebuilding of LLVM is not needed.
diff --git a/hipamd/hipify-clang/patches/patch_for_clang_7.1.0_bug_38811.zip b/hipamd/hipify-clang/patches/patch_for_clang_7.1.0_bug_38811.zip
new file mode 100644
index 0000000000..35bde0cce4
Binary files /dev/null and b/hipamd/hipify-clang/patches/patch_for_clang_7.1.0_bug_38811.zip differ
diff --git a/hipamd/include/hip/hcc_detail/hip_vector_types.h b/hipamd/include/hip/hcc_detail/hip_vector_types.h
index 49df41f5a7..0d424965ca 100644
--- a/hipamd/include/hip/hcc_detail/hip_vector_types.h
+++ b/hipamd/include/hip/hcc_detail/hip_vector_types.h
@@ -76,7 +76,151 @@ THE SOFTWARE.
template
struct HIP_vector_base {
- typedef T Native_vec_ __NATIVE_VECTOR__(3, T);
+ struct Native_vec_ {
+ T d[3];
+
+ __host__ __device__
+ constexpr
+ Native_vec_() = default;
+ __host__ __device__
+ explicit
+ constexpr
+ Native_vec_(T x) noexcept : d{x, x, x} {}
+ __host__ __device__
+ constexpr
+ Native_vec_(T x, T y, T z) noexcept : d{x, y, z} {}
+ __host__ __device__
+ constexpr
+ Native_vec_(const Native_vec_&) = default;
+ __host__ __device__
+ constexpr
+ Native_vec_(Native_vec_&&) = default;
+ __host__ __device__
+ ~Native_vec_() = default;
+
+ __host__ __device__
+ Native_vec_& operator=(const Native_vec_&) = default;
+ __host__ __device__
+ Native_vec_& operator=(Native_vec_&&) = default;
+
+ __host__ __device__
+ T& operator[](unsigned int idx) noexcept { return d[idx]; }
+ __host__ __device__
+ T operator[](unsigned int idx) const noexcept { return d[idx]; }
+
+ __host__ __device__
+ Native_vec_& operator+=(const Native_vec_& x) noexcept
+ {
+ for (auto i = 0u; i != 3u; ++i) d[i] += x.d[i];
+ return *this;
+ }
+ __host__ __device__
+ Native_vec_& operator-=(const Native_vec_& x) noexcept
+ {
+ for (auto i = 0u; i != 3u; ++i) d[i] -= x.d[i];
+ return *this;
+ }
+
+ __host__ __device__
+ Native_vec_& operator*=(const Native_vec_& x) noexcept
+ {
+ for (auto i = 0u; i != 3u; ++i) d[i] *= x.d[i];
+ return *this;
+ }
+ __host__ __device__
+ Native_vec_& operator/=(const Native_vec_& x) noexcept
+ {
+ for (auto i = 0u; i != 3u; ++i) d[i] /= x.d[i];
+ return *this;
+ }
+
+ template<
+ typename U = T,
+ typename std::enable_if{}>::type* = nullptr>
+ __host__ __device__
+ Native_vec_ operator-() const noexcept
+ {
+ auto r{*this};
+ for (auto&& x : r.d) x = -x;
+ return r;
+ }
+
+ template<
+ typename U = T,
+ typename std::enable_if{}>::type* = nullptr>
+ __host__ __device__
+ Native_vec_ operator~() const noexcept
+ {
+ auto r{*this};
+ for (auto&& x : r.d) x = ~x;
+ return r;
+ }
+ template<
+ typename U = T,
+ typename std::enable_if{}>::type* = nullptr>
+ __host__ __device__
+ Native_vec_& operator%=(const Native_vec_& x) noexcept
+ {
+ for (auto i = 0u; i != 3u; ++i) d[i] %= x.d[i];
+ return *this;
+ }
+ template<
+ typename U = T,
+ typename std::enable_if{}>::type* = nullptr>
+ __host__ __device__
+ Native_vec_& operator^=(const Native_vec_& x) noexcept
+ {
+ for (auto i = 0u; i != 3u; ++i) d[i] ^= x.d[i];
+ return *this;
+ }
+ template<
+ typename U = T,
+ typename std::enable_if{}>::type* = nullptr>
+ __host__ __device__
+ Native_vec_& operator|=(const Native_vec_& x) noexcept
+ {
+ for (auto i = 0u; i != 3u; ++i) d[i] |= x.d[i];
+ return *this;
+ }
+ template<
+ typename U = T,
+ typename std::enable_if{}>::type* = nullptr>
+ __host__ __device__
+ Native_vec_& operator&=(const Native_vec_& x) noexcept
+ {
+ for (auto i = 0u; i != 3u; ++i) d[i] &= x.d[i];
+ return *this;
+ }
+ template<
+ typename U = T,
+ typename std::enable_if{}>::type* = nullptr>
+ __host__ __device__
+ Native_vec_& operator>>=(const Native_vec_& x) noexcept
+ {
+ for (auto i = 0u; i != 3u; ++i) d[i] >>= x.d[i];
+ return *this;
+ }
+ template<
+ typename U = T,
+ typename std::enable_if{}>::type* = nullptr>
+ __host__ __device__
+ Native_vec_& operator<<=(const Native_vec_& x) noexcept
+ {
+ for (auto i = 0u; i != 3u; ++i) d[i] <<= x.d[i];
+ return *this;
+ }
+
+ using Vec3_cmp = T __NATIVE_VECTOR__(3, int);
+ __host__ __device__
+ Vec3_cmp operator==(const Native_vec_& x) const noexcept
+ {
+ Vec3_cmp r;
+ r[0] = d[0] == x.d[0];
+ r[1] = d[1] == x.d[1];
+ r[2] = d[2] == x.d[2];
+ return -r;
+ }
+ };
union {
Native_vec_ data;
@@ -632,7 +776,7 @@ THE SOFTWARE.
} CUDA_name##2;\
typedef struct {\
union {\
- CUDA_name##_impl3 data;\
+ T data[3];\
struct {\
T x;\
T y;\
diff --git a/hipamd/packaging/hip_hcc.txt b/hipamd/packaging/hip_hcc.txt
index 365af8ef1e..e2f912c7ca 100644
--- a/hipamd/packaging/hip_hcc.txt
+++ b/hipamd/packaging/hip_hcc.txt
@@ -33,9 +33,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}")
@@ -43,9 +43,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 >= 1.1")
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 >= 1.1")
endif()
set(CPACK_RPM_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "/opt")
set(CPACK_SOURCE_GENERATOR "TGZ")
diff --git a/hipamd/src/hip_module.cpp b/hipamd/src/hip_module.cpp
index bda6970298..9680299540 100644
--- a/hipamd/src/hip_module.cpp
+++ b/hipamd/src/hip_module.cpp
@@ -963,7 +963,7 @@ hipError_t ihipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* bloc
wavefrontsSGPRS = maxWavesWGLimited;
}
else {
- const size_t numSGPRsPerSIMD = (prop.gcnArch < 900) ? 512 : 800;
+ const size_t numSGPRsPerSIMD = (prop.gcnArch < 800) ? 512 : 800;
wavefrontsSGPRS = (numSGPRsPerSIMD / usedSGPRS) * numSIMD;
}
diff --git a/hipamd/src/program_state.inl b/hipamd/src/program_state.inl
index b01c22c6c6..f6cd290764 100644
--- a/hipamd/src/program_state.inl
+++ b/hipamd/src/program_state.inl
@@ -219,8 +219,8 @@ public:
dl_iterate_phdr([](dl_phdr_info* info, std::size_t, void* p) {
ELFIO::elfio tmp;
- const auto elf =
- info->dlpi_addr ? info->dlpi_name : "/proc/self/exe";
+ const auto elf = (info->dlpi_addr && std::strlen(info->dlpi_name) != 0) ?
+ info->dlpi_name : "/proc/self/exe";
if (!tmp.load(elf)) return 0;
@@ -277,8 +277,8 @@ public:
program_state_impl* t = static_cast(psi_ptr);
ELFIO::elfio tmp;
- const auto elf =
- info->dlpi_addr ? info->dlpi_name : "/proc/self/exe";
+ const auto elf = (info->dlpi_addr && std::strlen(info->dlpi_name) != 0) ?
+ info->dlpi_name : "/proc/self/exe";
if (!tmp.load(elf)) return 0;
@@ -503,8 +503,8 @@ public:
std::call_once(function_names.first, [this]() {
dl_iterate_phdr([](dl_phdr_info* info, std::size_t, void* p) {
ELFIO::elfio tmp;
- const auto elf =
- info->dlpi_addr ? info->dlpi_name : "/proc/self/exe";
+ const auto elf = (info->dlpi_addr && std::strlen(info->dlpi_name) != 0) ?
+ info->dlpi_name : "/proc/self/exe";
if (!tmp.load(elf)) return 0;
diff --git a/hipamd/tests/src/deviceLib/hipVectorTypes.cpp b/hipamd/tests/src/deviceLib/hipVectorTypes.cpp
index f05e6372d8..14479881ff 100644
--- a/hipamd/tests/src/deviceLib/hipVectorTypes.cpp
+++ b/hipamd/tests/src/deviceLib/hipVectorTypes.cpp
@@ -190,7 +190,7 @@ bool CheckVectorTypes() {
int main() {
static_assert(sizeof(float1) == 4, "");
static_assert(sizeof(float2) >= 8, "");
- static_assert(sizeof(float3) >= 12, "");
+ static_assert(sizeof(float3) == 12, "");
static_assert(sizeof(float4) >= 16, "");
if (CheckVectorTypes()) {