Not possible to use some math functions with the CUDA and HIP backends #5326
Comments
alexbatashev
added
cuda
|
could you call sycl::atan2 and sycl::hypot ? |
|
I don't believe that The |
|
Having said that, #include <CL/sycl.hpp>
#include <cmath>
#include <iostream>
int main() {
sycl::queue queue;
std::cout << "Running on device: "
<< queue.get_device().get_info<sycl::info::device::name>()
<< std::endl;
static constexpr int ARRAY_SIZES = 100;
sycl::buffer<float> a(ARRAY_SIZES), b(ARRAY_SIZES), c(ARRAY_SIZES);
{
auto acc_a = a.get_access<sycl::access::mode::write>();
auto acc_b = b.get_access<sycl::access::mode::write>();
for (int i = 0; i < ARRAY_SIZES; ++i) {
acc_a[i] = 1.23f;
acc_b[i] = 2.34f;
}
}
queue.submit([&](::sycl::handler& h) {
auto acc_a = a.get_access<sycl::access::mode::read>(h);
auto acc_b = b.get_access<sycl::access::mode::read>(h);
auto acc_c = c.get_access<sycl::access::mode::write>(h);
h.parallel_for<class ATan2Test>(sycl::range<1>(ARRAY_SIZES),
[=](sycl::item<1> i) {
acc_c[i] = sycl::hypot(acc_a[i], acc_b[i]);
});
});
{
static const float RESULT = sycl::hypot(1.23f, 2.34f);
auto acc_c = c.get_access<sycl::access::mode::read>();
for (int i = 0; i < ARRAY_SIZES; ++i) {
if (std::abs(acc_c[i] - RESULT) > 0.001) {
std::cerr << "Result at index " << i << " is " << acc_c[i] << " instead of "
<< RESULT << std::endl;
return 1;
}
}
}
std::cout << "All OK!" << std::endl;
return 0;
}And: So while being able to use |
|
🤦 I need to drink some coffee... You're right, So we can at least hack around the issue for now. But being able to use std::atan2 and std::hypot in our code instead of |
|
C++ standard functions are not supported by SYCL standard. Instead SYCL provides a set of "built-in" functions in There is an extension to support C++ standard functions, which is currently supported on Intel devices only. So, I think the right classification for that issue is "enhancement" - enable C++ standard library extension on HIP/CUDA backends. |
|
Hi Alexey, I have to say I was always confused by the trigonometric functions in the In our current development we are experimenting with declaring certain calculations in "accelerator agnostic" code in a way that they could be used either on the host or on a device. Being able to use "standard" trigonometric functions in such code is a big simplification. Since right now I'm working around this issue in our own code like: // SYCL include(s).
#if defined(CL_SYCL_LANGUAGE_VERSION) || defined(SYCL_LANGUAGE_VERSION)
#include <CL/sycl.hpp>
#endif
...
TRACCC_HOST_DEVICE
scalar phi() const {
#if defined(CL_SYCL_LANGUAGE_VERSION) || defined(SYCL_LANGUAGE_VERSION)
return cl::sycl::atan2(m_y, m_x);
#else
return std::atan2(m_y, m_x);
#endif // SYCL
}But what's worse is that even after fixing all of these function calls in this particular project, it turns out that we are still pulling in at least one Long story short, future SYCL standards should really officially support using the math functions from the C++ standard library in device code. Much like how CUDA supports the C Cheers, |
AerialMantis
added
compiler
|
Thanks for the update! |
|
No prob. I am advocating for us to add support for CXX stdlib for HIP backend as well, so hopefully this is something that happens in the next few months or so. I will update this ticket if/when this happens |
This is now implemented by #15055 Hopefully all of std::math now works on AMD too. I think that a lot of it is tested in the tests that were enabed by #15055 |
Describe the bug
This may have been reported already, I just couldn't easily find an existing ticket about it... When I try to use the atan2 and std::hypot functions in a kernel, I'm not able to build that kernel into "CUDA or HIP binaries".
To Reproduce
Take the following trivial example:
On Intel backends it seems to work fine.
But if I try to compile this code with the CUDA or HIP backends, those just really don't want to play ball... 😦
Environment (please complete the following information):
Additional context
I guess the answer may very well be that these two functions (atan2 and std::hypot) are just things that have not been implemented for these backends yet. Which is fair. In that case this would be a feature request to implement them. 😄
Pinging @ivorobts, @konradkusiak97 and @beomki-yeo.
The text was updated successfully, but these errors were encountered: