 #ifndef _REDUCE_KERNEL_H_
 #define _REDUCE_KERNEL_H_
 
-#include <device_functions.h>
+#ifdef __clang__
+#include "cuda.h"
+#endif
+
 #include <cuComplex.h>
 
-/*
-    Parallel sum reduction using shared memory
+
+/*
+    Parallel reduction using shared memory
     - takes log(n) steps for n input elements
     - uses n/2 threads
     - only works for power-of-2 arrays

        if (amLast)
        {
            int i = tid;
            double mySum = 0;

            while (i < gridDim.x)
            {
                mySum += g_odata[i];
                i += blockSize;
            }

            reduceBlock<blockSize>(smem, mySum, tid);

            if (tid==0)
            {

    {
        const unsigned int tid = threadIdx.x;
        __shared__ bool amLast;
        extern typ __shared__ smem[];

        // wait until all outstanding memory instructions in this thread are finished
        __threadfence();
