Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
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
57 changes: 10 additions & 47 deletions c2h/include/c2h/catch2_test_helper.h
Original file line number Diff line number Diff line change
Expand Up @@ -6,7 +6,6 @@
#include <cuda/std/detail/__config>

#include <cuda/__memory_resource/legacy_pinned_memory_resource.h>
#include <cuda/__nvtx/nvtx.h>
#include <cuda/buffer>
#include <cuda/std/bit>
#include <cuda/std/cmath>
Expand All @@ -15,14 +14,19 @@
#include <cuda/std/utility>

#include <cstdint>
#include <cstdlib>
#include <iomanip>
#include <optional>
#include <sstream>
#include <string>
#include <tuple>
#include <type_traits>
#include <vector>

#include <c2h/catch2_main.h>
#include <c2h/catch2_test_macros.h>
#include <c2h/checked_allocator.cuh>
#include <c2h/detail/catch2_nvtx.h>
#include <c2h/detail/catch2_seed.h>
#include <c2h/device_policy.h>
#include <c2h/extended_types.h>
#include <c2h/test_util_vec.h>
Expand Down Expand Up @@ -598,22 +602,6 @@ struct Catch::StringMaker<cudaError>
#include <c2h/custom_type.h>
#include <c2h/generators.h>

namespace detail
{
struct nvtx_c2h_domain
{
static constexpr const char* name = "C2H";
};

template <typename T>
class nvtx_fixture
{
#if _CCCL_HAS_NVTX3()
::nvtx3::v1::scoped_range_in<nvtx_c2h_domain> nvtx_range{Catch::getResultCapture().getCurrentTestName()};
#endif // _CCCL_HAS_NVTX3()
};
} // namespace detail

#define C2H_TEST_NAME_IMPL(NAME, PARAM) C2H_TEST_STR(NAME) "(" C2H_TEST_STR(PARAM) ")"

#define C2H_TEST_NAME(NAME) C2H_TEST_NAME_IMPL(NAME, VAR_IDX)
Expand All @@ -623,7 +611,8 @@ class nvtx_fixture

#define C2H_TEST_IMPL(ID, NAME, TAG, ...) \
using C2H_TEST_CONCAT(types_, ID) = c2h::cartesian_product<__VA_ARGS__>; \
CATCH_TEMPLATE_LIST_TEST_CASE_METHOD(::detail::nvtx_fixture, C2H_TEST_NAME(NAME), TAG, C2H_TEST_CONCAT(types_, ID))
CATCH_TEMPLATE_LIST_TEST_CASE_METHOD( \
::c2h::detail::nvtx_fixture, C2H_TEST_NAME(NAME), TAG, C2H_TEST_CONCAT(types_, ID))

#define C2H_TEST(NAME, TAG, ...) C2H_TEST_IMPL(__LINE__, NAME, TAG, __VA_ARGS__)

Expand All @@ -636,7 +625,8 @@ class nvtx_fixture

#define C2H_TEST_LIST_IMPL(ID, NAME, TAG, ...) \
using C2H_TEST_CONCAT(types_, ID) = c2h::type_list<__VA_ARGS__>; \
CATCH_TEMPLATE_LIST_TEST_CASE_METHOD(::detail::nvtx_fixture, C2H_TEST_NAME(NAME), TAG, C2H_TEST_CONCAT(types_, ID))
CATCH_TEMPLATE_LIST_TEST_CASE_METHOD( \
::c2h::detail::nvtx_fixture, C2H_TEST_NAME(NAME), TAG, C2H_TEST_CONCAT(types_, ID))

#define C2H_TEST_LIST(NAME, TAG, ...) C2H_TEST_LIST_IMPL(__LINE__, NAME, TAG, __VA_ARGS__)

Expand All @@ -648,30 +638,3 @@ class nvtx_fixture
C2H_TEST_LIST_WITH_FIXTURE_IMPL(__LINE__, FIXTURE, NAME, TAG, __VA_ARGS__)

#define C2H_TEST_STR(a) #a

