diff --git a/src/primitives.h b/src/primitives.h
index 4d7b86b..8f8fb69 100644
--- a/src/primitives.h
+++ b/src/primitives.h
@@ -177,29 +177,29 @@ class Primitives {
   template <typename... SYNC_Ts>
   static __device__ __forceinline__ void
   Copy(const T* src, T* dst,
-      int len, int step, SYNC_Ts... flags) {
-    GenericOp(src, nullptr, dst, nullptr, len, step, flags...);
+      int len, int maxOffset, int step, SYNC_Ts... flags) {
+    GenericOp(src, nullptr, dst, nullptr, len, maxOffset, step, flags...);
   }
 
   template <typename... SYNC_Ts>
   static __device__ __forceinline__ void
   DoubleCopy(const T* src, T* dst1, T* dst2,
-      int len, int step, SYNC_Ts... flags) {
-    GenericOp(src, nullptr, dst1, dst2, len, step, flags...);
+      int len, int maxOffset, int step, SYNC_Ts... flags) {
+    GenericOp(src, nullptr, dst1, dst2, len, maxOffset, step, flags...);
   }
 
   template <typename... SYNC_Ts>
   static __device__ __forceinline__ void
   Reduce(const T* src1, const T* src2, T* dst,
-      int len, int step, SYNC_Ts... flags) {
-    GenericOp(src1, src2, dst, nullptr, len, step, flags...);
+      int len, int maxOffset, int step, SYNC_Ts... flags) {
+    GenericOp(src1, src2, dst, nullptr, len, maxOffset, step, flags...);
   }
 
   template <typename... SYNC_Ts>
   static __device__ __forceinline__ void
   ReduceCopy(const T* src1, const T* src2, T* dst1, T* dst2,
-      int len, int step, SYNC_Ts... flags) {
-    GenericOp(src1, src2, dst1, dst2, len, step, flags...);
+      int len, int maxOffset, int step, SYNC_Ts... flags) {
+    GenericOp(src1, src2, dst1, dst2, len, maxOffset, step, flags...);
   }
 };