- Fix thrust::complex<T> to always be aligned to sizeof(T) * 2.

- Add unit tests for `thrust::complex<T>` alignment.
- Fix a harmless typo in a macro name in `<thrust/detail/alignment.h>`.
- Add a `DECLARE_UNITTEST_WITH_NAME` helper macro to the unit test framework.
- Make the test framework emit a diagnostic if a device is skipped because
  device code was not compiled for the device's architecture.

Bug 2509847

Author: brycelelbach

testing/complex.cu

@@ -11,6 +11,28 @@
    and takes a lot of time to run.   
  */
 
+template<typename T>
+struct TestComplexSizeAndAlignment
+{
+  void operator()()
+  {
+    THRUST_STATIC_ASSERT(
+      sizeof(thrust::complex<T>) == sizeof(T) * 2
+    );
+    THRUST_STATIC_ASSERT(
+      THRUST_ALIGNOF(thrust::complex<T>) == THRUST_ALIGNOF(T) * 2
+    );
+
+    THRUST_STATIC_ASSERT(
+      sizeof(thrust::complex<T const>) == sizeof(T) * 2
+    );
+    THRUST_STATIC_ASSERT(
+      THRUST_ALIGNOF(thrust::complex<T const>) == THRUST_ALIGNOF(T) * 2
+    );
+  }
+};
+SimpleUnitTest<TestComplexSizeAndAlignment, FloatingPointTypes> TestComplexSizeAndAlignmentInstance;
+
 template<typename T>
 struct TestComplexConstructors
 {
@@ -282,7 +304,6 @@ struct TestComplexStreamOperators
     ASSERT_ALMOST_EQUAL(a,b);
   }
 };
-
 SimpleUnitTest<TestComplexStreamOperators, FloatingPointTypes> TestComplexStreamOperatorsInstance;
 
 #if THRUST_CPP_DIALECT >= 2011
@@ -308,3 +329,4 @@ struct TestComplexStdComplexDeviceInterop
 };
 SimpleUnitTest<TestComplexStdComplexDeviceInterop, FloatingPointTypes> TestComplexStdComplexDeviceInteropInstance;
 #endif
+

testing/cuda/complex.cu

@@ -0,0 +1,53 @@
+#include <unittest/unittest.h>
+
+#include <thrust/complex.h>
+#include <thrust/detail/preprocessor.h>
+#include <thrust/detail/alignment.h>
+
+#include <cuda_fp16.h>
+
+template <typename T, typename VectorT>
+void TestComplexAlignment()
+{
+  THRUST_STATIC_ASSERT(
+    sizeof(thrust::complex<T>) == sizeof(VectorT)
+  );
+  THRUST_STATIC_ASSERT(
+    THRUST_ALIGNOF(thrust::complex<T>) == THRUST_ALIGNOF(VectorT)
+  );
+
+  THRUST_STATIC_ASSERT(
+    sizeof(thrust::complex<T const>) == sizeof(VectorT)
+  );
+  THRUST_STATIC_ASSERT(
+    THRUST_ALIGNOF(thrust::complex<T const>) == THRUST_ALIGNOF(VectorT)
+  );
+}
+DECLARE_UNITTEST_WITH_NAME(
+  THRUST_PP_EXPAND_ARGS(TestComplexAlignment<char, char2>)
+, TestComplexCharAlignment
+);
+DECLARE_UNITTEST_WITH_NAME(
+  THRUST_PP_EXPAND_ARGS(TestComplexAlignment<short, short2>)
+, TestComplexShortAlignment
+);
+DECLARE_UNITTEST_WITH_NAME(
+  THRUST_PP_EXPAND_ARGS(TestComplexAlignment<int, int2>)
+, TestComplexIntAlignment
+);
+DECLARE_UNITTEST_WITH_NAME(
+  THRUST_PP_EXPAND_ARGS(TestComplexAlignment<long, long2>)
+, TestComplexLongAlignment
+);
+DECLARE_UNITTEST_WITH_NAME(
+  THRUST_PP_EXPAND_ARGS(TestComplexAlignment<__half, __half2>)
+, TestComplexHalfAlignment
+);
+DECLARE_UNITTEST_WITH_NAME(
+  THRUST_PP_EXPAND_ARGS(TestComplexAlignment<float, float2>)
+, TestComplexFloatAlignment
+);
+DECLARE_UNITTEST_WITH_NAME(
+  THRUST_PP_EXPAND_ARGS(TestComplexAlignment<double, double2>)
+, TestComplexDoubleAlignment
+);

testing/unittest/cuda/testframework.cu

@@ -12,7 +12,12 @@ bool binary_exists_for_current_device()
   // we didn't compile a binary compatible with the current device
   cudaFuncAttributes attr;
   cudaError_t error = cudaFuncGetAttributes(&attr, dummy_kernel);
-  return error == cudaSuccess;
+
+  // clear the CUDA global error state if we just set it, so that
+  // check_cuda_error doesn't complain
+  if (cudaSuccess != error) (void)cudaGetLastError();
+
+  return cudaSuccess == error;
 }
 
 void list_devices(void)
@@ -159,13 +164,17 @@ bool CUDATestDriver::run_tests(const ArgumentSet &args, const ArgumentMap &kwarg
     // set the device
     cudaSetDevice(*device);
 
-    cudaDeviceSynchronize();
-
     // check if a binary exists for this device
     // if none exists, skip the device silently unless this is the only one we're targeting
     if(devices.size() > 1 && !binary_exists_for_current_device())
     {
-      continue;     
+      // note which device we're skipping
+      cudaDeviceProp deviceProp;
+      cudaGetDeviceProperties(&deviceProp, *device);
+      
+      std::cout << "Skipping Device " << *device << ": \"" << deviceProp.name << "\"" << std::endl;
+
+      continue;
     }
 
     if(!concise)

testing/unittest/testframework.h

@@ -351,6 +351,16 @@ class TEST##UnitTest : public UnitTest {                         \
 };                                                               \
 TEST##UnitTest TEST##Instance
 
+#define DECLARE_UNITTEST_WITH_NAME(TEST, NAME)                   \
+class NAME##UnitTest : public UnitTest {                         \
+    public:                                                      \
+    NAME##UnitTest() : UnitTest(#NAME) {}                        \
+    void run(){                                                  \
+            TEST();                                              \
+    }                                                            \
+};                                                               \
+NAME##UnitTest NAME##Instance
+
 // Macro to create host and device versions of a
 // unit test for a bunch of data types
 #define DECLARE_VECTOR_UNITTEST(VTEST)                          \

thrust/complex.h

@@ -63,6 +63,80 @@ namespace thrust
  *  \{
  */
 
+namespace detail
+{
+  
+template <typename T, std::size_t Align>
+struct complex_storage;
+
+#if __cplusplus >= 201103L                                                    \
+  && (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_GCC)                       \
+  && (THRUST_GCC_VERSION >= 40800)
+  // C++11 implementation, excluding GCC 4.7, which doesn't have `alignas`.
+  template <typename T, std::size_t Align>
+  struct complex_storage
+  {
+    struct alignas(Align) type { T x; T y; };
+  };
+#elif  (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_MSVC)                    \
+    || (   (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_GCC)                 \
+        && (THRUST_GCC_VERSION < 40300))
+  // C++03 implementation for MSVC and GCC <= 4.2.
+  // 
+  // We have to implement `aligned_type` with specializations for MSVC
+  // and GCC 4.2 and older because they require literals as arguments to 
+  // their alignment attribute.
+
+  #if (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_MSVC)
+    // MSVC implementation.
+    #define THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(X)                   \
+      template <typename T>                                                   \
+      struct complex_storage<T, X>                                            \
+      {                                                                       \
+        __declspec(align(X)) struct type { T x; T y; };                       \
+      };                                                                      \
+      /**/
+  #else
+    // GCC <= 4.2 implementation.
+    #define THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(X)                   \
+      template <typename T>                                                   \
+      struct complex_storage<T, X>                                            \
+      {                                                                       \
+        struct type { T x; T y; } __attribute__((aligned(X)));                \
+      };                                                                      \
+      /**/
+  #endif
+
+  // The primary template is a fallback, which doesn't specify any alignment.
+  // It's only used when T is very large and we're using an older compilers
+  // which we have to fully specialize each alignment case.
+  template <typename T, std::size_t Align>
+  struct complex_storage
+  {
+    T x; T y;
+  };
+  
+  THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(1);
+  THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(2);
+  THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(4);
+  THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(8);
+  THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(16);
+  THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(32);
+  THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(64);
+  THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(128);
+
+  #undef THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION
+#else
+  // C++03 implementation for GCC > 4.2, Clang, PGI, ICPC, and xlC.
+  template <typename T, std::size_t Align>
+  struct complex_storage
+  {
+    struct type { T x; T y; } __attribute__((aligned(Align)));
+  };
+#endif
+
+} // end namespace detail
+
   /*! \p complex is the Thrust equivalent to <tt>std::complex</tt>. It is
    *  functionally identical to it, but can also be used in device code which
    *  <tt>std::complex</tt> currently cannot.
@@ -377,31 +451,7 @@ struct complex
   operator std::complex<T>() const { return std::complex<T>(real(), imag()); }
 
 private:
-  /*! \cond
-   */
-  struct generic_storage_type { T x; T y; };
-  /*! \endcond
-   */
-
-#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
-  typedef typename detail::conditional<
-    detail::is_same<T, float>::value, float2,
-    typename detail::conditional<
-      detail::is_same<T, float const>::value, float2 const,
-      typename detail::conditional<
-        detail::is_same<T, double>::value, double2,
-        typename detail::conditional<
-          detail::is_same<T, double const>::value, double2 const,
-          generic_storage_type
-        >::type
-      >::type
-    >::type
-  >::type storage_type;
-#else
-  typedef generic_storage_type storage_type;
-#endif
-
-  storage_type data;
+  typename detail::complex_storage<T, sizeof(T) * 2>::type data;
 };
 
 

