Replace unsigned char/char by std::byte wherever possible

src/krepe/capture/kokkos_hooks.cpp

@@ -21,6 +21,7 @@
 #include <algorithm>
 #include <atomic>
 #include <charconv>
+#include <cstddef>
 #include <cstdint>
 #include <cstring>
 #include <optional>
@@ -74,8 +75,8 @@ std::unordered_map<std::string, std::uint64_t> kernel_invocations;
 Kokkos_Tools_toolInvokedFenceFunction tool_fence = nullptr;
 
 krepe::AllocationTracker allocation_tracker;
-std::vector<unsigned char> functor_data;
-std::vector<unsigned char> nvcc_inner_lambda_data;
+std::vector<std::byte> functor_data;
+std::vector<std::byte> nvcc_inner_lambda_data;
 std::unordered_map<std::string, std::string> metadata;
 std::variant<krepe::NoPolicyDesc, krepe::ScalarPolicyDesc,
              krepe::RangePolicyDesc, krepe::MDRangePolicyDesc,
@@ -202,7 +203,7 @@ bool should_track_allocation(const char* label, const void* ptr) {
 }
 
 const void* allocation_data_pointer(const void* ptr) {
-  return static_cast<const unsigned char*>(ptr) +
+  return static_cast<const std::byte*>(ptr) +
          KREPE_KOKKOS_ALLOCATION_HEADER_SIZE;
 }
 
@@ -369,16 +370,16 @@ KOKKOS_HOOKS_EXPORT void krepe_kernel_dump_add_metadata(const char* key,
   metadata[key] = std::string(value, size);
 }
 
