169 lines
7.8 KiB
Markdown
169 lines
7.8 KiB
Markdown
|
|
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
|
||
|
|
```
|
||
|
|
|
||
|
|
|
||
|
|
``` sh *** build ***
|
||
|
|
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 .
|
||
|
|
```
|
||
|
|
|
||
|
|
``` cfg *** build deps ***
|
||
|
|
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<T>)\
|
||
|
|
return val == 0;\
|
||
|
|
else if constexpr(std::is_floating_point_v<T> || std::is_same_v<T, __half>)\
|
||
|
|
return val == static_cast<T>(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<T> || 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<T>(0);/if constexpr(std::is_same_v<T, rocRoller::E8M0>) return false; else return val == static_cast<T>(0);/g' Register_impl.hpp
|
||
|
|
|
||
|
|
sed -i '/return val == static_cast<T>(0);/c\
|
||
|
|
else if constexpr(std::is_arithmetic_v<T> || std::is_same_v<T, __half>)\
|
||
|
|
return val == static_cast<T>(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<T> || 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<T>(0);\
|
||
|
|
else\
|
||
|
|
return false;\
|
||
|
|
},\
|
||
|
|
m_literalValue);\
|
||
|
|
}' Register_impl.hpp
|
||
|
|
|
||
|
|
|
||
|
|
sed -i '/return std::visit(/,/m_literalValue);/{
|
||
|
|
s/return static_cast<float>(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<T>)\
|
||
|
|
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<float>(val) == 0.0f;\
|
||
|
|
else if constexpr(std::is_arithmetic_v<T>)\
|
||
|
|
return val == 0;\
|
||
|
|
else\
|
||
|
|
return false;/
|
||
|
|
}' /pkg/gnu/rocm/libraries/rocroller/src/lib/include/rocRoller/InstructionValues/Register_impl.hpp
|