Remove conditional check in tsne iDGT2D1gpu kernel

Author: mgrabban

tsne/CUDA/src/kernels/nbodyfft.cu

@@ -351,15 +351,16 @@ void iDFT2D1gpu(thrust::complex<float>* din, thrust::complex<float>* dout, int n
     angle = 2.0f * PI * fdividef((float)i, (float)num_cols);
     sum = 0.0f;
 #pragma unroll
-    for (int k = 0; k < num_cols; ++k) {
+    for (int k = 0; k < num_cols/2 + 1; ++k) {
         // sincosf(angle * k, &sinf, &cosf);
         // twiddle = thrust::complex<float>(cosf, sinf);
         TWIDDLE();
-        if (k < (num_cols/2+1)) {
-            sum = sum + din[j * (num_cols/2+1) + k] * twiddle;
-        } else {
-            sum = sum + thrust::conj(din[((num_rows-j)%num_rows) * (num_cols/2+1) + ((num_cols-k)%num_cols)]) * twiddle;
-        }
+        sum = sum + din[j * (num_cols/2+1) + k] * twiddle;
+    }
+#pragma unroll
+    for (int k = num_cols/2 + 1; k < num_cols; ++k) {
+        TWIDDLE();
+        sum = sum + thrust::conj(din[((num_rows-j)%num_rows) * (num_cols/2+1) + ((num_cols-k)%num_cols)]) * twiddle;
     }
 
     dout[i * num_rows + j] = sum;

tsne/HIP/src/kernels/nbodyfft.cpp

@@ -360,15 +360,16 @@ void iDFT2D1gpu(thrust::complex<float>* din, thrust::complex<float>* dout, int n
     angle = 2.0f * PI * fdividef((float)i, (float)num_cols);
     sum = 0.0f;
 #pragma unroll
-    for (int k = 0; k < num_cols; ++k) {
+    for (int k = 0; k < num_cols/2 + 1; ++k) {
         // sincosf(angle * k, &sinf, &cosf);
         // twiddle = thrust::complex<float>(cosf, sinf);
         TWIDDLE();
-        if (k < (num_cols/2+1)) {
-            sum = sum + din[j * (num_cols/2+1) + k] * twiddle;
-        } else {
-            sum = sum + thrust::conj(din[((num_rows-j)%num_rows) * (num_cols/2+1) + ((num_cols-k)%num_cols)]) * twiddle;
-        }
+        sum = sum + din[j * (num_cols/2+1) + k] * twiddle;
+    }
+#pragma unroll
+    for (int k = num_cols/2 + 1; k < num_cols; ++k) {
+        TWIDDLE();
+        sum = sum + thrust::conj(din[((num_rows-j)%num_rows) * (num_cols/2+1) + ((num_cols-k)%num_cols)]) * twiddle;
     }
 
     dout[i * num_rows + j] = sum;

tsne/SYCL/src/kernels/nbodyfft.dp.cpp

@@ -358,17 +358,18 @@ void iDFT2D1gpu(std::complex<float>* din, std::complex<float>* dout, int num_row
     angle = 2.0f * PI * ((float)i / (float)num_cols);
     sum = 0.0f;
 #pragma unroll
-    for (int k = 0; k < num_cols; ++k) {
+    for (int k = 0; k < num_cols/2 + 1; ++k) {
         // sinf = sycl::sin(angle * k);
         // cosf = sycl::cos(angle * k);
         // sinf = sycl::sincos(angle * k, sycl::make_ptr<float, sycl::access::address_space::private_space>(&cosf));
         // twiddle = std::complex<float>(cosf, sinf);
         TWIDDLE();
-        if (k < (num_cols/2+1)) {
-            sum = sum + din[j * (num_cols/2+1) + k] * twiddle;
-        } else {
-            sum = sum + std::conj(din[((num_rows-j)%num_rows) * (num_cols/2+1) + ((num_cols-k)%num_cols)]) * twiddle;
-        }
+        sum = sum + din[j * (num_cols/2+1) + k] * twiddle;
+    }
+#pragma unroll
+    for (int k = num_cols/2 + 1; k < num_cols; ++k) {
+        TWIDDLE();
+        sum = sum + std::conj(din[((num_rows-j)%num_rows) * (num_cols/2+1) + ((num_cols-k)%num_cols)]) * twiddle;
     }
 
     dout[i * num_rows + j] = sum;

tsne/SYCL/src/utils/debug_utils.dp.cpp

@@ -37,6 +37,7 @@
  * 
  */
 
+#include <complex>
 #include <sycl/sycl.hpp>
 #include "include/utils/debug_utils.h"
 

tsne/SYCL/src/utils/matrix_broadcast_utils.dp.cpp

@@ -37,6 +37,7 @@
  * Copyright (c) 2018, Regents of the University of California
  */
 
+#include <complex>
 #include <sycl/sycl.hpp>
 #include "include/utils/matrix_broadcast_utils.h"