To avoid splitting increases maintainability burden greatly, use -DENABLE_FP64 macro to split 'double' type code as following:

```
// test.cpp
// RUN: ... -DENABLE_FP64=false
constexpr bool EnableFP64 = ENABLE_FP64;

void test() {
   // non-fp64-case
   do_smth();

   if constexpr (EnableFP64)
     do_smth64();
}

// test-fp64.cpp
; RUN: ... -DENABLE_FP64=true
\#include <test-fp.cpp>
```

Author: myler

SYCL/AtomicRef/assignment_atomic64.cpp

@@ -17,6 +17,9 @@ int main() {
   }
 
   constexpr int N = 32;
+#ifdef ENABLE_FP64
+  assignment_test<double>(q, N);
+#endif
 
   // Include long tests if they are 64 bits wide
   if constexpr (sizeof(long) == 8) {

SYCL/AtomicRef/assignment_atomic64_aspect-fp64.cpp

@@ -1,5 +1,5 @@
 // REQUIRES: aspect-fp64
-// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DENABLE_FP64 %s -o %t.out
 // RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
@@ -8,23 +8,4 @@
 // XFAIL: hip
 // Expected failure because hip does not have atomic64 check implementation
 
-#include "assignment.h"
-#include <iostream>
-using namespace sycl;
-
-int main() {
-  queue q;
-  if (!q.get_device().has(aspect::fp64)) {
-    std::cout << "Skipping test\n";
-    return 0;
-  }
-  if (!q.get_device().has(aspect::atomic64)) {
-    std::cout << "Skipping test\n";
-    return 0;
-  }
-
-  constexpr int N = 32;
-  assignment_test<double>(q, N);
-
-  std::cout << "Test passed." << std::endl;
-}
+#include "assignment_atomic64.cpp"

SYCL/AtomicRef/assignment_atomic64_generic.cpp

@@ -20,7 +20,9 @@ int main() {
   }
 
   constexpr int N = 32;
-
+#ifdef ENABLE_FP64
+  assignment_generic_test<double>(q, N);
+#endif
   // Include long tests if they are 64 bits wide
   if constexpr (sizeof(long) == 8) {
     assignment_generic_test<long>(q, N);
@@ -37,6 +39,5 @@ int main() {
   if constexpr (sizeof(char *) == 8) {
     assignment_generic_test<char *>(q, N);
   }
-
   std::cout << "Test passed." << std::endl;
 }

SYCL/AtomicRef/assignment_atomic64_generic_aspect-fp64.cpp

@@ -1,5 +1,5 @@
 // REQUIRES: aspect-fp64
-// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DENABLE_FP64 %s -o %t.out
 // RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
@@ -8,23 +8,4 @@
 // CUDA backend has had no support for the generic address space yet
 // XFAIL: cuda || hip
 
-#include "assignment.h"
-#include <iostream>
-using namespace sycl;
-
-int main() {
-  queue q;
-  if (!q.get_device().has(aspect::fp64)) {
-    std::cout << "Skipping test\n";
-    return 0;
-  }
-  if (!q.get_device().has(aspect::atomic64)) {
-    std::cout << "Skipping test\n";
-    return 0;
-  }
-
-  constexpr int N = 32;
-  assignment_generic_test<double>(q, N);
-
-  std::cout << "Test passed." << std::endl;
-}
+#include "assignment_atomic64_generic.cpp"

SYCL/Basic/buffer/buffer.cpp

@@ -25,6 +25,7 @@ int main() {
   int data = 5;
   bool failed = false;
   buffer<int, 1> buf(&data, range<1>(1));
+
   {
     int data1[10] = {-1, -1, -1, -1, -1, -1, -1, -1, -1, -1};
     {
@@ -508,16 +509,26 @@ int main() {
     size_t size = 32;
     const size_t dims = 1;
     sycl::range<dims> r(size);
-
     std::shared_ptr<bool> bool_shrd(new bool[size],
                                     [](bool *data) { delete[] data; });
     std::shared_ptr<int> int_shrd(new int[size],
                                   [](int *data) { delete[] data; });
+#ifdef ENABLE_FP64
+    std::shared_ptr<double> double_shrd(new double[size],
+                                        [](double *data) { delete[] data; });
+#endif
 
     std::vector<bool> bool_vector;
     std::vector<int> int_vector;
+#ifdef ENABLE_FP64
+    std::vector<double> double_vector;
+#endif
+
     bool_vector.reserve(size);
     int_vector.reserve(size);
+#ifdef ENABLE_FP64
+    double_vector.reserve(size);
+#endif
 
     sycl::queue Queue;
     std::mutex m;
@@ -528,30 +539,55 @@ int main() {
       sycl::buffer<int, dims> buf_int_shrd(
           int_shrd, r,
           sycl::property_list{sycl::property::buffer::use_mutex(m)});
+#ifdef ENABLE_FP64
+      sycl::buffer<double, dims> buf_double_shrd(
+          double_shrd, r,
+          sycl::property_list{sycl::property::buffer::use_mutex(m)});
+#endif
       m.lock();
       std::fill(bool_shrd.get(), (bool_shrd.get() + size), bool());
       std::fill(int_shrd.get(), (int_shrd.get() + size), int());
+#ifdef ENABLE_FP64
+      std::fill(double_shrd.get(), (double_shrd.get() + size), double());
+#endif
       m.unlock();
-
       buf_bool_shrd.set_final_data(bool_vector.begin());
       buf_int_shrd.set_final_data(int_vector.begin());
+#ifdef ENABLE_FP64
+      buf_double_shrd.set_final_data(double_vector.begin());
+#endif
+
       buf_bool_shrd.set_write_back(true);
       buf_int_shrd.set_write_back(true);
+#ifdef ENABLE_FP64
+      buf_double_shrd.set_write_back(true);
+#endif
 
       Queue.submit([&](sycl::handler &cgh) {
         auto Accessor_bool =
             buf_bool_shrd.get_access<sycl::access::mode::write>(cgh);
         auto Accessor_int =
             buf_int_shrd.get_access<sycl::access::mode::write>(cgh);
+#ifdef ENABLE_FP64
+        auto Accessor_double =
+            buf_double_shrd.get_access<sycl::access::mode::write>(cgh);
+#endif
         cgh.parallel_for<class FillBuffer>(r, [=](sycl::id<1> WIid) {
           Accessor_bool[WIid] = true;
           Accessor_int[WIid] = 3;
+#ifdef ENABLE_FP64
+          Accessor_double[WIid] = 7.5;
+#endif
         });
       });
     } // Data is copied back
 
     for (size_t i = 0; i < size; i++) {
       if (bool_vector[i] != true || int_vector[i] != 3) {
+#ifdef ENABLE_FP64
+      if (bool_vector[i] != true || int_vector[i] != 3 ||
+          double_vector[i] != 7.5) {
+#endif
         assert(false && "Data was not copied back");
         return 1;
       }

SYCL/Basic/buffer/buffer_aspect-fp64.cpp

@@ -1,5 +1,5 @@
 // REQUIRES: aspect-fp64
-// RUN: %clangxx %cxx_std_optionc++17 %s -o %t1.out %sycl_options
+// RUN: %clangxx %cxx_std_optionc++17 -DENABLE_FP64 %s -o %t1.out %sycl_options
 // RUN: %HOST_RUN_PLACEHOLDER %t1.out
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t2.out
 // RUN: %HOST_RUN_PLACEHOLDER %t2.out
@@ -15,56 +15,4 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include <sycl/sycl.hpp>
-
-#include <cassert>
-#include <memory>
-
-using namespace sycl;
-
-template <class T> constexpr T write_back_result = T(3);
-template <> constexpr double write_back_result<double> = double(7.5);
-template <class T> class fill_buffer_for_write_back {};
-
-template <class T, int D> void check_set_write_back() {
-  size_t size = 32;
-  sycl::range<D> r(size);
-  std::shared_ptr<T> shrd(new T[size], [](T *data) { delete[] data; });
-  std::vector<T> vector;
-  vector.reserve(size);
-  sycl::queue Queue;
-  std::mutex m;
-  {
-    sycl::buffer<T, D> buf_shrd(
-        shrd, r, sycl::property_list{sycl::property::buffer::use_mutex(m)});
-    m.lock();
-    std::fill(shrd.get(), (shrd.get() + size), T());
-    m.unlock();
-    buf_shrd.set_final_data(vector.begin());
-    buf_shrd.set_write_back(true);
-    Queue.submit([&](sycl::handler &cgh) {
-      auto Accessor =
-          buf_shrd.template get_access<sycl::access::mode::write>(cgh);
-      cgh.parallel_for<class fill_buffer_for_write_back<T>>(
-          r, [=](sycl::id<1> WIid) { Accessor[WIid] = write_back_result<T>; });
-    });
-  } // Data is copied back
-  for (size_t i = 0; i < size; i++) {
-    if (vector[i] != write_back_result<T>) {
-      assert(false && "Data was not copied back");
-    }
-  }
-}
-
-int main() {
-  // Check that data is copied back after forcing write-back using
-  // set_write_back
-  queue q;
-  if (!q.get_device().has(aspect::fp64)) {
-    std::cout << "Skipping test\n";
-    return 0;
-  }
-
-  check_set_write_back<double, 1>();
-  return 0;
-}
+#include "buffer.cpp"

SYCL/DeprecatedFeatures/SpecConsts1.2.1/specialization_constants.cpp

@@ -23,7 +23,6 @@
 #define HALF 0 // FIXME Spec constants do not support half type yet
 
 class SpecializedKernel;
-
 class MyBoolConst;
 class MyInt8Const;
 class MyUInt8Const;
@@ -35,6 +34,9 @@ class MyInt64Const;
 class MyUInt64Const;
 class MyHalfConst;
 class MyFloatConst;
+#ifdef ENABLE_FP64
+class MyDoubleConst;
+#endif
 
 using namespace sycl;
 
@@ -53,6 +55,9 @@ int64_t int64_ref = rnd() % std::numeric_limits<int64_t>::max();
 uint64_t uint64_ref = rnd() % std::numeric_limits<uint64_t>::max();
 half half_ref = rnd() % std::numeric_limits<uint16_t>::max();
 float float_ref = rnd() % std::numeric_limits<uint32_t>::max();
+#ifdef ENABLE_FP64
+double double_ref = rnd() % std::numeric_limits<uint64_t>::max();
+#endif
 
 template <typename T1, typename T2>
 bool check(const T1 &test, const T2 &ref, std::string type) {
@@ -107,9 +112,11 @@ int main(int argc, char **argv) {
 #endif
     ext::oneapi::experimental::spec_constant<float, MyFloatConst> f32 =
         prog.set_spec_constant<MyFloatConst>(float_ref);
-
+#ifdef ENABLE_FP64
+    ext::oneapi::experimental::spec_constant<double, MyDoubleConst> f64 =
+        prog.set_spec_constant<MyDoubleConst>(double_ref);
+#endif
     prog.build_with_kernel_type<SpecializedKernel>();
-
     bool bool_test = 0;
     int8_t int8_test = 0;
     uint8_t uint8_test = 0;
@@ -121,6 +128,9 @@ int main(int argc, char **argv) {
     uint64_t uint64_test = 0;
     half half_test = 0;
     float float_test = 0;
+#ifdef ENABLE_FP64
+    double double_test = 0;
+#endif
 
     {
       buffer<bool> bool_buf(&bool_test, 1);
@@ -134,6 +144,9 @@ int main(int argc, char **argv) {
       buffer<uint64_t> uint64_buf(&uint64_test, 1);
       buffer<half> half_buf(&half_test, 1);
       buffer<float> float_buf(&float_test, 1);
+#ifdef ENABLE_FP64
+      buffer<double> double_buf(&double_test, 1);
+#endif
 
       q.submit([&](handler &cgh) {
         auto bool_acc = bool_buf.get_access<access::mode::write>(cgh);
@@ -147,6 +160,9 @@ int main(int argc, char **argv) {
         auto uint64_acc = uint64_buf.get_access<access::mode::write>(cgh);
         auto half_acc = half_buf.get_access<access::mode::write>(cgh);
         auto float_acc = float_buf.get_access<access::mode::write>(cgh);
+#ifdef ENABLE_FP64
+        auto double_acc = double_buf.get_access<access::mode::write>(cgh);
+#endif
         cgh.single_task<SpecializedKernel>(prog.get_kernel<SpecializedKernel>(),
                                            [=]() {
                                              bool_acc[0] = i1.get();
@@ -162,6 +178,9 @@ int main(int argc, char **argv) {
                                              half_acc[0] = f16.get();
 #endif
                                              float_acc[0] = f32.get();
+#ifdef ENABLE_FP64
+                                             double_acc[0] = f64.get();
+#endif
                                            });
       });
     }
@@ -189,6 +208,10 @@ int main(int argc, char **argv) {
 #endif
     if (!check(float_test, float_ref, "float"))
       return 1;
+#ifdef ENABLE_FP64
+    if (!check(double_test, double_ref, "double"))
+      return 1;
+#endif
   } catch (const exception &e) {
     std::cout << "an async SYCL exception was caught: "
               << std::string(e.what());