                               sb1[5 * threadIdx.x + 1],
                               sb1[5 * threadIdx.x + 2],
                               sb1[5 * threadIdx.x + 3]);
+#if 1
+        // Streaming store
+        asm volatile("st.global.cs.v4.u32 [%0], {%1,%2,%3,%4};"
+                     :: PTR_REG((uint4*)hp2_d + offset2), "r"(hp.x), "r"(hp.y), "r"(hp.z), "r"(hp.w));
+#else
         ((uint4*)hp2_d)[offset2] = hp;
+#endif
         sb2_d[offset2] = hp.x + hp.y + hp.z + hp.w + sb1[5 * threadIdx.x + 4];
       }
     }


    __shared__ uint32_t sb1[160];

    for (uint32_t i = 0; i < 20; i++) {
      const uint32_t offset0 = offset0_base + 256 * i;
      const uint32_t value = offset0 < n0 ? sb0_d[offset0] : 0;
      const uint32_t warpMask = __ballot(value != 0);
      if (lane == 0) {
        const uint32_t offset1 = offset1_base + 8 * i + warp;
        if (offset1 < n1) {
          hp1_d[offset1] = warpMask;
        }
        sb1[8 * i + warp] = __popc(warpMask);
      }

#include <device_launch_parameters.h>

#include "HP5.h"

#if defined(__LP64__) || defined(_WIN64)
#define PTR_REG "l"
#define PTR_REG_DST "=l"
#else
#define PTR_REG "r"
#define PTR_REG_DST "=r"
#endif

#define HP5_WARP_COUNT 5
#define HP5_WARP_SIZE 32
