Skip to content
Merged
Show file tree
Hide file tree
Changes from 1 commit
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion src/krepe/common/extended_lambda_utils.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -72,7 +72,7 @@ std::size_t hdl_host_lambda_size(
// We assume that it is the last member of __nv_hdl_wrapper_t.
template <class Functor>
void* hdl_host_lambda_pointer(Functor&& f) {
unsigned char* data = reinterpret_cast<unsigned char*>(&f);
std::byte* data = reinterpret_cast<std::byte*>(&f);
return *reinterpret_cast<void**>((data + sizeof(f)) - sizeof(void*));
}

Expand Down
13 changes: 7 additions & 6 deletions src/krepe/replay/allocation.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -3,6 +3,7 @@

#pragma once

#include <cstddef>
#include <string>
#include <tuple>

Expand All @@ -22,7 +23,7 @@

namespace krepe::impl {

inline std::pair<char*, std::size_t> get_allocation_address(
inline std::pair<std::byte*, std::size_t> get_allocation_address(
void* target_address, std::size_t size, MemorySpaceType space) {
std::size_t granularity;
if (space == MemorySpaceType::HOST) {
Expand All @@ -37,18 +38,18 @@ inline std::pair<char*, std::size_t> get_allocation_address(
#endif
}

char* aligned_address = reinterpret_cast<char*>(
std::byte* aligned_address = reinterpret_cast<std::byte*>(
((reinterpret_cast<std::uintptr_t>(target_address)) / granularity) *
Comment thread
science-enthusiast marked this conversation as resolved.
granularity);
size += reinterpret_cast<char*>(target_address) - aligned_address;
size += static_cast<std::byte*>(target_address) - aligned_address;

std::size_t total_size = (((size - 1) / granularity) + 1) * granularity;

return {aligned_address, total_size};
}

inline void copy_data(MemorySpaceType space, char* address, char* data,
std::size_t size) {
inline void copy_data(MemorySpaceType space, std::byte* address,
const std::byte* data, std::size_t size) {
if (space == MemorySpaceType::HOST) {
host_copy_data(address, data, size);
} else {
Expand All @@ -69,7 +70,7 @@ struct Allocation {

Allocation() = default;

Allocation(MemorySpaceType memory_space, char* target_address,
Allocation(MemorySpaceType memory_space, std::byte* target_address,
std::size_t requested_size)
: memory_space(memory_space) {
if (memory_space == MemorySpaceType::HOST) {
Expand Down
6 changes: 4 additions & 2 deletions src/krepe/replay/cuda_allocation.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -3,6 +3,7 @@

#pragma once

#include <cstddef>
#include <cstring>
#include <iostream>
#include <sstream>
Expand Down Expand Up @@ -39,7 +40,7 @@ inline void throw_error(CUresult error, const char* expr, const char* file,
#define CHECK_CUDA_CALL(expr) \
::krepe::impl::throw_error((expr), #expr, __FILE__, __LINE__)

inline auto regular_device_allocate(std::size_t size, char* data) {
inline auto regular_device_allocate(std::size_t size, const std::byte* data) {
void* ptr;
KOKKOS_IMPL_CUDA_SAFE_CALL(cudaMalloc(&ptr, size));
KOKKOS_IMPL_CUDA_SAFE_CALL(
Expand All @@ -48,7 +49,8 @@ inline auto regular_device_allocate(std::size_t size, char* data) {
ptr, [](void* ptr) { KOKKOS_IMPL_CUDA_SAFE_CALL(cudaFree(ptr)); });
}

inline void device_copy_data(char* address, char* data, std::size_t size) {
inline void device_copy_data(std::byte* address, const std::byte* data,
std::size_t size) {
KOKKOS_IMPL_CUDA_SAFE_CALL(
cudaMemcpy(address, data, size, cudaMemcpyHostToDevice));
}
Expand Down
6 changes: 4 additions & 2 deletions src/krepe/replay/hip_allocation.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -3,6 +3,7 @@

#pragma once

#include <cstddef>
#include <cstring>
#include <tuple>
#include <utility>
Expand All @@ -12,15 +13,16 @@
#define CHECK_HIP_CALL(expr) KOKKOS_IMPL_HIP_SAFE_CALL(expr)

namespace krepe::impl {
inline auto regular_device_allocate(std::size_t size, char* data) {
inline auto regular_device_allocate(std::size_t size, const std::byte* data) {
void* ptr;
CHECK_HIP_CALL(hipMalloc(&ptr, size));
CHECK_HIP_CALL(hipMemcpy(ptr, data, size, hipMemcpyHostToDevice));
return std::unique_ptr<void, void (*)(void*)>(
ptr, [](void* ptr) { CHECK_HIP_CALL(hipFree(ptr)); });
}

inline void device_copy_data(char* address, char* data, std::size_t size) {
inline void device_copy_data(std::byte* address, const std::byte* data,
std::size_t size) {
CHECK_HIP_CALL(hipMemcpy(address, data, size, hipMemcpyHostToDevice));
}

Expand Down
5 changes: 3 additions & 2 deletions src/krepe/replay/host_allocation.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -14,15 +14,16 @@

namespace krepe::impl {

inline auto regular_host_allocate(std::size_t size, char* data) {
inline auto regular_host_allocate(std::size_t size, const std::byte* data) {
void* ptr = std::malloc(size);
std::memcpy(ptr, data, size);
return std::unique_ptr<void, void (*)(void*)>(ptr, std::free);
}

inline std::size_t host_allocation_granularity() { return getpagesize(); }

inline void host_copy_data(char* address, char* data, std::size_t size) {
inline void host_copy_data(std::byte* address, const std::byte* data,
std::size_t size) {
std::memcpy(address, data, size);
}

Expand Down
58 changes: 31 additions & 27 deletions src/krepe/replay/kernel_replayer.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -4,6 +4,7 @@
#include <Kokkos_Core.hpp>
#include <algorithm>
#include <cassert>
#include <cstddef>
#include <exception>
#include <iterator>
#include <map>
Expand Down Expand Up @@ -94,9 +95,9 @@ void device_init() {
}
}

static std::vector<char> functor_data;
static std::vector<std::byte> functor_data;
#if defined(KERNEL_REPLAYER_USE_NVCC_HDL_WORKAROUND)
static std::vector<char> nvcc_inner_lambda_data;
static std::vector<std::byte> nvcc_inner_lambda_data;

namespace krepe::impl {
void* copy_extended_lambda_inner_lambda(void* inner_lambda_ptr,
Expand Down Expand Up @@ -250,7 +251,7 @@ void validate_comparison(const ReplayAllocation& allocation,
}
}

void init_functor(char* buffer, std::size_t size) {
void init_functor(std::byte* buffer, std::size_t size) {
if (functor_data.size() < size) {
throw std::runtime_error(
"The stored functor is smaller than the requested functor, expected at "
Expand All @@ -261,21 +262,21 @@ void init_functor(char* buffer, std::size_t size) {
std::memcpy(buffer, functor_data.data(), size);
}

char* allocate_host_buffer(std::size_t size, MemorySpaceType mem) {
std::byte* allocate_host_buffer(std::size_t size, MemorySpaceType mem) {
if (mem == MemorySpaceType::HOST) {
return new char[size];
return new std::byte[size];
} else {
void* ptr = nullptr;
#if defined(KOKKOS_ENABLE_CUDA)
KOKKOS_IMPL_CUDA_SAFE_CALL(cudaMallocHost(&ptr, size));
#elif defined(KOKKOS_ENABLE_HIP)
KOKKOS_IMPL_HIP_SAFE_CALL(hipHostMalloc(&ptr, size));
#endif
return static_cast<char*>(ptr);
return static_cast<std::byte*>(ptr);
}
}

void free_host_buffer(char* ptr, MemorySpaceType mem) {
void free_host_buffer(std::byte* ptr, MemorySpaceType mem) {
if (mem == MemorySpaceType::HOST) {
delete[] ptr;
} else {
Expand All @@ -290,14 +291,14 @@ void free_host_buffer(char* ptr, MemorySpaceType mem) {
struct SnapshotAllocation {
std::string label;
std::string space;
char* allocation_id;
char* address;
std::byte* allocation_id;
std::byte* address;
std::size_t size;
bool bytes_dumped;
};

using hdf5_iterate_fun_t =
std::function<void(const SnapshotAllocation&, char*)>;
std::function<void(const SnapshotAllocation&, const std::byte*)>;

void iterate_allocations(hid_t file, const char* path, H5L_iterate_t callback,
void* data) {
Expand Down Expand Up @@ -395,7 +396,7 @@ std::uint64_t get_hdf5_uint64_attribute(hid_t group, const char* name,

SnapshotAllocation read_allocation_info(hid_t group, const char* name) {
const auto pointer_attribute = [&](const char* attribute) {
return reinterpret_cast<char*>(std::stoull(
return reinterpret_cast<std::byte*>(std::stoull(
get_hdf5_string_attribute(group, name, attribute), nullptr, 16));
};
const auto size = get_hdf5_uint64_attribute(group, name, "size");
Expand Down Expand Up @@ -442,7 +443,8 @@ herr_t get_hdf5_dataset_alloc_info(hid_t group, const char* name,
check_byte_dataset(group, name, entry.size);
if (entry.size != 0) {
auto& allocations =
*static_cast<std::set<std::pair<char*, std::size_t>>*>(allocation_set);
*static_cast<std::set<std::pair<std::byte*, std::size_t>>*>(
allocation_set);
allocations.insert(
get_allocation_address(entry.address, entry.size, space));
}
Expand Down Expand Up @@ -470,8 +472,10 @@ herr_t allocate_hdf5_dataset(hid_t group, const char* name, const H5L_info_t*,
}

const auto space = memory_space_type_from_string(entry.space);
const auto free_buffer = [space](char* ptr) { free_host_buffer(ptr, space); };
std::unique_ptr<char, decltype(free_buffer)> buffer(
const auto free_buffer = [space](std::byte* ptr) {
free_host_buffer(ptr, space);
};
std::unique_ptr<std::byte, decltype(free_buffer)> buffer(
allocate_host_buffer(entry.size, space), free_buffer);
const std::string dataset = std::string(name) + "/bytes";
CHECK_HDF5_CALL(
Expand All @@ -490,7 +494,7 @@ void read_functor_from_hdf5(hid_t file) {
std::size_t datatype_size;
CHECK_HDF5_CALL(H5LTget_dataset_info(file, "functor/functor", &dim, &datatype,
&datatype_size));
// We only deal with arrays of chars
// We only deal with arrays of bytes
assert(datatype == H5T_INTEGER);
assert(datatype_size == 1);

Expand Down Expand Up @@ -718,13 +722,13 @@ void read_policy_from_hdf5(hid_t file) {
}
}

std::vector<std::pair<char*, std::size_t>> compute_allocations(
const std::set<std::pair<char*, std::size_t>>& allocations) {
std::vector<std::pair<std::byte*, std::size_t>> compute_allocations(
const std::set<std::pair<std::byte*, std::size_t>>& allocations) {
if (allocations.empty()) {
return {};
}

std::vector<std::pair<char*, std::size_t>> allocs;
std::vector<std::pair<std::byte*, std::size_t>> allocs;
auto it = allocations.begin();
auto [current_address, current_size] = *(it++);
for (; it != allocations.end(); ++it) {
Expand Down Expand Up @@ -758,7 +762,7 @@ ScopeGuard::ScopeGuard(int& argc, char* argv[], bool enable_input_reset)
CHECK_HDF5_ID(H5Fopen(hdf5_filename.data(), H5F_ACC_RDONLY, H5P_DEFAULT)),
H5Fclose);

std::set<std::pair<char*, std::size_t>> host_allocation_locs;
std::set<std::pair<std::byte*, std::size_t>> host_allocation_locs;

impl::iterate_allocations(
file.get(), "in/views",
Expand All @@ -775,7 +779,7 @@ ScopeGuard::ScopeGuard(int& argc, char* argv[], bool enable_input_reset)
// we do it after doing the host allocations
device_init();

std::set<std::pair<char*, std::size_t>> device_allocation_locs;
std::set<std::pair<std::byte*, std::size_t>> device_allocation_locs;

impl::iterate_allocations(
file.get(), "in/views",
Expand All @@ -791,12 +795,12 @@ ScopeGuard::ScopeGuard(int& argc, char* argv[], bool enable_input_reset)
allocate(impl::MemorySpaceType::DEVICE, address, size);
}

std::map<std::pair<char*, std::string>, impl::StoredAllocation*>
std::map<std::pair<std::byte*, std::string>, impl::StoredAllocation*>
allocation_index;

impl::hdf5_iterate_fun_t copy_data_wrapper =
[this, &allocation_index](const impl::SnapshotAllocation& entry,
char* data) {
const std::byte* data) {
if (!entry.bytes_dumped) {
throw std::runtime_error("Missing input bytes for allocation '" +
entry.label + "'");
Expand All @@ -813,7 +817,7 @@ ScopeGuard::ScopeGuard(int& argc, char* argv[], bool enable_input_reset)
if (enable_input_reset_) {
input_snapshots_.push_back(
{space, entry.address,
std::vector<char>(data, data + entry.size)});
std::vector<std::byte>(data, data + entry.size)});
}
}
const auto allocation = replay_allocations_[space].emplace(
Expand Down Expand Up @@ -843,7 +847,7 @@ ScopeGuard::ScopeGuard(int& argc, char* argv[], bool enable_input_reset)

impl::hdf5_iterate_fun_t allocate_wrapper =
[this, &allocation_index](const impl::SnapshotAllocation& entry,
char* data) {
const std::byte* data) {
const auto it = allocation_index.find(
std::make_pair(entry.allocation_id, entry.space));
if (it == allocation_index.end()) {
Expand Down Expand Up @@ -902,8 +906,8 @@ std::optional<std::string> get_metadata(const std::string& key) {
}
}

void ScopeGuard::allocate(impl::MemorySpaceType memory_space, char* address,
std::size_t size) {
void ScopeGuard::allocate(impl::MemorySpaceType memory_space,
std::byte* address, std::size_t size) {
if (memory_space == impl::MemorySpaceType::HOST) {
host_raw_allocations.emplace_back(memory_space, address, size);
} else {
Expand All @@ -918,7 +922,7 @@ void ScopeGuard::allocate(impl::MemorySpaceType memory_space, char* address,

void ScopeGuard::allocate_output(impl::StoredAllocation& allocation,
const impl::SnapshotAllocation& snapshot,
char* data) {
const std::byte* data) {
if (allocation.output_seen || allocation.descriptor.label != snapshot.label) {
throw std::runtime_error("Duplicate or inconsistent output allocation '" +
snapshot.label + "'");
Expand Down
Loading
Loading