diff --git a/CUDA/cudapi.cu b/CUDA/cudapi.cu
old mode 100755
new mode 100644
index 2586924fe8b629ab7f9ba8b44e961b18ee6ca969..b755b845b2f2d0466db1e9087724f86058ee0d0d
--- a/CUDA/cudapi.cu
+++ b/CUDA/cudapi.cu
@@ -6,8 +6,8 @@
 //
 #include <iostream>
 #include <chrono>
-#include <cstdint>
 #include <string>
+using namespace std;
 uint64_t blocks = 1024;
 uint64_t threads = 1024;
 uint64_t nloop = 1000000;
@@ -18,16 +18,6 @@ void cudaCheck(string msg) {
    if (cudaSuccess != err)
    cerr << msg << ": " << cudaGetErrorString(err) << endl;
    }
-void reduce(double *arr) {
-   uint64_t len = npts >> 1;
-   while (1) {
-      reduce_sum<<<blocks,threads>>>(arr,len);
-      cudaCheck("reduce");
-      len = len >> 1;
-      if (len == 0)
-         return;
-      }
-   }
 __global__ void init(double *arr,uint64_t nloop) {
    uint64_t i = blockIdx.x*blockDim.x+threadIdx.x;
    uint64_t start = nloop*i+1;
@@ -41,6 +31,16 @@ __global__ void reduce_sum(double *arr,uint64_t len) {
    if (i < len)
       arr[i] += arr[i+len];
    }
+void reduce(double *arr) {
+   uint64_t len = npts >> 1;
+   while (1) {
+      reduce_sum<<<blocks,threads>>>(arr,len);
+      cudaCheck("reduce");
+      len = len >> 1;
+      if (len == 0)
+         return;
+      }
+   }
 int main(void) {
    double harr[1],*darr;
    cudaMalloc(&darr,npts*sizeof(double));