thrust icon indicating copy to clipboard operation
thrust copied to clipboard

thrust::complex `constexpr` constructors

Open Jacobfaib opened this issue 3 years ago • 2 comments

Is it possible to make the thrust::complex constructors constexpr? This would allow them to be used in constant-initialized __constant__ and/or __device__ variables. Currently such usage:

static __constant__ thrust::complex<double> One = 1.0;

barfs with

error: dynamic initialization is not supported for a __constant__ variable

swapping thrust::complex<double> for a constexpr equivalent:

template <typename T>
struct complex
{
  T r;
  T i;

  constexpr complex(T r_ = T{0}, T i_ = T{0}) : r(r_), i(i_) { }
};

works.

Jacobfaib avatar May 02 '22 17:05 Jacobfaib

Eventually thrust::complex will just be a wrapper around libcu++'s complex implementation, which does provide constexpr constructors. This will happen sometime after Thrust 2.0.

In the meantime, would switching to cuda::std::complex work for your usecase? https://www.godbolt.org/z/xWc1xeoPa

alliepiper avatar May 02 '22 17:05 alliepiper

would switching to cuda::std::complex work for your usecase?

It seems from the documentation that <cuda/std/complex> is a very recent addition, requiring at least CUDA 11.2. AFAIK we don't have a hard minimum requirement set -- but we check for compute capability as low as 30 so I think this would be a hard pill to swallow.

I suppose in the meantime I can just go ahead and use the work-around I listed above, as these instances are few and far between.


To put this in perspective, we don't "use" thrust::complex directly, rather it is the complex back-end type to a generic scalar type:

#if LIB_HAS_DEVICE
#  define libcomplex_backend thrust
#  include <thrust/complex.h>
#else
#  define libcomplex_backend std
#  include <complex>
#endif

typedef complex_backend::<lib_real_type> LibScalar 

Jacobfaib avatar May 02 '22 17:05 Jacobfaib

@Jacobfaib minimal supported CUDA version for Thrust is 11.0, which doesn't support sm30. To my understanding, your use case is not supported by thrust as well. Please, give libcu++ a try, there should be nothing architecture specific. If you have issues using complex from libcu++, feel free to open an issue there.

gevtushenko avatar Feb 23 '23 16:02 gevtushenko