[SYCLomatic oneapi-src#507] Enable test for complex, array, tuple in …

…libcu (oneapi-src#204)

Signed-off-by: Ni, Wenhui <wenhui.ni@intel.com>

features/feature_case/libcu/libcu_array.cu

@@ -0,0 +1,45 @@
+#include <cuda_runtime.h>
+#include <cuda/std/array>
+
+template <class T>
+__host__ __device__ void
+test(T *res)
+{
+  cuda::std::array<T, 3> arr = {1, 2, 3.5};
+  *(res) = arr.at(0); 
+  *(res+1) = arr.at(1); 
+  *(res+2) = arr.at(2);
+  *(res+3) = *(arr.begin());
+  *(res+4) = arr.size();
+}
+
+__global__ void test_global(float * res)
+{
+  test<float>(res);
+}
+
+int main(int, char **)
+{
+
+  float *floatRes = (float *)malloc(5 * sizeof(float));
+  test<float>(floatRes);
+  //test<double>(doubleRes);
+  float *hostRes = (float *)malloc(5 * sizeof(float));
+  float *deviceRes;
+  cudaMalloc((float **)&deviceRes, 5 * sizeof(float));
+  test_global<<<1, 1>>>(deviceRes);
+  cudaMemcpy(hostRes, deviceRes, sizeof(float) * 5, cudaMemcpyDeviceToHost);
+  cudaFree(deviceRes);
+
+  for (int i = 0;i<5;++i){
+    if(hostRes[i]!=floatRes[i]){
+      free(hostRes);
+      free(floatRes);
+      return 1;
+    }
+  }
+  free(hostRes);
+  free(floatRes);
+  return 0;
+
+}

features/feature_case/libcu/libcu_complex.cu

@@ -0,0 +1,66 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+// <cuda/std/complex>
+
+#include <cuda/std/complex>
+
+template <class T>
+__host__ __device__ void
+test(T *res)
+{
+  cuda::std::complex<T> x(1.5, 2.5);
+  cuda::std::complex<T> y(2.5, 3);
+  T *a = (T *)&x;
+  a[0] = 5;
+  a[1] = 6;
+  *(res) = x.real() * x.imag();
+  cuda::std::complex<T> z = x / y;
+  *(res +1) =z.real();
+  *(res + 2) = z.imag();
+  z = x + y;
+  *(res +3) =z.real();
+  *(res + 4) = z.imag();
+  z = x - y;
+  *(res +5) =z.real();
+  *(res + 6) = z.imag();
+  z = x * y;
+  *(res +7) =z.real();
+  *(res + 8) = z.imag();
+}
+
+__global__ void test_global(float * res)
+{
+  test<float>(res);
+}
+
+int main(int, char **)
+{
+
+  float *floatRes = (float *)malloc(9 * sizeof(float));
+  test<float>(floatRes);
+  //test<double>(doubleRes);
+  float *hostRes = (float *)malloc(9 * sizeof(float));
+  float *deviceRes;
+  cudaMalloc((float **)&deviceRes, 9 * sizeof(float));
+  test_global<<<1, 1>>>(deviceRes);
+  cudaMemcpy(hostRes, deviceRes, sizeof(float) * 9, cudaMemcpyDeviceToHost);
+  cudaFree(deviceRes);
+
+  for (int i = 0;i<9;++i){
+    if(hostRes[i]!=floatRes[i]){
+      free(hostRes);
+      free(floatRes);
+      return 1;
+    }
+  }
+  free(hostRes);
+  free(floatRes);
+  return 0;
+
+}

features/feature_case/libcu/libcu_tuple.cu

@@ -0,0 +1,50 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+
+#include <cuda/std/tuple>
+
+template <class T>
+__host__ __device__ void
+test(T *res)
+{
+  cuda::std::tuple<T, T, T> t = cuda::std::make_tuple(2.0f, 3.0f, 4.0f);
+  *(res) = cuda::std::get<0>(t);
+  *(res+1) = cuda::std::get<1>(t);
+  *(res+2) = cuda::std::get<2>(t);
+}
+
+__global__ void test_global(float * res)
+{
+  test<float>(res);
+}
+
+int main(int, char **)
+{
+
+  float *floatRes = (float *)malloc(3 * sizeof(float));
+  test<float>(floatRes);
+  float *hostRes = (float *)malloc(3 * sizeof(float));
+  float *deviceRes;
+  cudaMalloc((float **)&deviceRes, 3 * sizeof(float));
+  test_global<<<1, 1>>>(deviceRes);
+  cudaMemcpy(hostRes, deviceRes, sizeof(float) * 3, cudaMemcpyDeviceToHost);
+  cudaFree(deviceRes);
+
+  for (int i = 0;i<3;++i){
+    if(hostRes[i]!=floatRes[i]){
+      free(hostRes);
+      free(floatRes);
+      return 1;
+    }
+  }
+  free(hostRes);
+  free(floatRes);
+  return 0;
+
+}

features/features.xml

@@ -397,6 +397,9 @@
     <test testName="ccl" configFile="config/TEMPLATE_ccl.xml" />
     <test testName="cufft_test" configFile="config/TEMPLATE_fft_runable.xml" splitGroup="double"/>
     <test testName="libcu_atomic" configFile="config/TEMPLATE_libcu.xml" />
+    <test testName="libcu_array" configFile="config/TEMPLATE_libcu.xml" />
+    <test testName="libcu_complex" configFile="config/TEMPLATE_libcu.xml" />
+    <test testName="libcu_tuple" configFile="config/TEMPLATE_libcu.xml" />
     <test testName="pointer_attributes" configFile="config/TEMPLATE_pointer_attributes.xml" />
     <test testName="image" configFile="config/TEMPLATE_image.xml" />
     <test testName="defaultStream" configFile="config/TEMPLATE_defaultStream.xml" />

features/test_feature.py

@@ -28,7 +28,7 @@
               'cub_device_scan_inclusive_sum', 'cub_device_scan_exclusive_sum', 'cub_device_select_unique', 'cub_device_radix_sort_keys', 'cub_device_radix_sort_pairs',
               'cub_device_select_flagged', 'cub_device_run_length_encide_encode', 'cub_counting_iterator', 'cub_arg_index_input_iterator',
               'cub_device_inclusive_sum_by_key', 'cub_device_exclusive_sum_by_key', 'cub_device_inclusive_scan_by_key', 'cub_device_exclusive_scan_by_key',
-              'cub_transform_iterator', 'activemask', 'complex', 'thrust-math'
+              'cub_transform_iterator', 'activemask', 'complex', 'thrust-math', 'libcu_array', 'libcu_complex', 'libcu_tuple',
               'user_defined_rules', 'math-exec', 'math-habs', 'math-ext-half', 'math-ext-half2', 'cudnn-activation',
               'cudnn-fill', 'cudnn-lrn', 'cudnn-memory', 'cudnn-pooling', 'cudnn-reorder', 'cudnn-scale', 'cudnn-softmax',
               'cudnn-sum', 'math-funnelshift', 'ccl', 'thrust-sort_by_key', 'thrust-find', 'thrust-inner_product', 'thrust-reduce_by_key',