namespace c2h
{
inline std::size_t get_override_seed_count()
{
// Setting this environment variable forces a fixed number of seeds to be generated, regardless of the requested
// count. Set to 1 to reduce redundant, expensive testing when using sanitizers, etc.
static std::optional<std::string> override_str = c2h::detail::get_env("C2H_SEED_COUNT_OVERRIDE");
static const int override_seeds =
override_str ? static_cast<int>(std::strtol(override_str->c_str(), nullptr, 10)) : 0;
return override_seeds;
}

inline std::size_t adjust_seed_count(std::size_t requested)
{
static const std::size_t override_seeds = get_override_seed_count();
return override_seeds != 0 ? override_seeds : requested;
}
} // namespace c2h

#define C2H_SEED(N) \
c2h::seed_t \
{ \
GENERATE_COPY(take(c2h::adjust_seed_count(N), \
random(::cuda::std::numeric_limits<unsigned long long int>::min(), \
::cuda::std::numeric_limits<unsigned long long int>::max()))) \
}
131 changes: 3 additions & 128 deletions c2h/include/c2h/checked_allocator.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -9,140 +9,15 @@
#include <thrust/system/cuda/memory_resource.h>
#include <thrust/system/cuda/pointer.h>

#include <cstdlib>
#include <iostream>
#include <cstddef>
#include <new>
#include <optional>
#include <string>

#include <cuda_runtime_api.h>

