Skip to content
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.

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

Replace C++ 'std' methods with Axom core methods in host-device functions. #1195

Open
rhornung67 opened this issue Oct 4, 2023 · 0 comments · May be fixed by #1196
Open

Replace C++ 'std' methods with Axom core methods in host-device functions. #1195

rhornung67 opened this issue Oct 4, 2023 · 0 comments · May be fixed by #1196
Assignees
Labels
Core Issues related to Axom's 'core' component design Issues related to design and software engineering Reviewed

Comments

@rhornung67
Copy link
Member

rhornung67 commented Oct 4, 2023

For example, users see this warning:

/usr/gapps/bdiv/blueos_3_ppc64le_ib_p9/clang-14-cuda-11.8-gcc831/axom/0.8.1-cuda/include/axom/spin/policy/LinearBVH.hpp(82): warning #20013-D: calling a constexpr __host__ function("max") from a __host__ __device__ function("operator()") is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this.

The code generating the warning is: https://github.com/LLNL/axom/blob/develop/src/axom/spin/policy/LinearBVH.hpp#L82

We should replace use of 'std' utility methods in our host-device methods with methods we support in core that are appropriate for use in host-device code, for example something like this: https://github.com/LLNL/axom/blob/develop/src/axom/core/utilities/Utilities.hpp#L79

Here is a reference describing how std::numeric_limits is supported in CUDA: https://saturncloud.io/blog/is-there-a-cuda-equivalent-to-stdnumericlimits/

For HIP, it looks like no extra namespace is needed and the compiler does the right thing; i.e., std::numeric_limits::max() will work in host and device code. For example, the following works

`#include <hip/hip_runtime.h>

void global
kernel(int *a, size_t *b)
{
a[0] = std::numeric_limits::max();
b[0] = std::numeric_limits<size_t>::max();
}

int main(void)
{
int * a;
(void) hipMalloc(&a,sizeof(int));
(void) hipMemset(a,0,sizeof(int));
size_t * b;
(void) hipMalloc(&b,sizeof(size_t));
(void) hipMemset(b,0,sizeof(size_t));

kernel<<<1,1,0,0>>>(a,b);

int ha;
size_t hb;
(void) hipMemcpy(&ha,a,sizeof(int),hipMemcpyDeviceToHost);
(void) hipMemcpy(&hb,b,sizeof(size_t),hipMemcpyDeviceToHost);

printf("a = %i, b = %zu\n",ha,hb);

assert(ha == std::numeric_limits::max());
assert(hb == std::numeric_limits<size_t>::max());

}`

@rhornung67 rhornung67 added Core Issues related to Axom's 'core' component design Issues related to design and software engineering labels Oct 4, 2023
@rhornung67 rhornung67 linked a pull request Oct 5, 2023 that will close this issue
@rhornung67 rhornung67 self-assigned this Oct 23, 2023
@rhornung67 rhornung67 added this to the April 2024 Release milestone Oct 23, 2023
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
Core Issues related to Axom's 'core' component design Issues related to design and software engineering Reviewed
Projects
None yet
Development

Successfully merging a pull request may close this issue.

2 participants