diff --git a/src/include/mscclpp.h b/src/include/mscclpp.h
index 7ff606d6..48544911 100644
--- a/src/include/mscclpp.h
+++ b/src/include/mscclpp.h
@@ -139,8 +139,8 @@ struct mscclppDevConn
   __forceinline__ __device__ void putDirect(uint64_t dstDataOffset, uint64_t srcDataOffset, uint64_t dataSize,
                                             uint32_t threadId, uint32_t numThreads)
   {
-    uint64_t* src = (uint64_t*)localBuff + srcDataOffset;
-    uint64_t* dst = (uint64_t*)remoteBuff + dstDataOffset;
+    uint64_t* src = (uint64_t*)((char*)localBuff + srcDataOffset);
+    uint64_t* dst = (uint64_t*)((char*)remoteBuff + dstDataOffset);
     // assume the memory is aligned to 8 bytes
     size_t nElem =
       dataSize % sizeof(uint64_t) ? (dataSize + sizeof(uint64_t)) / sizeof(uint64_t) : dataSize / sizeof(uint64_t);
diff --git a/tests/common.cu b/tests/common.cu
index 1d424452..c9cecadf 100644
--- a/tests/common.cu
+++ b/tests/common.cu
@@ -224,6 +224,11 @@ testResult_t CheckData(struct testArgs* args, int in_place, int64_t* wrongElts)
   CUDACHECK(cudaMemcpy(dataHostRecv, args->recvbuff, args->expectedBytes, cudaMemcpyDeviceToHost));
   CUDACHECK(cudaMemcpy(dataHostExpected, args->expected, args->expectedBytes, cudaMemcpyDeviceToHost));
 
+  for (size_t i = 0; i < count; i++) {
+    if (dataHostRecv[i] != dataHostExpected[i]) {
+      *wrongElts += 1;
+    }
+  }
   if (args->reportErrors && *wrongElts) {
     (args->error)++;
   }
diff --git a/tests/ring_send_recv_test.cu b/tests/ring_send_recv_test.cu
index 886aeae9..86a4339f 100644
--- a/tests/ring_send_recv_test.cu
+++ b/tests/ring_send_recv_test.cu
@@ -129,6 +129,6 @@ testResult_t RingSendRecvRunTest(struct testArgs* args)
   return testSuccess;
 }
 
-struct testEngine ringSendRecvTestEngine = {RingSendRecvGetBuffSize, RingSendRecvRunTest};
+struct testEngine ringSendRecvTestEngine = {RingSendRecvGetBuffSize, RingSendRecvRunTest, nullptr, nullptr};
 
 #pragma weak mscclppTestEngine = ringSendRecvTestEngine