Skip to content

Improve <cuda/std/ctime> implementation #4682

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 1 commit into
base: main
Choose a base branch
from
Open

Improve <cuda/std/ctime> implementation #4682

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
7 changes: 7 additions & 0 deletions docs/libcudacxx/standard_api/c_library.rst
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -80,3 +80,10 @@ Any Standard C++ header not listed below is omitted.
- CCCL 3.0.0
- CUDA 13.0
- `\<cstring\> <https://en.cppreference.com/w/cpp/header/cstring>`_

* - ``<cuda/std/ctime>``
- Provides ``clock``, ``difftime``, ``time`` and ``timespec_get`` functions
- Defines ``CCCL_CLOCKS_PER_SEC`` alternative to ``CLOCKS_PER_SEC`` working in both host and device code
- CCCL 3.1.0
- CUDA 13.1
- `\<ctime\> <https://en.cppreference.com/w/cpp/header/ctime>`_
55 changes: 13 additions & 42 deletions libcudacxx/include/cuda/std/ctime
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,7 +24,7 @@
#include <cuda/__ptx/instructions/get_sreg.h>

#if !_CCCL_COMPILER(NVRTC)
# include <time.h>
# include <ctime>
#else
# define TIME_UTC 1

Expand All @@ -46,36 +46,25 @@ using ::size_t;
using ::time_t;
using ::timespec;

// glibc defines clock and other functions as macros, we need to prevent our symbols being redefined by them
// CCCL_CLOCKS_PER_SEC

[[nodiscard]] _CCCL_FORCEINLINE _CCCL_VISIBILITY_HIDDEN _CCCL_HOST_DEVICE clock_t __cccl_clocks_per_sec_impl() noexcept
{
NV_IF_ELSE_TARGET(NV_IS_HOST, (return CLOCKS_PER_SEC;), (return clock_t(1'000'000'000);))
}

#define CCCL_CLOCKS_PER_SEC _CUDA_VSTD::__cccl_clocks_per_sec_impl()
Comment on lines +51 to +56
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Alternatively, we could introduce cuda::clocks_per_sec() function, that would be probably a more suitable solution..

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I believe we want something that is valid both on host device if we ever need it

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yeah, that's what I meant, making this function public in cuda:: namespace, the implementation could stay the same


// clock

#if defined(clock)
# pragma push_macro("clock")
# undef clock
# define _LIBCUDACXX_POP_CLOCK_MACRO
#endif // clock
[[nodiscard]] _LIBCUDACXX_HIDE_FROM_ABI clock_t clock() noexcept
#ifdef _LIBCUDACXX_POP_CLOCK_MACRO
# pragma pop_macro("clock")
# undef _LIBCUDACXX_POP_CLOCK_MACRO
#endif // _LIBCUDACXX_POP_CLOCK_MACRO
{
NV_IF_ELSE_TARGET(NV_IS_HOST, (return ::clock();), (return static_cast<clock_t>(_CUDA_VPTX::get_sreg_clock64());))
NV_IF_ELSE_TARGET(NV_IS_HOST, (return ::std::clock();), (return static_cast<clock_t>(_CUDA_VPTX::get_sreg_clock64());))
}

// difftime

#if defined(difftime)
# pragma push_macro("difftime")
# undef difftime
# define _LIBCUDACXX_POP_DIFFTIME_MACRO
#endif // difftime
[[nodiscard]] _LIBCUDACXX_HIDE_FROM_ABI constexpr double difftime(time_t __end, time_t __start) noexcept
#ifdef _LIBCUDACXX_POP_DIFFTIME_MACRO
# pragma pop_macro("difftime")
# undef _LIBCUDACXX_POP_DIFFTIME_MACRO
#endif // _LIBCUDACXX_POP_DIFFTIME_MACRO
{
return static_cast<double>(__end - __start);
}
Expand All @@ -94,26 +83,17 @@ using ::timespec;
}
#endif // _CCCL_CUDA_COMPILATION()

#if defined(time)
# pragma push_macro("time")
# undef time
# define _LIBCUDACXX_POP_TIME_MACRO
#endif // time
_LIBCUDACXX_HIDE_FROM_ABI time_t time(time_t* __v) noexcept
#ifdef _LIBCUDACXX_POP_TIME_MACRO
# pragma pop_macro("time")
# undef _LIBCUDACXX_POP_TIME_MACRO
#endif // _LIBCUDACXX_POP_TIME_MACRO
{
NV_IF_ELSE_TARGET(NV_IS_HOST, (return ::time(__v);), (return _CUDA_VSTD::__cccl_time_impl_device(__v);))
NV_IF_ELSE_TARGET(NV_IS_HOST, (return ::std::time(__v);), (return _CUDA_VSTD::__cccl_time_impl_device(__v);))
}

// timespec_get

#if _CCCL_CUDA_COMPILATION()
[[nodiscard]] _CCCL_HIDE_FROM_ABI _CCCL_DEVICE int __cccl_timespec_get_impl_device(timespec* __ts, int __base) noexcept
{
if (__base != TIME_UTC)
if (__ts == nullptr || __base != TIME_UTC)
{
return 0;
}
Expand All @@ -124,19 +104,10 @@ _LIBCUDACXX_HIDE_FROM_ABI time_t time(time_t* __v) noexcept
}
#endif // _CCCL_CUDA_COMPILATION()

#if defined(timespec_get)
# pragma push_macro("timespec_get")
# undef timespec_get
# define _LIBCUDACXX_POP_TIMESPEC_GET_MACRO
#endif // timespec_get
[[nodiscard]] _LIBCUDACXX_HIDE_FROM_ABI int timespec_get(timespec* __ts, int __base) noexcept
#ifdef _LIBCUDACXX_POP_TIMESPEC_GET_MACRO
# pragma pop_macro("timespec_get")
# undef _LIBCUDACXX_POP_TIMESPEC_GET_MACRO
#endif // _LIBCUDACXX_POP_TIMESPEC_GET_MACRO
{
NV_IF_ELSE_TARGET(NV_IS_HOST,
(return ::timespec_get(__ts, __base);),
(return ::std::timespec_get(__ts, __base);),
(return _CUDA_VSTD::__cccl_timespec_get_impl_device(__ts, __base);))
}

Expand Down
34 changes: 16 additions & 18 deletions libcudacxx/test/libcudacxx/std/time/ctime/ctime.pass.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,32 +13,30 @@
#include <cuda/std/type_traits>
#include <cuda/std/utility>

// Undefine macros that conflict with the tested symbols

#if defined(clock)
# undef clock
#endif // clock

#if defined(difftime)
# undef difftime
#endif // difftime

#if defined(time)
# undef time
#endif // time

#if defined(timespec_get)
# undef timespec_get
#endif // timespec_get

#ifndef TIME_UTC
# error TIME_UTC not defined
#endif

#ifndef CCCL_CLOCKS_PER_SEC
# error CCCL_CLOCKS_PER_SEC not defined
#endif

static_assert(TIME_UTC != 0);

__host__ __device__ cuda::std::clock_t ref_clocks_per_sec()
{
NV_IF_ELSE_TARGET(NV_IS_HOST, (return CLOCKS_PER_SEC;), (return 1'000'000'000;));
}

__host__ __device__ bool test()
{
// CCCL_CLOCKS_PER_SEC

{
static_assert(cuda::std::is_same_v<cuda::std::clock_t, decltype(CCCL_CLOCKS_PER_SEC)>);
assert(CCCL_CLOCKS_PER_SEC == ref_clocks_per_sec());
}

// struct timespec

{
Expand Down
Loading