This isn't atomic? You need to return csum directly, otherwise the value at addr_as_ull may change underneath you.

In fact this is also wrong as atomic read-modify-write ops return the old value, not the new value. So this should be bit-casting assumed.

In fact this is also wrong as atomic read-modify-write ops return the old value, not the new value. So this should be bit-casting assumed.

Sorry for the oversight. Could you help me understand? I know that atomicCAS returns the old value but with what in mind are you referring that to?

I understand that addr_as_ull sholdn't be returned, as as another thread can change it correct? Why are we to use assumed though and not csum?

assumed is the value before performing the update, which is what is returned by normal atomicAdd, atomicMax, etc.

See the CAS implementation for half as an example:

Oh you're right. I got tunnel visioned on the line on the actual call of atomicCAS, yes it should be assumed. i forgot that I am actually implementing an atomic operation here and that it should follow suit lol

The build failure is totally not your fault as it can't be seen from external CI, we only see it when internal workflows run after the PR is merged

Looking at the failure and pattern matching a bit, it looks like maybe we need __host__ __device__ here as well as for complex_max on 433.

Does this change make sense? I can import the PR and see whether this fixes the internal build tomorrow morning

Unfortunately since complex_min and complex_max both use CUDA intrinsics, they won't compile if you make it a __host__ function as well. The use of __fsqrt_rn for example should lead to more performant code/better CUDA assembly. CUDA intrinsics should be taken advantage of imo because it's kernel code and complex numbers are heavier computations to make in general.

The easiest solution would be to add an overload for complex when compiled with CUDA to have an operator*= available with __host__ __device__.

Just adding this

#if defined(__CUDACC__) || defined(__HIPCC__)
  template <typename U>
  C10_HOST_DEVICE constexpr complex<T>& operator*=(const complex<U>& rhs) {
    // (a + bi) * (c + di) = (a*c - b*d) + (a * d + b * c) i
    T a = real_;
    T b = imag_;
    U c = rhs.real();
    U d = rhs.imag();
    real_ = a * c - b * d;
    imag_ = a * d + b * c;
    return *this;
  }
#endif

in complex.h should fix this problem.

Also on second look I made an oversight in the complex_max and complex_min functions. They should be using regular comparisons and not std::max given it's a __device__ function. So on that note, it's actually good that this PR got reverted! I will push those changes and things should build on your end.

Is there any precedence for this definition of complex max/min in PyTorch?

I am not experienced enough with Pytorch to answer that. Aside from using magnitudes how else would you order them? I followed convention from other ecosystems and from my research this is how it is done across different disciplines/domains, is it not?

Exactly, they cannot be ordered (in mathematical terms, complex numbers are not an ordered field)
We should error in these cases, same as we error when we call max on a complex tensor. If people want to use these ops on complex tensors, they can do a view_as_real and perform some transformations on the output to define the order they want.

I was under the impression that some contexts will use magnitude for ordering complex numbers, like spectral analysis for DSP. I also took motivation from https://www.mathworks.com/help/matlab/ref/max.html as well

@Franklalalala

Could you comment on whether or not you had a use case for scattering complex numbers? What kind of work were you trying to do? Would you know if ordering of complex numbers is practically useful?

It may be implemented and it may be useful, but we don't implement that in PyTorch at a kernel level. As mentioned above, all these orderings can often be simulated with the current API and a bit of imagination :)

I see - in that case I'll wait until @mikaylagawarecki has a chance to review again. Thanks for taking a look!

@ZelboK I am working on Tensor_network, which require a series of matrix multiplication. In the case of complex elements, the torch scatter connot be used in GPU. As far as I concerned right now, we do not use sortage here, just elements' multiplication.
By the way, we have worked a way out, that is, we transform the complex number through Euler transformation and turns the multiplication to addition of angles and multiplication of magnitude.
The excellent work of you guys has reached out of my knowledge base, I connot give anymore advices. But thanks!

@ZelboK I am working on Tensor_network, which require a series of matrix multiplication. In the case of complex elements, the torch scatter connot be used in GPU. As far as I concerned right now, we do not use sortage here, just elements' multiplication. By the way, we have worked a way out, that is, we transform the complex number through Euler transformation and turns the multiplication to addition of angles and multiplication of magnitude. The excellent work of you guys has reached out of my knowledge base, I connot give anymore advices. But thanks!

Thanks a lot for responding, I was genuinely curious. This helps give me perspective :)