fixing knn sycl for single precision execution

Author: adarshyoga

dpbench/benchmarks/knn/knn_sycl_native_ext/knn_sycl/_knn_kernel.hpp

@@ -5,13 +5,15 @@
 #include <CL/sycl.hpp>
 #include <cmath>
 
-struct neighbors
+template <typename FpTy, typename IntTy> class theKernel;
+
+template <typename FpTy> struct neighbors
 {
-    double dist;
+    FpTy dist;
     size_t label;
 };
 
-template <typename FpTy>
+template <typename FpTy, typename IntTy>
 sycl::event knn_impl(sycl::queue q,
                      FpTy *d_train,
                      size_t *d_train_labels,
@@ -20,18 +22,18 @@ sycl::event knn_impl(sycl::queue q,
                      size_t classes_num,
                      size_t train_size,
                      size_t test_size,
-                     size_t *d_predictions,
+                     IntTy *d_predictions,
                      FpTy *d_votes_to_classes,
                      size_t data_dim)
 {
     sycl::event partial_hists_ev = q.submit([&](sycl::handler &h) {
-        h.parallel_for<class theKernel>(
+        h.parallel_for<theKernel<FpTy, IntTy>>(
             sycl::range<1>{test_size}, [=](sycl::id<1> myID) {
                 size_t i = myID[0];
 
                 // here k has to be 5 in order to match with numpy no. of
                 // neighbors
-                struct neighbors queue_neighbors[5];
+                struct neighbors<FpTy> queue_neighbors[5];
 
                 // count distances
                 for (size_t j = 0; j < k; ++j) {
@@ -102,10 +104,10 @@ sycl::event knn_impl(sycl::queue q,
                                        queue_neighbors[j].label]++;
                 }
 
-                size_t max_ind = 0;
+                IntTy max_ind = 0;
                 FpTy max_value = 0.0;
 
-                for (size_t j = 0; j < classes_num; ++j) {
+                for (IntTy j = 0; j < (IntTy)classes_num; ++j) {
                     if (d_votes_to_classes[i * classes_num + j] > max_value) {
                         max_value = d_votes_to_classes[i * classes_num + j];
                         max_ind = j;

dpbench/benchmarks/knn/knn_sycl_native_ext/knn_sycl/_knn_sycl.cpp

@@ -5,8 +5,6 @@
 #include "_knn_kernel.hpp"
 #include <dpctl4pybind11.hpp>
 
-#include <iostream>
-
 template <typename... Args> bool ensure_compatibility(const Args &...args)
 {
     std::vector<dpctl::tensor::usm_ndarray> arrays = {args...};
@@ -41,16 +39,28 @@ void knn_sync(dpctl::tensor::usm_ndarray x_train,
                               votes_to_classes))
         throw std::runtime_error("Input arrays are not acceptable.");
 
-    if (x_train.get_typenum() != UAR_DOUBLE) {
-        throw std::runtime_error("Expected a double precision FP array.");
+    auto typenum = x_train.get_typenum();
+    if (typenum == UAR_FLOAT) {
+        sycl::event res_ev = knn_impl<float, unsigned int>(
+            x_train.get_queue(), x_train.get_data<float>(),
+            y_train.get_data<size_t>(), x_test.get_data<float>(), k,
+            classes_num, train_size, test_size,
+            predictions.get_data<unsigned int>(),
+            votes_to_classes.get_data<float>(), data_dim);
+        res_ev.wait();
+    }
+    else if (typenum == UAR_DOUBLE) {
+        sycl::event res_ev = knn_impl<double, size_t>(
+            x_train.get_queue(), x_train.get_data<double>(),
+            y_train.get_data<size_t>(), x_test.get_data<double>(), k,
+            classes_num, train_size, test_size, predictions.get_data<size_t>(),
+            votes_to_classes.get_data<double>(), data_dim);
+        res_ev.wait();
+    }
+    else {
+        throw std::runtime_error(
+            "Expected a double or single precision FP array.");
     }
-
-    sycl::event res_ev = knn_impl(
-        x_train.get_queue(), x_train.get_data<double>(),
-        y_train.get_data<size_t>(), x_test.get_data<double>(), k, classes_num,
-        train_size, test_size, predictions.get_data<size_t>(),
-        votes_to_classes.get_data<double>(), data_dim);
-    res_ev.wait();
 }
 
 PYBIND11_MODULE(_knn_sycl, m)