Adds a CUDA "sleep" kernel

Adds a CUDA "sleep" kernel which spins for the given number of
iterations. This is useful for testing correct synchronization with
streams.

init.c

@@ -2,6 +2,7 @@
 #include "luaT.h"
 #include "THCGeneral.h"
 #include "THCCachingAllocator.h"
+#include "THCSleep.h"
 #include "THCTensorRandom.h"
 #include "THCHalf.h" // for CUDA_HALF_TENSOR
 
@@ -938,6 +939,15 @@ static int cutorch_hasFastHalfInstructions(lua_State *L) {
   return 1;
 }
 
+static int cutorch_sleep(lua_State *L) {
+  THCState *state = cutorch_getstate(L);
+  if (!luaT_checklong(L, 1)) {
+      THError("expected number 'cycles'");
+  }
+  THC_sleep(state, luaT_tolong(L, 1));
+  return 0;
+}
+
 static const struct luaL_Reg cutorch_stuff__ [] = {
   {"synchronize", cutorch_synchronize},
   {"synchronizeAll", cutorch_synchronizeAll},
@@ -972,6 +982,7 @@ static const struct luaL_Reg cutorch_stuff__ [] = {
   {"initialSeed", cutorch_initialSeed},
   {"manualSeed", cutorch_manualSeed},
   {"manualSeedAll", cutorch_manualSeedAll},
+  {"_sleep", cutorch_sleep},
   {"getRNGState", cutorch_getRNGState},
   {"setRNGState", cutorch_setRNGState},
   {"getState", cutorch_getState},

lib/THC/CMakeLists.txt

@@ -137,6 +137,7 @@ SET(src
 SET(src-cuda
   THCReduceApplyUtils.cu
   THCBlas.cu
+  THCSleep.cu
   THCStorage.cu
   THCStorageCopy.cu
   THCTensor.cu
@@ -199,6 +200,7 @@ INSTALL(FILES
           THC.h
           ${CMAKE_CURRENT_BINARY_DIR}/THCGeneral.h
           THCBlas.h
+          THCSleep.h
           THCStorage.h
           THCStorageCopy.h
           THCStream.h

lib/THC/THC.h

@@ -5,6 +5,7 @@
 #include "THCAllocator.h"
 #include "THCBlas.h"
 #include "THCCachingAllocator.h"
+#include "THCSleep.h"
 #include "THCStorage.h"
 #include "THCStorageCopy.h"
 #include "THCStream.h"

lib/THC/THCSleep.cu

@@ -0,0 +1,21 @@
+#include "THCSleep.h"
+
+
+__global__ void spin_kernel(long long cycles)
+{
+  // see concurrentKernels CUDA sampl
+  long long start_clock = clock64();
+  long long clock_offset = 0;
+  while (clock_offset < cycles)
+  {
+    clock_offset = clock64() - start_clock;
+  }
+}
+
+THC_API void THC_sleep(THCState* state, long long cycles)
+{
+  dim3 grid(1);
+  dim3 block(1);
+  spin_kernel<<<grid, block, 0, THCState_getCurrentStream(state)>>>(cycles);
+  THCudaCheck(cudaGetLastError());
+}

lib/THC/THCSleep.h

@@ -0,0 +1,10 @@
+#ifndef THC_SPIN_INC
+#define THC_SPIN_INC
+
+#include "THCGeneral.h"
+#include <time.h>
+
+// enqueues a kernel that spins for the specified number of cycles
+THC_API void THC_sleep(THCState* state, long long cycles);
+
+#endif