Update Readme to include Kokkos

README.md

@@ -11,16 +11,17 @@ The top-level `makedeck` directory contains an input deck generation program and
 Each other subdirectory contains a separate implementation:
 
 - [OpenMP](openmp/) for CPUs
-- [SYCL](sycl/) for CPUs
 - [OpenMP target](openmp-target/) for GPUs
-- [OpenCL](opencl/) for GPUs
 - [CUDA](cuda/) for GPUs
+- [OpenCL](opencl/) for GPUs
 - [OpenACC](openacc/) for GPUs
+- [SYCL](sycl/) for CPUs and GPUs
+- [Kokkos](kokkos/) for CPUs and GPUs
 
 ## Building
 
 To build with the default options, type `make` in an implementation directory.
-There are options to choose the compiler used and the architecture targetted.
+There are options to choose the compiler used and the architecture targeted.
 
 Refer to each implementation's README for further build instructions.
 

cuda/bude.c

@@ -94,7 +94,7 @@ int main(int argc, char *argv[])
       // printf ("Maxdiff: %.2f (%.3f vs %.3f)\n", maxdiff, resultsRef[i], resultsCUDA[i]);
     }
 
-    if (i < 8)
+    if (i < 1024)
       printf("%7.2f    vs   %7.2f  (%5.2f%%)\n", resultsRef[i], resultsCUDA[i], 100*diff);
   }
   printf("\nLargest difference was %.3f%%\n\n", maxdiff*100);

cuda/bude_cuda.cu

@@ -84,6 +84,8 @@ void runCUDA(float* results)
   cudaDeviceSynchronize();
 
   double start = getTimestamp();
+  printf("<<%d,%d>> -> %d\n",global, local,params.nposes);
+
 
   for(int ii = 0; ii < params.iterations; ++ii)
   {
@@ -133,10 +135,12 @@ __device__ void compute_transformation_matrix(const float transform_0,
   transform[0].y = sx*sy*cz - cx*sz;
   transform[0].z = cx*sy*cz + sx*sz;
   transform[0].w = transform_3;
+
   transform[1].x = cy*sz;
   transform[1].y = sx*sy*sz + cx*cz;
   transform[1].z = cx*sy*sz - sx*cz;
   transform[1].w = transform_4;
+
   transform[2].x = -sy;
   transform[2].y = sx*cy;
   transform[2].z = cx*cy;
@@ -191,12 +195,18 @@ __global__ void fasten_main(const int natlig,
         transforms_5[index],
         transform[i]);
     etot[i] = ZERO;
+    // etotals[ix] = transform[i][0].x;
   }
 
 #ifdef USE_SHARED
   __syncthreads();
 #endif
 
+    // etotals[ix] =  ix ;//  transform[3][1].y; // (transforms_0[ix + 1 * lsz]) ;// transform[0][0].x ;// transform[0][0].x;
+    // etot[1] = transforms_1[0] ;// transform[0][0].y;
+    // etot[2] = transforms_2[0] ;// transform[0][0].z;
+    // etot[3] = transforms_3[0] ;// transform[0][0].w;
+
   // Loop over ligand atoms
   int il = 0;
   do
@@ -221,6 +231,7 @@ __global__ void fasten_main(const int natlig,
         linitpos.y*transform[i][2].y + linitpos.z*transform[i][2].z;
     }
 
+
     // Loop over protein atoms
     int ip = 0;
     do
@@ -290,6 +301,7 @@ __global__ void fasten_main(const int natlig,
       etotals[td_base+i*blockDim.x] = etot[i]*HALF;
     }
   }
+  // etotals[ix] = transform[0][0].x;
 } //end of fasten_main
 
 

kokkos/CMakeLists.txt

@@ -1,7 +1,7 @@
 cmake_minimum_required(VERSION 3.12 FATAL_ERROR)
 
 project(bude_kokkos)
-SET(CMAKE_VERBOSE_MAKEFILE OFF)
+set(CMAKE_VERBOSE_MAKEFILE ON)
 
 set(CMAKE_CXX_STANDARD 14)
 set(CMAKE_CXX_EXTENSIONS OFF)

kokkos/src/bude.h

@@ -11,7 +11,7 @@
 #endif
 
 #ifndef WG_SIZE
-#define WG_SIZE 128
+#define WG_SIZE 256
 #endif
 
 #define DEFAULT_ITERS  8