[SYCL][fp64][SYCLLowerIR] Fix an issue when double type inside multip…

…le nested structs is not identified (#14523)

There are cases where some instructions use data structures that contain
double data type deeply nested inside multiple levels.
Example:
struct T1 {
  double x;
};
struct T2 {
  struct T1 y;
};
T2 x;

In such case, the fp64 emulation support in SYCLPropagateAspectUsage
pass was not correctly identifying T2 as a type that has doubles. This
led to wrong behavior when adding fp64 aspect.
This PR fixes the issue.

Thanks

---------

Signed-off-by: Arvind Sudarsanam <arvind.sudarsanam@intel.com>

llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp

@@ -267,7 +267,7 @@ bool hasDoubleType(const Type *T) {
   if (T->isDoubleTy())
     return true;
   for (const Type *TT : T->subtypes())
-    if (TT->isDoubleTy())
+    if (hasDoubleType(TT))
       return true;
   return false;
 }

sycl/test/check_device_code/optional_kernel_features/fp64-conv-emu-convert.cpp

@@ -8,14 +8,23 @@
 
 #include <sycl/sycl.hpp>
 
+struct T1 {
+  double x;
+};
+struct T2 {
+  struct T1 y;
+};
+
 int main() {
   sycl::queue q;
   q.submit([&](sycl::handler &h) {
     h.single_task([=]() {
       double a[10];
       double b[10];
+      struct T2 tmp;
       int i = 4;
       b[i] = (double)((float)(a[i]));
+      tmp.y.x = (double)((float)(a[i]));
     });
   });
   return 0;

sycl/test/check_device_code/optional_kernel_features/fp64-conv-emu-test-call-convert.cpp

@@ -8,6 +8,13 @@
 
 #include <sycl/sycl.hpp>
 
+struct T1 {
+  double x;
+};
+struct T2 {
+  struct T1 y;
+};
+
 double bar_convert(double a) { return (double)((float)(a)); }
 
 int main() {
@@ -16,8 +23,10 @@ int main() {
     h.single_task([=]() {
       double a[10];
       double b[10];
+      struct T2 tmp;
       int i = 4;
       b[i] = bar_convert(a[i]);
+      tmp.y.x = bar_convert(a[i]);
     });
   });
   return 0;

sycl/test/check_device_code/optional_kernel_features/fp64-conv-emu-test-call.cpp

@@ -11,6 +11,13 @@
 
 #include <sycl/sycl.hpp>
 
+struct T1 {
+  double x;
+};
+struct T2 {
+  struct T1 y;
+};
+
 SYCL_EXTERNAL double foo(double a);
 double bar_compute(double a) { return a + 1.0; }
 
@@ -20,8 +27,10 @@ int main() {
     h.single_task([=]() {
       double a[10];
       double b[10];
+      struct T2 tmp;
       int i = 4;
       b[i] = foo(a[i]);
+      tmp.y.x = foo(a[i]);
     });
   });
   q.submit([&](sycl::handler &h) {

sycl/test/check_device_code/optional_kernel_features/fp64-conv-emu-with-recursive-type-decls.cpp

@@ -0,0 +1,29 @@
+// UNSUPPORTED: cuda, hip
+
+// RUN: %clangxx -O0 -fsycl -fsycl-device-only -fsycl-targets=spir64_gen -fsycl-fp64-conv-emu %s -c -S -emit-llvm -o- | FileCheck %s
+
+// CHECK-NOT: define {{.*}} spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_(){{.*}} !sycl_used_aspects
+
+// Tests if -fsycl-fp64-conv-emu option helps to correctly generate fp64 aspect
+// in the presence of possibly recursive type declaration.
+
+#include <sycl/sycl.hpp>
+
+struct T2;
+struct T1 {
+  struct T2 *x;
+};
+struct T2 {
+  struct T1 *y;
+};
+
+int main() {
+  sycl::queue q;
+  q.submit([&](sycl::handler &h) {
+    h.single_task([=]() {
+      struct T1 tmp1;
+      struct T2 tmp2;
+    });
+  });
+  return 0;
+}

sycl/test/check_device_code/optional_kernel_features/fp64-conv-emu.cpp

@@ -9,14 +9,22 @@
 
 #include <sycl/sycl.hpp>
 
+struct T1 {
+  double x;
+};
+struct T2 {
+  struct T1 y;
+};
 int main() {
   sycl::queue q;
   q.submit([&](sycl::handler &h) {
     h.single_task([=]() {
       double a[10];
       double b[10];
+      struct T2 tmp;
       int i = 4;
       b[i] = a[i] + 1.0;
+      tmp.y.x = a[i] + 1.0;
     });
   });
   return 0;