Update register keyword handling for C++17 (#100)

* Update register keyword handling for C++17

The keyword 'register' for storage class is removed in C++17,
so keeping it active for only c++14 and lower.

* Updates to the code

apex/contrib/csrc/peer_memory/peer_memory_cuda.cu

@@ -27,6 +27,13 @@ namespace cg = cooperative_groups;
   }                                                 \
 } while(0)
 
+// C++17 removes 'register' storage keyword
+#if __cplusplus < 201703L
+#define REGISTER register
+#else
+#define REGISTER
+#endif
+
 namespace {
 
 /* Basic deleter function for from_blob function.
@@ -184,7 +191,7 @@ __device__ void checked_signal(
 	// flush all writes to global memory
 	__threadfence_system();
 	// wait for top or bottom neighbor to clear signal
-	register int r1, r2, r3, r4;
+	REGISTER int r1, r2, r3, r4;
 	if (!(top_zero || btm_zero)) {
 	    bool top_zeroed=false, top_done=false;
 	    bool btm_zeroed=false, btm_done=false;
@@ -308,7 +315,7 @@ __device__ void wait_for(
 {
     bool is_main_thread = (blockIdx.x == 0 && threadIdx.x == 0) ? true : false;
     if (is_main_thread) {
-    	register int r1, r2, r3, r4;
+    	REGISTER int r1, r2, r3, r4;
 	// wait for senders to signal their output is read
 	do {
 #ifdef __HIP_PLATFORM_HCC__
@@ -332,7 +339,7 @@ __device__ void clear_flag(
     cg::this_grid().sync();  // wait for all threads in kernel to finish
     bool is_main_thread = (blockIdx.x == 0 && threadIdx.x == 0) ? true : false;
     if (is_main_thread) {
-	register int r1, r2, r3, r4;
+	REGISTER int r1, r2, r3, r4;
 	r1 = 0;  r2 = 0;  r3 = 0;  r4 = 0;
 #ifdef __HIP_PLATFORM_HCC__
 	__builtin_nontemporal_store(r1, wait_flag);