thrust/detail/alignment.h

@@ -100,7 +100,7 @@ struct aligned_type;
 #if __cplusplus >= 201103L                                                     \
   && (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_GCC)                        \
   && (THRUST_GCC_VERSION >= 40800)
-    // GCC 4.7 doesn't have `alignas`.
+    // C++11 implementation, excluding GCC 4.7, which doesn't have `alignas`.
     template <std::size_t Align>
     struct aligned_type
     {
@@ -109,39 +109,44 @@ struct aligned_type;
 #elif  (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_MSVC)                    \
     || (   (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_GCC)                 \
         && (THRUST_GCC_VERSION < 40300))
+    // C++03 implementation for MSVC and GCC <= 4.2.
+    // 
     // We have to implement `aligned_type` with specializations for MSVC
     // and GCC 4.2.x and older because they require literals as arguments to 
     // their alignment attribute.
 
     #if (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_MSVC)
-        #define THRUST_DEFINE_ALIGNED_BYTE_SPECIALIZATION(X)                  \
+        // MSVC implementation.
+        #define THRUST_DEFINE_ALIGNED_TYPE_SPECIALIZATION(X)                  \
             template <>                                                       \
-            struct aligned_type<X>                                    \
+            struct aligned_type<X>                                            \
             {                                                                 \
                 __declspec(align(X)) struct type {};                          \
             };                                                                \
             /**/
     #else
-        #define THRUST_DEFINE_ALIGNED_BYTE_SPECIALIZATION(X)                  \
+        // GCC <= 4.2 implementation.
+        #define THRUST_DEFINE_ALIGNED_TYPE_SPECIALIZATION(X)                  \
             template <>                                                       \
-            struct aligned_type<X>                                    \
+            struct aligned_type<X>                                            \
             {                                                                 \
                 struct type {} __attribute__((aligned(X)));                   \
             };                                                                \
             /**/
     #endif
 
-    THRUST_DEFINE_ALIGNED_BYTE_SPECIALIZATION(1);
-    THRUST_DEFINE_ALIGNED_BYTE_SPECIALIZATION(2);
-    THRUST_DEFINE_ALIGNED_BYTE_SPECIALIZATION(4);
-    THRUST_DEFINE_ALIGNED_BYTE_SPECIALIZATION(8);
-    THRUST_DEFINE_ALIGNED_BYTE_SPECIALIZATION(16);
-    THRUST_DEFINE_ALIGNED_BYTE_SPECIALIZATION(32);
-    THRUST_DEFINE_ALIGNED_BYTE_SPECIALIZATION(64);
-    THRUST_DEFINE_ALIGNED_BYTE_SPECIALIZATION(128);
-
-    #undef THRUST_DEFINE_ALIGNED_BYTE_SPECIALIZATION
+    THRUST_DEFINE_ALIGNED_TYPE_SPECIALIZATION(1);
+    THRUST_DEFINE_ALIGNED_TYPE_SPECIALIZATION(2);
+    THRUST_DEFINE_ALIGNED_TYPE_SPECIALIZATION(4);
+    THRUST_DEFINE_ALIGNED_TYPE_SPECIALIZATION(8);
+    THRUST_DEFINE_ALIGNED_TYPE_SPECIALIZATION(16);
+    THRUST_DEFINE_ALIGNED_TYPE_SPECIALIZATION(32);
+    THRUST_DEFINE_ALIGNED_TYPE_SPECIALIZATION(64);
+    THRUST_DEFINE_ALIGNED_TYPE_SPECIALIZATION(128);
+
+    #undef THRUST_DEFINE_ALIGNED_TYPE_SPECIALIZATION
 #else
+    // C++03 implementation for GCC > 4.2, Clang, PGI, ICPC, and xlC.
     template <std::size_t Align>
     struct aligned_type
     {