HIP
HIP copied to clipboard
ROCm
[hip_complex.h] Problem with how HIP registers overloads
Hello, the research group I work for is porting their software to HIP, and while we've been porting, we've ran up against an issue with HIP's operator overloading for complex types:
For example:
test_1.c:
#include <hip/hip_runtime.h>
#include <hip/hip_complex.h>
int main(int argc, char ** argv) {
hipDoubleComplex x = make_hipDoubleComplex(2.0, 3.0);
hipDoubleComplex y = make_hipDoubleComplex(4.0, 2.0);
y = y * x;
printf("%f%+fi\n", hipCreal(y), hipCimag(y));
return 0;
}
Compile and run with hipcc test_1.c -o test_1.o && ./test_1.o works as expected:
$ hipcc test_1.c -o test_1.o && ./test_1.o
clang-10: warning: treating 'c' input as 'c++' when in C++ mode, this behavior is deprecated [-Wdeprecated]
2.000000+16.000000i
However, using the *= operator for shorter syntax (and the ability for the compiler to optimize):
test_1.c:
#include <hip/hip_runtime.h>
#include <hip/hip_complex.h>
int main(int argc, char ** argv) {
hipDoubleComplex x = make_hipDoubleComplex(2.0, 3.0);
hipDoubleComplex y = make_hipDoubleComplex(4.0, 2.0);
//y = y * x;
y *= x;
printf("%f%+fi\n", hipCreal(y), hipCimag(y));
return 0;
}
And doing the same command:
$ hipcc test_1.c -o test_1.o && ./test_1.o
clang-10: warning: treating 'c' input as 'c++' when in C++ mode, this behavior is deprecated [-Wdeprecated]
test_1.c:13:7: error: use of overloaded operator '*=' is ambiguous (with operand types 'hipDoubleComplex' (aka 'HIP_vector_type<double, 2>') and 'hipDoubleComplex')
y *= x;
~ ^ ~
/opt/rocm/hip/include/hip/hcc_detail/hip_vector_types.h:344:26: note: candidate function
HIP_vector_type& operator*=(const HIP_vector_type& x) noexcept
^
/opt/rocm/hip/include/hip/hcc_detail/hip_complex.h:253:1: note: candidate function
COMPLEX_MUL_PREOP_OVERLOAD(hipDoubleComplex)
^
/opt/rocm/hip/include/hip/hcc_detail/hip_complex.h:105:45: note: expanded from macro 'COMPLEX_MUL_PREOP_OVERLOAD'
__device__ __host__ static inline type& operator*=(type& lhs, const type& rhs) { \
^
1 error generated.
test_1.c:13:7: error: use of overloaded operator '*=' is ambiguous (with operand types 'hipDoubleComplex' (aka 'HIP_vector_type<double, 2>') and 'hipDoubleComplex')
y *= x;
~ ^ ~
/opt/rocm/hip/include/hip/hcc_detail/hip_vector_types.h:344:26: note: candidate function
HIP_vector_type& operator*=(const HIP_vector_type& x) noexcept
^
/opt/rocm/hip/include/hip/hcc_detail/hip_complex.h:253:1: note: candidate function
COMPLEX_MUL_PREOP_OVERLOAD(hipDoubleComplex)
^
/opt/rocm/hip/include/hip/hcc_detail/hip_complex.h:105:45: note: expanded from macro 'COMPLEX_MUL_PREOP_OVERLOAD'
__device__ __host__ static inline type& operator*=(type& lhs, const type& rhs) { \
^
1 error generated.
Basically, the issue seems to be that hip*Complex types are defined as HIP_vector_types as well, which already have overloads for +=, *=, which conflicts with the hip*Complex types (and for some operations, like *=, they are different mathematical functions which will give different results!).
Obviously, we can just use y = y * x, but the *= would be nice to use. What can be done from here?
Have Raised PR#1585, Let me know if you still see any issue with hip complex number implementation
@SarbojitAMD There is another similar problem (although not with operator overloading).
Say I have the cuda code:
#include <cuComplex.h>
#include <stdio.h>
int main(int argc, char** argv) {
cuDoubleComplex rC[16]= {0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0};
printf("%f%+fi\n", cuCreal(rC[5]), cuCimag(rC[5]));
return 0;
}
This is valid, as cuComplex can be understood as an integer (being just the real part, and imag being 0 implicitly).
Therefore, the array rC contains 16 complex numbers with zeros.
If I am to run hipify (or just write the code in HIP):
#include <hip/hip_complex.h>
#include <stdio.h>
int main(int argc, char** argv) {
hipDoubleComplex rC[16]={0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0};
printf("%f%+fi\n", hipCreal(rC[5]), hipCimag(rC[5]));
return 0;
}
This gives an error during compilation (only the relevant part included for brevity):
test_complex.hip.cpp:7:58: error: no viable conversion from 'int' to 'hipDoubleComplex' (aka 'HIP_vector_type<double, 2>')
hipDoubleComplex rC[16]={0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0};
^
Is this intended? It would be very useful to be able to initialize them this way.
With a slight modification, this works:
#include <hip/hip_complex.h>
#include <stdio.h>
int main(int argc, char** argv) {
hipDoubleComplex rC[16]={{0,0},{0,0},{0,0},{0,0},{0,0},{0,0},{0,0},{0,0},{0,0},{0,0},{0,0},{0,0},{0,0},{0,0},{0,0},{0,0}};
printf("%f%+fi\n", hipCreal(rC[5]), hipCimag(rC[5]));
return 0;
}
Because each of the {0,0} gets mapped to the HIP_vector_type<double,2> constructor taking 2 numbers.
But, could you consider adding support for initialization implicitly using a single number?
This is still broken. It's been almost 2 years. In HPC applications involving quantum mechanics complex arithmetic is common and important. Typesafe overloads for complex types should be possible. There seems to have been a fix in #1585 but it was rejected it seems because of a conflicting design goal.
For CUDA we could write our own (operator overloads)[https://github.com/CompFUSE/DCA/blob/master/include/dca/linalg/util/complex_operators_cuda.cu.hpp] while using cuComplex types. However HIP provides broken complex operator overloads. For future searchers who use MAGMA just use magma(Float|Double)Complex and magma's overloads, pretend hip_complex.h doesn't exist.
Hi @cadebrown and @PDoakORNL, apologies for the lack of response. Please try latest ROCm 6.0.2 (HIP 6.0.32831) to see if your issue still exists? Thanks.
Closing the ticket as it is stale. Please re-open if the issue still exists with latest ROCm 6.0.2 (HIP 6.0.32831). Thanks.
Sorry @nartmada only stumbled across this today (and started encountering it yesterday) but issue is still there.
hipcc --version
HIP version: 6.0.32831-204d35d16
AMD clang version 17.0.0 (https://github.com/RadeonOpenCompute/llvm-project roc-6.0.2 24012 af27734ed982b52a9f1be0f035ac91726fc697e4)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /opt/rocm-6.0.2/llvm/bin
Configuration file: /opt/rocm-6.0.2/lib/llvm/bin/clang++.cfg
Objective
Have a header that can distinguish between CUDA, HIP and plain old C complex numbers and uses the C complex number syntax.
The HIP complex number section looks like
33 #elif defined __HIP_DEVICE_COMPILE__
34 #include <hip/hip_complex.h>
35 #warning "Compiling with hip_complex"
36 ///@brief Single precision complex number
37 typedef hipFloatComplex Complex_f;
38 ///@brief Double precision complex number
39 typedef hipDoubleComplex Complex;
40
41 ///@brief Extract Real Component
42 #define creal(z) z.x
43 ///@brief Extract Imaginary Component
44 #define cimag(z) z.y
45 ///@brief Complex Conjugation
46 ///@param z
47 __HOST_DEVICE__ static __inline__ Complex conj(Complex z){
48 return hipConj(z);
49 }
50 ///@brief Single precision Complex Conjugation
51 ///@param z
52 __HOST_DEVICE__ static __inline__ Complex_f conj(Complex_f z){
53 return hipConjf(z);
54 }
55
56 #ifndef _CEXP
57 #define _CEXP
58 ///@brief HIP does not contain a cexp routine so I've taken the one from GCC
59 ///and tweaked it slightly.
60 ///@param z
61 __HOST_DEVICE__ static __inline__ Complex cexp(Complex z)
62 {
63 double a = creal(z); double b = cimag(z);
64 Complex v;
65 v.x=cos(b);v.y=sin(b);
66 return exp (a) * v;
67 }
68 __HOST_DEVICE__ static __inline__ Complex_f cexp(Complex_f z)
69 {
70
71 float a = creal(z); float b = cimag(z);
72 Complex_f v;
73 v.x=cos(b);v.y=sin(b);
74 return exp (a) * v;
75 }
The issue
When I try to compile a routine with complex multiplication in it we get
hipcc -I../INCLUDE/ -I /opt/AMD/aocl/aocl-linux-aocc-4.1.0/aocc/include -I/usr/include -I /opt/rocm/include/hipblas -I /opt/rocm/include/hip -famd-opt -std=c++14 -err-no -w --offload-arch=gfx1034 -G -Xcompiler "-I../INCLUDE/ -I /opt/AMD/aocl/aocl-linux-aocc-4.1.0/aocc/include -I/usr/include -I /opt/rocm/include/hipblas -I /opt/rocm/include/hip -std=gnu11 -fcommon -D__HIPCC__ -march=native -fopenmp -fopenmp-simd -static-libgcc -DSA3AT -flto -D__HIP_PLATFORM_AMD__ -foffload=disable -Ofast -funroll-all-loops -fcf-protection=none " -c ../production_hip/../production_hip/hipbosonic.hip -o ../production_hip/hipbosonic.o
../production_hip/../production_hip/hipbosonic.hip:87:35: error: use of overloaded operator '*' is ambiguous (with operand types 'Complex_f' (aka 'HIP_vector_type<float, 2>') and 'Complex_f')
Complex_f Sigma11=u11t[i*ndim+mu]*u11t[uidm*ndim+nu]-u12t[i*ndim+mu]*conj(u12t[uidm*ndim+nu]);
~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~
/opt/rocm-6.0.2/include/hip/amd_detail/amd_hip_vector_types.h:548:65: note: candidate function
friend __HOST_DEVICE__ inline constexpr HIP_vector_type operator*(
^
/opt/rocm-6.0.2/include/hip/amd_detail/amd_hip_complex.h:239:1: note: candidate function
COMPLEX_MUL_OP_OVERLOAD(hipFloatComplex)
^
/opt/rocm-6.0.2/include/hip/amd_detail/amd_hip_complex.h:80:40: note: expanded from macro 'COMPLEX_MUL_OP_OVERLOAD'
__HOST_DEVICE__ static inline type operator*(const type& lhs, const type& rhs) {