-KOKKOS_HOOKS_EXPORT void krepe_kernel_dump_copy_functor(
-    const unsigned char* data, std::size_t size) {
+KOKKOS_HOOKS_EXPORT void krepe_kernel_dump_copy_functor(const std::byte* data,
+                                                        std::size_t size) {
   functor_data.resize(size);
   std::memcpy(functor_data.data(), data, size);
   nvcc_inner_lambda_data.clear();
 }
 
 KOKKOS_HOOKS_EXPORT void krepe_kernel_dump_copy_nvcc_lambda(
-    const unsigned char* data, std::size_t size,
-    const unsigned char* inner_functor_data, std::size_t inner_functor_size) {
+    const std::byte* data, std::size_t size,
+    const std::byte* inner_functor_data, std::size_t inner_functor_size) {
   functor_data.resize(size);
   std::memcpy(functor_data.data(), data, size);
   nvcc_inner_lambda_data.resize(inner_functor_size);

src/krepe/capture/memory_copy.cpp

@@ -3,6 +3,7 @@
 
 #include "memory_copy.hpp"
 
+#include <cstddef>
 #include <cstring>
 #include <exception>
 #include <limits>
@@ -32,7 +33,7 @@ bool is_cuda_device_space(const std::string& space) { return space == "Cuda"; }
 bool is_hip_device_space(const std::string& space) { return space == "HIP"; }
 
 std::string allocate_staging_buffer(const ActiveAllocation& allocation,
-                                    std::vector<unsigned char>& bytes) {
+                                    std::vector<std::byte>& bytes) {
   try {
     bytes.resize(static_cast<std::size_t>(allocation.record.size));
   } catch (const std::exception& error) {
@@ -46,7 +47,7 @@ std::string allocate_staging_buffer(const ActiveAllocation& allocation,
 }  // namespace
 
 std::string copy_allocation_bytes(const ActiveAllocation& allocation,
-                                  std::vector<unsigned char>& bytes) {
+                                  std::vector<std::byte>& bytes) {
   if (!allocation.record.data_size_known) {
     return "allocation data size could not be bounded safely";
   }

src/krepe/capture/memory_copy.hpp

@@ -5,12 +5,13 @@
 
 #include "allocation_tracker.hpp"
 
+#include <cstddef>
 #include <string>
 #include <vector>
 
 namespace krepe {
 
 std::string copy_allocation_bytes(const ActiveAllocation& allocation,
-                                  std::vector<unsigned char>& bytes);
+                                  std::vector<std::byte>& bytes);
 
 }  // namespace krepe

src/krepe/capture/view_dump.cpp

@@ -105,7 +105,7 @@ void write_int_attribute(hid_t object, const char* name, int value) {
 
 template <class T>
 auto Hdf5Datatype() {
-  if constexpr (std::is_same_v<T, unsigned char>) {
+  if constexpr (std::is_same_v<T, std::byte>) {
     return H5T_NATIVE_UCHAR;
   } else if constexpr (std::is_same_v<T, std::int64_t>) {
     return H5T_NATIVE_INT64;
@@ -211,7 +211,7 @@ void write_allocation_group(hid_t views_group,
   write_uint64_attribute(group.get(), "reported_size",
                          allocation.record.reported_size);
 
-  std::vector<unsigned char> bytes;
+  std::vector<std::byte> bytes;
   const std::string skip_reason = copy_allocation_bytes(allocation, bytes);
   write_int_attribute(group.get(), "bytes_dumped", skip_reason.empty() ? 1 : 0);
   if (!skip_reason.empty()) {
@@ -250,8 +250,8 @@ void write_snapshot(hid_t file, const char* name,
 
 ViewDumpResult create_kernel_dump(
     const AllocationSnapshot& input_snapshot,
-    const std::vector<unsigned char>& functor_data,
-    const std::vector<unsigned char>& nvcc_inner_lambda_data,
+    const std::vector<std::byte>& functor_data,
+    const std::vector<std::byte>& nvcc_inner_lambda_data,
     const std::unordered_map<std::string, std::string>& metadata,
     const std::variant<krepe::NoPolicyDesc, krepe::ScalarPolicyDesc,
                        krepe::RangePolicyDesc, krepe::MDRangePolicyDesc,

src/krepe/capture/view_dump.hpp

@@ -5,6 +5,7 @@
 
 #include "allocation_tracker.hpp"
 
+#include <cstddef>
 #include <cstdint>
 #include <string>
 #include <string_view>
@@ -68,8 +69,8 @@ struct TeamPolicyDesc {
 
 ViewDumpResult create_kernel_dump(
     const AllocationSnapshot& input_snapshot,
-    const std::vector<unsigned char>& functor_data,
-    const std::vector<unsigned char>& nvcc_inner_lambda_data,
+    const std::vector<std::byte>& functor_data,
+    const std::vector<std::byte>& nvcc_inner_lambda_data,
     const std::unordered_map<std::string, std::string>& metadata,
     const std::variant<krepe::NoPolicyDesc, krepe::ScalarPolicyDesc,
                        krepe::RangePolicyDesc, krepe::MDRangePolicyDesc,

src/krepe/common/extended_lambda_utils.hpp

@@ -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*));
 }
 

src/krepe/extract/kernel_extractor.cpp

@@ -2,6 +2,7 @@
 // SPDX-FileCopyrightText: Copyright Contributors to the KREPE project
 
 #include "kernel_extractor.hpp"
+#include <cstddef>
 #include <iostream>
 #if defined(KOKKOS_ENABLE_LIBDL)
 #include <dlfcn.h>
@@ -13,9 +14,9 @@ namespace krepe {
 namespace impl {
 
 void (*add_metadata_function)(const char*, const char*, std::size_t) = nullptr;
-void (*copy_functor_function)(const unsigned char*, std::size_t)     = nullptr;
-void (*copy_nvcc_lambda_function)(const unsigned char*, std::size_t,
-                                  const unsigned char*, std::size_t) = nullptr;
+void (*copy_functor_function)(const std::byte*, std::size_t)         = nullptr;
+void (*copy_nvcc_lambda_function)(const std::byte*, std::size_t,
+                                  const std::byte*, std::size_t)     = nullptr;
 void (*register_scalar_policy_function)(std::uint64_t)               = nullptr;
 void (*register_range_policy_function)(const char*, const char*, std::size_t,
                                        bool, std::uint64_t, std::uint64_t,
@@ -81,9 +82,9 @@ void init_internal_functions() {
       has_kernel_dump_tool = true;
     } else {
       add_metadata_function     = [](const char*, const char*, std::size_t) {};
-      copy_functor_function     = [](const unsigned char*, std::size_t) {};
-      copy_nvcc_lambda_function = [](const unsigned char*, std::size_t,
-                                     const unsigned char*, std::size_t) {};
+      copy_functor_function     = [](const std::byte*, std::size_t) {};
+      copy_nvcc_lambda_function = [](const std::byte*, std::size_t,
+                                     const std::byte*, std::size_t) {};
       register_scalar_policy_function = [](std::uint64_t) {};
       register_range_policy_function = [](const char*, const char*, std::size_t,
                                           bool, std::uint64_t, std::uint64_t,
@@ -103,13 +104,13 @@ void init_internal_functions() {
   }
 }
 
-void copy_functor(const unsigned char* functor_data, std::size_t functor_size) {
+void copy_functor(const std::byte* functor_data, std::size_t functor_size) {
   init_internal_functions();
   copy_functor_function(functor_data, functor_size);
 }
 
-void copy_functor(const unsigned char* functor_data, std::size_t functor_size,
-                  const unsigned char* inner_functor_data,
+void copy_functor(const std::byte* functor_data, std::size_t functor_size,
+                  const std::byte* inner_functor_data,
                   std::size_t inner_functor_size) {
   init_internal_functions();
   copy_nvcc_lambda_function(functor_data, functor_size, inner_functor_data,

src/krepe/extract/kernel_extractor.hpp

@@ -16,9 +16,9 @@ struct scratch_description {
   int level1;
 };
 
-void copy_functor(const unsigned char* functor_data, std::size_t functor_size);
-void copy_functor(const unsigned char* functor_data, std::size_t functor_size,
-                  const unsigned char* inner_functor_data,
+void copy_functor(const std::byte* functor_data, std::size_t functor_size);
+void copy_functor(const std::byte* functor_data, std::size_t functor_size,
+                  const std::byte* inner_functor_data,
                   std::size_t inner_functor_size);
 void register_scalar_policy(std::uint64_t N);
 void register_range_policy(const char* space, const char* schedule,
@@ -171,15 +171,15 @@ template <class Functor>
 Functor replay_functor(Functor&& functor) {
 #if defined(KERNEL_REPLAYER_USE_NVCC_HDL_WORKAROUND)
   if constexpr (krepe::hdl_utils::lambda_is_hdl<Functor>()) {
-    impl::copy_functor(reinterpret_cast<const unsigned char*>(&functor),
+    impl::copy_functor(reinterpret_cast<const std::byte*>(&functor),
                        sizeof(functor),
-                       reinterpret_cast<const unsigned char*>(
+                       static_cast<const std::byte*>(
                            krepe::hdl_utils::hdl_host_lambda_pointer(functor)),
                        krepe::hdl_utils::hdl_host_lambda_size(functor));
   } else
 #endif
   {
-    impl::copy_functor(reinterpret_cast<const unsigned char*>(&functor),
+    impl::copy_functor(reinterpret_cast<const std::byte*>(&functor),
                        sizeof(Functor));
   }
   return functor;

src/krepe/replay/allocation.hpp

@@ -3,6 +3,7 @@
 
 #pragma once
 
+#include <cstddef>
 #include <string>
 #include <tuple>
 
@@ -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) {
@@ -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) *
       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 {
@@ -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) {

src/krepe/replay/cuda_allocation.hpp

@@ -3,6 +3,7 @@
 
 #pragma once
 
+#include <cstddef>
 #include <cstring>
 #include <iostream>
 #include <sstream>
@@ -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(
@@ -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));
 }

src/krepe/replay/hip_allocation.hpp

@@ -3,6 +3,7 @@
 
 #pragma once
 
+#include <cstddef>
 #include <cstring>
 #include <tuple>
 #include <utility>
@@ -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));
 }
 

src/krepe/replay/host_allocation.hpp

@@ -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);
 }