bytes.cuh

Go to the documentation of this file.

// SPDX-License-Identifier: Apache-2.0
 
#pragma once
#include <fss/group.cuh>
#include <cuda_runtime.h>
#include <cassert>
#include <fss/util.cuh>
 
namespace fss::group {

class Bytes {
public:
  // Groupable interface.

  __host__ __device__ Bytes operator+(Bytes rhs) const {
    return util::Xor(val_, rhs.val_);
  }

  __host__ __device__ Bytes operator-() const {
    return *this;
  }

  __host__ __device__ Bytes() : val_({0, 0, 0, 0}) {}

  __host__ __device__ static Bytes From(int4 buf) {
    util::RequireArg((buf.w & 1) == 0, "Bytes::From: the block must have a zero LSB");
    return Bytes(buf);
  }

  __host__ __device__ int4 Into() const {
    return val_;
  }
 
private:
  int4 val_;
 
  __host__ __device__ Bytes(int4 buf) : val_(buf) {}
};
static_assert(Groupable<Bytes>);
 
}  // namespace fss::group