[SYCL] Cleanup/refactor device_code_tests

sycl/test/check_device_code/ap_fixed.cpp

@@ -8,103 +8,59 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include "CL/__spirv/spirv_ops.hpp"
+#include <sycl/sycl.hpp>
 
-template <int W, int rW, bool S, int I, int rI>
-void sqrt() {
-  sycl::detail::ap_int<W> a;
-  auto ap_fixed_Sqrt = __spirv_FixedSqrtINTEL<W, rW>(a, S, I, rI);
+SYCL_EXTERNAL auto test_sqrt(sycl::detail::ap_int<13> a) {
+  return __spirv_FixedSqrtINTEL<13, 5>(a, false, 2, 2);
   // CHECK: %{{.*}} = call spir_func signext i5 @_Z[[#]]__spirv_FixedSqrtINTEL{{.*}}(i13 signext  %[[#]], i1 zeroext false, i32 2, i32 2, i32 0, i32 0)
 }
 
-template <int W, int rW, bool S, int I, int rI>
-void recip() {
-  sycl::detail::ap_int<W> a;
-  auto ap_fixed_Recip = __spirv_FixedRecipINTEL<W, rW>(a, S, I, rI);
+SYCL_EXTERNAL auto test_recip(sycl::detail::ap_int<3> a) {
+  return __spirv_FixedRecipINTEL<3, 8>(a, true, 4, 4);
   // CHECK: %{{.*}} = call spir_func signext i8 @_Z[[#]]__spirv_FixedRecipINTEL{{.*}}(i3 signext %[[#]], i1 zeroext true, i32 4, i32 4, i32 0, i32 0)
 }
 
-template <int W, int rW, bool S, int I, int rI>
-void rsqrt() {
-  sycl::detail::ap_int<W> a;
-  auto ap_fixed_Rsqrt = __spirv_FixedRsqrtINTEL<W, rW>(a, S, I, rI);
+SYCL_EXTERNAL auto test_rsqrt(sycl::detail::ap_int<11> a) {
+  return __spirv_FixedRsqrtINTEL<11, 10>(a, false, 8, 6);
   // CHECK: %{{.*}} = call spir_func signext i10 @_Z[[#]]__spirv_FixedRsqrtINTEL{{.*}}(i11 signext %[[#]], i1 zeroext false, i32 8, i32 6, i32 0, i32 0)
 }
 
-template <int W, int rW, bool S, int I, int rI>
-void sin() {
-  sycl::detail::ap_int<W> a;
-  auto ap_fixed_Sin = __spirv_FixedSinINTEL<W, rW>(a, S, I, rI);
+SYCL_EXTERNAL auto test_sin(sycl::detail::ap_int<17> a) {
+  return __spirv_FixedSinINTEL<17, 11>(a, true, 7, 5);
   // CHECK: %{{.*}} = call spir_func signext i11 @_Z[[#]]__spirv_FixedSinINTEL{{.*}}(i17 signext %[[#]], i1 zeroext true, i32 7, i32 5, i32 0, i32 0)
 }
 
-template <int W, int rW, bool S, int I, int rI>
-void cos() {
-  sycl::detail::ap_int<W> a;
-  auto ap_fixed_Cos = __spirv_FixedCosINTEL<W, rW>(a, S, I, rI);
+SYCL_EXTERNAL auto test_cos(sycl::detail::ap_int<35> a) {
+  return __spirv_FixedCosINTEL<35, 28>(a, false, 9, 3);
   // CHECK: %{{.*}} = call spir_func signext i28 @_Z[[#]]__spirv_FixedCosINTEL{{.*}}(i35 %[[#]], i1 zeroext false, i32 9, i32 3, i32 0, i32 0)
 }
 
-template <int W, int rW, bool S, int I, int rI>
-void sin_cos() {
-  sycl::detail::ap_int<W> a;
-  auto ap_fixed_SinCos = __spirv_FixedSinCosINTEL<W, rW>(a, S, I, rI);
+SYCL_EXTERNAL auto test_sin_cos(sycl::detail::ap_int<31> a) {
+  return __spirv_FixedSinCosINTEL<31, 20>(a, true, 10, 12);
   // CHECK: %{{.*}} = call spir_func i40 @_Z[[#]]__spirv_FixedSinCosINTEL{{.*}}(i31 signext %[[#]], i1 zeroext true, i32 10, i32 12, i32 0, i32 0)
 }
 
-template <int W, int rW, bool S, int I, int rI>
-void sin_pi() {
-  sycl::detail::ap_int<W> a;
-  auto ap_fixed_SinPi = __spirv_FixedSinPiINTEL<W, rW>(a, S, I, rI);
+SYCL_EXTERNAL auto test_sin_pi(sycl::detail::ap_int<60> a) {
+  return __spirv_FixedSinPiINTEL<60, 5>(a, false, 2, 2);
   // CHECK: %{{.*}} = call spir_func signext i5 @_Z[[#]]__spirv_FixedSinPiINTEL{{.*}}(i60 %[[#]], i1 zeroext false, i32 2, i32 2, i32 0, i32 0)
 }
 
-template <int W, int rW, bool S, int I, int rI>
-void cos_pi() {
-  sycl::detail::ap_int<W> a;
-  auto ap_fixed_CosPi = __spirv_FixedCosPiINTEL<W, rW>(a, S, I, rI);
+SYCL_EXTERNAL auto test_cos_pi(sycl::detail::ap_int<28> a) {
+  return __spirv_FixedCosPiINTEL<28, 16>(a, false, 8, 5);
   // CHECK: %{{.*}} = call spir_func signext i16 @_Z[[#]]__spirv_FixedCosPiINTEL{{.*}}(i28 signext %[[#]], i1 zeroext false, i32 8, i32 5, i32 0, i32 0)
 }
 
-template <int W, int rW, bool S, int I, int rI>
-void sin_cos_pi() {
-  sycl::detail::ap_int<W> a;
-  auto ap_fixed_SinCosPi = __spirv_FixedSinCosPiINTEL<W, rW>(a, S, I, rI);
+SYCL_EXTERNAL auto test_sin_cos_pi(sycl::detail::ap_int<13> a) {
+  return __spirv_FixedSinCosPiINTEL<13, 5>(a, false, 2, 2);
   // CHECK: %{{.*}} = call spir_func signext i10 @_Z[[#]]__spirv_FixedSinCosPiINTEL{{.*}}(i13 signext %[[#]], i1 zeroext false, i32 2, i32 2, i32 0, i32 0)
 }
 
-template <int W, int rW, bool S, int I, int rI>
-void log() {
-  sycl::detail::ap_int<W> a;
-  auto ap_fixed_Log = __spirv_FixedLogINTEL<W, rW>(a, S, I, rI);
+SYCL_EXTERNAL auto test_log(sycl::detail::ap_int<64> a) {
+  return __spirv_FixedLogINTEL<64, 44>(a, true, 24, 22);
   // CHECK: %{{.*}} = call spir_func i44 @_Z[[#]]__spirv_FixedLogINTEL{{.*}}(i64 %[[#]], i1 zeroext true, i32 24, i32 22, i32 0, i32 0)
 }
 
-template <int W, int rW, bool S, int I, int rI>
-void exp() {
-  sycl::detail::ap_int<W> a;
-  auto ap_fixed_Exp = __spirv_FixedExpINTEL<W, rW>(a, S, I, rI);
+SYCL_EXTERNAL auto test_exp(sycl::detail::ap_int<44> a) {
+  return __spirv_FixedExpINTEL<44, 34>(a, false, 20, 20);
   // CHECK: %{{.*}} = call spir_func i34 @_Z[[#]]__spirv_FixedExpINTEL{{.*}}(i44 %[[#]], i1 zeroext false, i32 20, i32 20, i32 0, i32 0)
-}
-
-template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
-  kernelFunc();
-}
-
-int main() {
-  kernel_single_task<class kernel_function>([]() {
-    sqrt<13, 5, false, 2, 2>();
-    recip<3, 8, true, 4, 4>();
-    rsqrt<11, 10, false, 8, 6>();
-    sin<17, 11, true, 7, 5>();
-    cos<35, 28, false, 9, 3>();
-    sin_cos<31, 20, true, 10, 12>();
-    sin_pi<60, 5, false, 2, 2>();
-    cos_pi<28, 16, false, 8, 5>();
-    sin_cos_pi<13, 5, false, 2, 2>();
-    log<64, 44, true, 24, 22>();
-    exp<44, 34, false, 20, 20>();
-  });
-  return 0;
-}
+}

sycl/test/check_device_code/device_global_const_eval_use.cpp

@@ -80,42 +80,38 @@ constexpr device_global<TestStruct2, decltype(properties(device_image_scope))>
     dg_constexpr_constructor_struct{TS4};
 // CHECK: @{{[A-Za-z0-9_]*}}dg_constexpr_constructor_struct = internal addrspace(1) constant { %struct.TestStruct2 } { %struct.TestStruct2 { i32 4 } }, align 4, !spirv.Decorations
 
-int main() {
-  sycl::queue Q;
-  Q.submit([&](sycl::handler &h) {
-    // Simple kernel that just copies over the values from the device_globals so
-    // that we can observe the GlobalVariables that are created to represent
-    // them in the IR
-    h.single_task([=] {
-      // Int and array of ints
-      std::ignore = dg_int;
-      std::ignore = dg_int_arr[0];
-
-      // Char and array of chars
-      std::ignore = dg_char;
-      std::ignore = dg_char_arr[0];
-
-      // Multidimensional array of integers
-      std::ignore = dg_multi_dim_arr[1][1];
-
-      // Float and array of floats
-      std::ignore = dg_float;
-      std::ignore = dg_float_arr[0];
-
-      // Double and array of doubles
-      std::ignore = dg_double;
-      std::ignore = dg_double_arr[0];
-
-      // Bool and array of bools
-      std::ignore = dg_bool;
-      std::ignore = dg_bool_arr[0];
-
-      // Struct and array of structs
-      std::ignore = dg_struct.get().field1;
-      std::ignore = dg_struct_arr[0];
-
-      // Struct with constexpr constructor
-      std::ignore = dg_constexpr_constructor_struct.get().value;
-    });
-  });
+SYCL_EXTERNAL void device_global_const_eval_use() {
+  // Simple kernel that just copies over the values from the device_globals so
+  // that we can observe the GlobalVariables that are created to represent
+  // them in the IR
+
+  // Int and array of ints
+  std::ignore = dg_int;
+  std::ignore = dg_int_arr[0];
+
+  // Char and array of chars
+  std::ignore = dg_char;
+  std::ignore = dg_char_arr[0];
+
+  // Multidimensional array of integers
+  std::ignore = dg_multi_dim_arr[1][1];
+
+  // Float and array of floats
+  std::ignore = dg_float;
+  std::ignore = dg_float_arr[0];
+
+  // Double and array of doubles
+  std::ignore = dg_double;
+  std::ignore = dg_double_arr[0];
+
+  // Bool and array of bools
+  std::ignore = dg_bool;
+  std::ignore = dg_bool_arr[0];
+
+  // Struct and array of structs
+  std::ignore = dg_struct.get().field1;
+  std::ignore = dg_struct_arr[0];
+
+  // Struct with constexpr constructor
+  std::ignore = dg_constexpr_constructor_struct.get().value;
 }

sycl/test/check_device_code/device_has.cpp

@@ -9,10 +9,9 @@
 #include <sycl/sycl.hpp>
 
 using namespace sycl;
-queue q;
 
-// CHECK-ASPECTS: define weak_odr dso_local spir_kernel void @{{.*}}kernel_name_1{{.*}} !sycl_declared_aspects ![[ASPECTS1:[0-9]+]] {{.*}}
-// CHECK-SRCLOC: define weak_odr dso_local spir_kernel void @{{.*}}kernel_name_1{{.*}} !srcloc ![[SRCLOC1:[0-9]+]] {{.*}}
+// CHECK-ASPECTS: define dso_local spir_func void @{{.*}}kernel_name_1{{.*}} !sycl_declared_aspects ![[ASPECTS1:[0-9]+]] {{.*}}
+// CHECK-SRCLOC: define dso_local spir_func void @{{.*}}kernel_name_1{{.*}} !srcloc ![[SRCLOC1:[0-9]+]] {{.*}}
 
 // CHECK-ASPECTS: define {{.*}}spir_func void @{{.*}}func1{{.*}} !sycl_declared_aspects ![[ASPECTS1]]
 // CHECK-ASPECTS-SAME: !sycl_used_aspects ![[ASPECTS1]]
@@ -45,29 +44,19 @@ constexpr sycl::aspect getAspect() { return sycl::aspect::cpu; }
 // CHECK-SRCLOC: define {{.*}}spir_func void @{{.*}}func6{{.*}} !srcloc ![[SRCLOC7:[0-9]+]]
 [[sycl::device_has(getAspect())]] void func6() {}
 
-class KernelFunctor {
-public:
-  [[sycl::device_has(sycl::aspect::cpu)]] void operator()() const {
-    func1();
-    func2();
-    func3();
-    func4<sycl::aspect::host>();
-    func5();
-    func6();
-  }
-};
-
-void foo() {
-  q.submit([&](handler &h) {
-    KernelFunctor f1;
-    h.single_task<class kernel_name_1>(f1);
-    // CHECK-ASPECTS: define weak_odr dso_local spir_kernel void @{{.*}}kernel_name_2{{.*}} !sycl_declared_aspects ![[ASPECTS4:[0-9]+]]
-    // CHECK-SRCLOC: define weak_odr dso_local spir_kernel void @{{.*}}kernel_name_2{{.*}} !srcloc ![[SRCLOC8:[0-9]+]] {{.*}}
-    h.single_task<class kernel_name_2>(
-        []() [[sycl::device_has(sycl::aspect::gpu)]] {});
-  });
+SYCL_EXTERNAL [[sycl::device_has(sycl::aspect::cpu)]] void kernel_name_1() {
+  func1();
+  func2();
+  func3();
+  func4<sycl::aspect::host>();
+  func5();
+  func6();
 }
 
+// CHECK-ASPECTS: define dso_local spir_func void @{{.*}}kernel_name_2{{.*}} !sycl_declared_aspects ![[ASPECTS4:[0-9]+]]
+// CHECK-SRCLOC: define dso_local spir_func void @{{.*}}kernel_name_2{{.*}} !srcloc ![[SRCLOC8:[0-9]+]] {{.*}}
+SYCL_EXTERNAL [[sycl::device_has(sycl::aspect::gpu)]] void kernel_name_2() {}
+
 // CHECK-ASPECTS-DAG: [[ASPECTS1]] = !{![[ASPECTCPU:[0-9]+]]}
 // CHECK-ASPECTS-DAG: [[ASPECTCPU]] = !{!"cpu", i32 1}
 // CHECK-SRCLOC-DAG: [[SRCLOC1]] = !{i32 {{[0-9]+}}}