"git@developer.sourcefind.cn:cnjsdfcy/simbricks.git" did not exist on "c35116cfda06871a82c143f418232da4f76378db"
Commit e83439ad authored by Po-Yen, Chen's avatar Po-Yen, Chen
Browse files

Merge branch 'feature/support-readfirstlane-for-object-types' into...

Merge branch 'feature/support-readfirstlane-for-object-types' into feature/simplify-karg-for-device-gemm-xdl
parents fb130f76 a609bfaa
...@@ -7,6 +7,7 @@ ...@@ -7,6 +7,7 @@
#include "ck/utility/functional2.hpp" #include "ck/utility/functional2.hpp"
#include "ck/utility/math.hpp" #include "ck/utility/math.hpp"
#include <cstddef>
#include <cstdint> #include <cstdint>
#include <type_traits> #include <type_traits>
...@@ -54,9 +55,9 @@ __device__ auto readfirstlane(const Object& obj) ...@@ -54,9 +55,9 @@ __device__ auto readfirstlane(const Object& obj)
using Sgpr = detail::get_signed_int_t<SgprSize>; using Sgpr = detail::get_signed_int_t<SgprSize>;
alignas(Object) unsigned char memory[ObjectSize]; alignas(Object) std::byte memory[ObjectSize];
const auto* from = reinterpret_cast<const unsigned char*>(&obj); const auto* from = reinterpret_cast<const std::byte*>(&obj);
static_for<0, ObjectSize, SgprSize>{}([&](auto offset) { static_for<0, ObjectSize, SgprSize>{}([&](auto offset) {
*reinterpret_cast<Sgpr*>(memory + offset) = *reinterpret_cast<Sgpr*>(memory + offset) =
readfirstlane(*reinterpret_cast<const Sgpr*>(from + offset)); readfirstlane(*reinterpret_cast<const Sgpr*>(from + offset));
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment