rocm/libraries/rocroller 7.2.4

https://github.com/ROCm/rocm-libraries/releases/download/rocm-7.2.4/rocroller.tar.gz

cp /pkg/gnu/rocm/libraries/rocroller/src/lib/include/rocRoller/InstructionValues/Register_impl.hpp /pkg/gnu/rocm/libraries/rocroller/src/lib/include/rocRoller/InstructionValues/Register_impl.hpp.bkp
cp -f /pkg/gnu/rocm/libraries/rocroller/src/lib/include/rocRoller/InstructionValues/Register_impl.hpp.bkp /pkg/gnu/rocm/libraries/rocroller/src/lib/include/rocRoller/InstructionValues/Register_impl.hpp
sed -i 's/return val == 0;/return val == static_cast<T>(0);/g' \
    /pkg/gnu/rocm/libraries/rocroller/src/lib/include/rocRoller/InstructionValues/Register_impl.hpp

# 1. Восстанавливаем оригинал
cp -f /pkg/gnu/rocm/libraries/rocroller/src/lib/include/rocRoller/InstructionValues/Register_impl.hpp.bkp \
      /pkg/gnu/rocm/libraries/rocroller/src/lib/include/rocRoller/InstructionValues/Register_impl.hpp

# 2. Применяем исправленный патч
# Заменяем "return val == 0;" на "return val == T{};"
# T{} для арифметических типов = 0, для __half = __half(), для E8M0 = E8M0() (default ctor)
sed -i 's/return val == 0;/return val == T{};/g' \
    /pkg/gnu/rocm/libraries/rocroller/src/lib/include/rocRoller/InstructionValues/Register_impl.hpp

mkdir -p build; cd build
HIP_DEVICE_LIB_PATH=/pkg/gnu/rocm/llvm/amd/device-libs/amdgcn/bitcode \
HIP_PATH=/pkg/gnu/rocm/systems/clr \
CC=/pkg/gnu/rocm/llvm/amd/hipcc/lib/llvm/bin/clang \
CXX=/pkg/gnu/rocm/llvm/amd/hipcc/bin/hipcc \
cmake .. \
  -DCMAKE_INSTALL_PREFIX:PATH=/pkg/gnu/rocm/libraries/rocroller \
  -DCMAKE_BUILD_TYPE=Release \
  -DROCROLLER_BUILD_TESTING=OFF \
  -DROCROLLER_ENABLE_CLIENT=OFF \
  -DGPU_BUILD_TARGETS=gfx1150 \
  -DCMAKE_CXX_COMPILER=/pkg/gnu/rocm/llvm/amd/hipcc/bin/hipcc \
  -DCMAKE_LINKER=/pkg/gnu/rocm/llvm/amd/hipcc/lib/llvm/bin/ld.lld \
  -DCMAKE_CXX_FLAGS="-fuse-ld=lld -Wno-deprecated-declarations"
CC=/pkg/gnu/rocm/llvm/amd/hipcc/lib/llvm/bin/clang \
CXX=/pkg/gnu/rocm/llvm/amd/hipcc/bin/hipcc \
HIP_DEVICE_LIB_PATH=/pkg/gnu/rocm/llvm/amd/device-libs/amdgcn/bitcode \
HIP_PATH=/pkg/gnu/rocm/systems/clr \
cmake --build . -j$(nproc)
cmake --install .

rocm/llvm/libcxx
fmt
spdlog
libdivide
yaml-cpp
msgpack-cxx

sed -i '/inline bool Value::isZeroLiteral() const/,/^ }$/{ /inline bool Value::isZeroLiteral() const/,/^ }$/c
inline bool Value::isZeroLiteral() const
{
if(m_regType != Type::Literal)
return false;

return std::visit(
[](auto const& val) -> bool {
using T = std::decay_t<decltype(val)>;
if constexpr(std::is_same_v<T, Raw32>)
return static_cast<uint32_t>(val) == 0;
else if constexpr(std::is_same_v<T, Buffer>)
return val.desc0 == 0 && val.desc1 == 0 && val.desc2 == 0 && val.desc3 == 0;
else if constexpr(std::is_integral_v)
return val == 0;
else if constexpr(std::is_floating_point_v || std::is_same_v<T, __half>)
return val == static_cast(0);
else
return false;
},
m_literalValue);
} }' Register_impl.hpp