namespace c2h
{
namespace detail
{
inline std::optional<std::string> get_env(const char* name)
{
#ifdef _WIN32
char* buf = nullptr;
std::size_t len = 0;
if (_dupenv_s(&buf, &len, name) || !buf)
{
return std::nullopt;
}
std::string val(buf);
free(buf);
return val;
#else
if (const char* v = std::getenv(name))
{
return std::string(v);
}
return std::nullopt;
#endif
}

struct memory_info
{
std::size_t free{};
std::size_t total{};
bool override{false};
};

// If the environment variable C2H_DEVICE_MEMORY_LIMIT is set, the total device memory
// will be limited to this number of bytes.
inline std::size_t get_device_memory_limit()
{
static std::optional<std::string> override_str = get_env("C2H_DEVICE_MEMORY_LIMIT");
static const std::size_t result =
override_str ? static_cast<std::size_t>(std::strtoll(override_str->c_str(), nullptr, 10)) : 0;
return result;
}

inline bool get_debug_checked_allocs()
{
static std::optional<std::string> debug_checked_allocs = get_env("C2H_DEBUG_CHECKED_ALLOC_FAILURES");
static const bool result = debug_checked_allocs && (std::strtol(debug_checked_allocs->c_str(), nullptr, 10) != 0);
return result;
}

inline cudaError_t get_device_memory(memory_info& info)
{
static const std::size_t device_memory_limit = get_device_memory_limit();

const cudaError_t status = cudaMemGetInfo(&info.free, &info.total);
if (status != cudaSuccess)
{
return status;
}

if (device_memory_limit > 0)
{
info.free = (std::max) (std::size_t{0}, static_cast<std::size_t>(info.free - (info.total - device_memory_limit)));
info.total = device_memory_limit;
info.override = true;
}
#include <c2h/detail/checked_memory.cuh>

return cudaSuccess;
}

inline cudaError_t check_free_device_memory(std::size_t bytes)
{
memory_info info;
const cudaError_t status = get_device_memory(info);
if (status != cudaSuccess)
{
return status;
}

// Avoid allocating all available memory:
constexpr std::size_t padding = 16 * 1024 * 1024; // 16 MiB
if (info.free < (bytes + padding))
{
if (get_debug_checked_allocs())
{
const double total_GiB = static_cast<double>(info.total) / (1024 * 1024 * 1024);
const double free_GiB = static_cast<double>(info.free) / (1024 * 1024 * 1024);
const double requested_GiB = static_cast<double>(bytes) / (1024 * 1024 * 1024);
const double padded_GiB = static_cast<double>(bytes + padding) / (1024 * 1024 * 1024);

std::cerr << "Device memory allocation failed due to insufficient free device memory.\n";

if (info.override)
{
std::cerr
<< "Available device memory has been limited (env var C2H_DEVICE_MEMORY_LIMIT=" << get_device_memory_limit()
<< ").\n";
}

std::cerr
<< "Total device mem: " << total_GiB << " GiB\n" //
<< "Free device mem: " << free_GiB << " GiB\n" //
<< "Requested device mem: " << requested_GiB << " GiB\n" //
<< "Padded device mem: " << padded_GiB << " GiB\n";
}

return cudaErrorMemoryAllocation;
}

return cudaSuccess;
}

// Check available memory prior to calling cudaMalloc.
// This avoids hangups and slowdowns from allocating swap / non-device memory
// on some platforms, namely tegra.
inline cudaError_t checked_cuda_malloc(void** ptr, std::size_t bytes)
namespace c2h
{
auto status = check_free_device_memory(bytes);
if (status != cudaSuccess)
{
return status;
}

return cudaMalloc(ptr, bytes);
}
} // namespace detail

using checked_cuda_memory_resource = THRUST_NS_QUALIFIER::system::cuda::detail::
cuda_memory_resource<detail::checked_cuda_malloc, cudaFree, THRUST_NS_QUALIFIER::cuda::pointer<void>>;

Expand Down
26 changes: 26 additions & 0 deletions c2h/include/c2h/detail/catch2_nvtx.h
Original file line number Diff line number Diff line change
@@ -0,0 +1,26 @@
// SPDX-FileCopyrightText: Copyright (c) 2011-2026, NVIDIA CORPORATION. All rights reserved.
// SPDX-License-Identifier: BSD-3-Clause

#pragma once

#include <cuda/std/detail/__config>

#include <cuda/__nvtx/nvtx.h>

#include <catch2/interfaces/catch_interfaces_capture.hpp>

namespace c2h::detail
{
struct nvtx_c2h_domain
{
static constexpr const char* name = "C2H";
};

template <typename T>
class nvtx_fixture
{
#if _CCCL_HAS_NVTX3()
::nvtx3::v1::scoped_range_in<nvtx_c2h_domain> nvtx_range{Catch::getResultCapture().getCurrentTestName()};
#endif // _CCCL_HAS_NVTX3()
};
} // namespace c2h::detail
42 changes: 42 additions & 0 deletions c2h/include/c2h/detail/catch2_seed.h
Original file line number Diff line number Diff line change
@@ -0,0 +1,42 @@
// SPDX-FileCopyrightText: Copyright (c) 2011-2026, NVIDIA CORPORATION. All rights reserved.
// SPDX-License-Identifier: BSD-3-Clause

#pragma once

#include <cuda/std/limits>

#include <cstddef>
#include <cstdlib>
#include <optional>
#include <string>

#include <c2h/detail/env.cuh>
#include <c2h/generator_common.h>
#include <catch2/generators/catch_generators_all.hpp>

namespace c2h
{
inline std::size_t get_override_seed_count()
{
// Setting this environment variable forces a fixed number of seeds to be generated, regardless of the requested
// count. Set to 1 to reduce redundant, expensive testing when using sanitizers, etc.
static const std::optional<std::string> override_str = c2h::detail::get_env("C2H_SEED_COUNT_OVERRIDE");
static const int override_seeds =
override_str ? static_cast<int>(std::strtol(override_str->c_str(), nullptr, 10)) : 0;
return override_seeds;
}

inline std::size_t adjust_seed_count(std::size_t requested)
{
static const std::size_t override_seeds = get_override_seed_count();
return override_seeds != 0 ? override_seeds : requested;
}
} // namespace c2h

#define C2H_SEED(N) \
c2h::seed_t \
{ \
GENERATE_COPY(take(c2h::adjust_seed_count(N), \
random(::cuda::std::numeric_limits<unsigned long long int>::min(), \
::cuda::std::numeric_limits<unsigned long long int>::max()))) \
}
Loading
Loading