Reduce register pressure

Author: vmarkovtsev

kernel.cu

@@ -42,8 +42,8 @@ __global__ void weighted_minhash_cuda(
   const uint32_t sample_offset = sample_index * sample_delta;
   const uint32_t samples = blockDim.x * sample_delta;
   extern __shared__ float shmem[];
-  float *lnmins = &shmem[(threadIdx.y * blockDim.x + sample_index) * 3 * sample_delta];
-  uint2 *dtmins = reinterpret_cast<uint2 *>(lnmins + sample_delta);
+  float *volatile lnmins = &shmem[(threadIdx.y * blockDim.x + sample_index) * 3 * sample_delta];
+  uint2 *volatile dtmins = reinterpret_cast<uint2 *>(lnmins + sample_delta);
   int32_t row = -1;
   for (uint32_t index = 0, border = 0;; index++) {
     if (index >= border) {