sed -i '/inline bool Value::isZeroLiteral() const/,/^ }$/{ /inline bool Value::isZeroLiteral() const/,/^ }$/c
inline bool Value::isZeroLiteral() const
{
if(m_regType != Type::Literal)
return false;

return std::visit(
[](auto const& val) -> bool {
using T = std::decay_t<decltype(val)>;
if constexpr(std::is_same_v<T, Raw32>)
return static_cast<uint32_t>(val) == 0;
else if constexpr(std::is_same_v<T, Buffer>)
return val.desc0 == 0 && val.desc1 == 0 && val.desc2 == 0 && val.desc3 == 0;
else if constexpr(std::is_arithmetic_v || std::is_same_v<T, __half>)
return val == T{0};
else
return false;
},
m_literalValue);
} }' Register_impl.hpp

-DCMAKE_C_COMPILER=/pkg/gnu/rocm/llvm/amd/hipcc/lib/llvm/bin/clang
-DCMAKE_CXX_FLAGS="-fuse-ld=lld -stdlib=libc++"

sed -i 's/return val == static_cast(0);/if constexpr(std::is_same_v<T, rocRoller::E8M0>) return false; else return val == static_cast(0);/g' Register_impl.hpp

sed -i '/return val == static_cast(0);/c
else if constexpr(std::is_arithmetic_v || std::is_same_v<T, __half>)
return val == static_cast(0);
else
return false;' Register_impl.hpp

sed -i '/inline bool Value::isZeroLiteral/,/^ }$/c
inline bool Value::isZeroLiteral() const
{
if(m_regType != Type::Literal)
return false;

return std::visit(
[](auto const& val) -> bool {
using T = std::decay_t<decltype(val)>;
if constexpr(std::is_same_v<T, Raw32>)
return static_cast<uint32_t>(val) == 0;
else if constexpr(std::is_same_v<T, Buffer>)
return val.desc0 == 0 && val.desc1 == 0 && val.desc2 == 0 && val.desc3 == 0;
else if constexpr(std::is_arithmetic_v || std::is_same_v<T, __half>
|| std::is_same_v<T, rocRoller::BFloat16>
|| std::is_same_v<T, rocRoller::FP8>
|| std::is_same_v<T, rocRoller::BF8>
|| std::is_same_v<T, rocRoller::FP6>
|| std::is_same_v<T, rocRoller::BF6>
|| std::is_same_v<T, rocRoller::FP4>
|| std::is_same_v<T, rocRoller::E8M0>)
return val == static_cast(0);
else
return false;
},
m_literalValue);
}' Register_impl.hpp

sed -i '/return std::visit(/,/m_literalValue);/{ s/return static_cast(val) == 0.0f;/using T = std::decay_t<decltype(val)>;
if constexpr(std::is_same_v<T, Raw32>)
return static_cast<uint32_t>(val) == 0;
else if constexpr(std::is_same_v<T, Buffer>)
return val.desc0 == 0 && val.desc1 == 0 && val.desc2 == 0 && val.desc3 == 0;
else if constexpr(std::is_arithmetic_v)
return val == decltype(val)(0);
else
return false;/ }' /pkg/gnu/rocm/libraries/rocroller/src/lib/include/rocRoller/InstructionValues/Register_impl.hpp

sed -i '/return std::visit(/,/m_literalValue);/{ s/return val == 0;/using T = std::decay_t<decltype(val)>;
if constexpr(std::is_same_v<T, Raw32>)
return static_cast<uint32_t>(val) == 0;
else if constexpr(std::is_same_v<T, Buffer>)
return val.desc0 == 0 && val.desc1 == 0 && val.desc2 == 0 && val.desc3 == 0;
else if constexpr(std::is_same_v<T, __half>)
return static_cast(val) == 0.0f;
else if constexpr(std::is_arithmetic_v)
return val == 0;
else
return false;/ }' /pkg/gnu/rocm/libraries/rocroller/src/lib/include/rocRoller/InstructionValues/Register_impl.hpp

7.8 KiB Raw Permalink Blame History

rocm/libraries/rocroller 7.2.4

7.8 KiB

Raw Permalink Blame History