AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge 'master' into 'amd-master'

Browse Source 
Change-Id: Ied8c47046b006252096e0381960b62f5aa95d6d3
This commit is contained in:
Jenkins
2019-06-29 21:09:16 -07:00
parent 178f06d22f 54dab53c04
commit 04cbe15857
9 changed files with 162 additions and 16 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/hip-config-clang.cmake.in
+1
View File
@@ -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)
hipamd/hip-config-hcc.cmake.in
+1
View File
@@ -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)
hipamd/hipify-clang/README.md
+2 -2
View File
@@ -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) | + <br/> **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) | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_7.0.0_bug_38811.zip)*</br> | - <br/> not working due to <br/> 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) | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_7.0.1_bug_38811.zip)*</br> | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
| 7.1.0 | 9.2 (?) | - <br/> LLVM 7.1.0 <br/> is not yet released | - <br/> LLVM 7.1.0 <br/> 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) | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_7.1.0_bug_38811.zip)*</br> | - <br/> not working due to <br/> 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) | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_8.0.0_bug_38811.zip)*</br> | + <br/> **LATEST STABLE RELEASE** |
| 8.0.1 | 10.1 (?) | - <br/> LLVM 8.0.1 <br/> is not yet released | - <br/> LLVM 8.0.1 <br/> is not yet released |
| 8.0.1 | 10.2 | - <br/> LLVM 8.0.1 <br/> is not yet released | - <br/> LLVM 8.0.1 <br/> 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.
hipamd/hipify-clang/patches/patch_for_clang_7.1.0_bug_38811.zip
BIN
View File
Binary file not shown.
hipamd/include/hip/hcc_detail/hip_vector_types.h
+146 -2
View File
@@ -76,7 +76,151 @@ THE SOFTWARE.
template<typename T>
struct HIP_vector_base<T, 3> {
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<std::is_signed<U>{}>::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<std::is_integral<U>{}>::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<std::is_integral<U>{}>::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<std::is_integral<U>{}>::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<std::is_integral<U>{}>::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<std::is_integral<U>{}>::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<std::is_integral<U>{}>::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<std::is_integral<U>{}>::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;\
hipamd/packaging/hip_hcc.txt
+4 -4
View File
@@ -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")
hipamd/src/hip_module.cpp
+1 -1
View File
@@ -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;
}
hipamd/src/program_state.inl
+6 -6
View File
@@ -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<program_state_impl*>(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;
hipamd/tests/src/deviceLib/hipVectorTypes.cpp
+1 -1
View File
@@ -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()) {