Skip to content

Commit

Permalink
lint
Browse files Browse the repository at this point in the history
  • Loading branch information
Saeed Maleki committed Aug 1, 2023
1 parent 9d6e51e commit f217876
Show file tree
Hide file tree
Showing 2 changed files with 9 additions and 10 deletions.
18 changes: 9 additions & 9 deletions include/mscclpp/sm_channel_device.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -4,9 +4,9 @@
#ifndef MSCCLPP_SM_CHANNEL_DEVICE_HPP_
#define MSCCLPP_SM_CHANNEL_DEVICE_HPP_

#include "semaphore_device.hpp"
#include "poll.hpp"
#include "packet.hpp"
#include "poll.hpp"
#include "semaphore_device.hpp"

namespace mscclpp {

Expand Down Expand Up @@ -105,19 +105,19 @@ struct Element<uint4> {
using T = uint4;
static __forceinline__ __device__ void load(T& v, const T* p) {
asm volatile("ld.volatile.global.v4.u32 {%0,%1,%2,%3}, [%4];"
: "=r"(v.w), "=r"(v.x), "=r"(v.y), "=r"(v.z)
: "l"(p)
: "memory");
: "=r"(v.w), "=r"(v.x), "=r"(v.y), "=r"(v.z)
: "l"(p)
: "memory");
}

static __forceinline__ __device__ void store(T* p, const T& v) {
asm volatile("st.volatile.global.v4.u32 [%0], {%1,%2,%3,%4};"
:
: "l"(p), "r"(v.w), "r"(v.x), "r"(v.y), "r"(v.z)
: "memory");
:
: "l"(p), "r"(v.w), "r"(v.x), "r"(v.y), "r"(v.z)
: "memory");
}
};
#endif // __CUDACC__
#endif // __CUDACC__

/// Channel for accessing peer memory directly from SM.
struct SmChannelDeviceHandle {
Expand Down
1 change: 0 additions & 1 deletion python/sm_channel_py.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -29,5 +29,4 @@ void register_sm_channel(nb::module_& m) {
.def_prop_ro("raw", [](const SmChannel::DeviceHandle& self) -> nb::bytes {
return nb::bytes(reinterpret_cast<const char*>(&self), sizeof(self));
});

};

0 comments on commit f217876

Please sign in to comment.