[STF] Rename the redux access mode into relaxed (NVIDIA#2776)

* Rename the redux access mode into relaxed

* change redux to relaxed in the documentation

* clang-format

Author: caugonnet

cudax/include/cuda/experimental/__stf/internal/acquire_release.cuh

@@ -137,9 +137,9 @@ inline event_list task::acquire(backend_ctx_untyped& ctx)
     const data_place& dplace = it->get_dplace() == data_place::affine ? get_affine_data_place() : it->get_dplace();
 
     const instance_id_t instance_id =
-      mode == access_mode::redux ? d.find_unused_instance_id(dplace) : d.find_instance_id(dplace);
+      mode == access_mode::relaxed ? d.find_unused_instance_id(dplace) : d.find_instance_id(dplace);
 
-    if (mode == access_mode::redux)
+    if (mode == access_mode::relaxed)
     {
       d.get_data_instance(instance_id).set_redux_op(it->get_redux_op());
     }
@@ -187,7 +187,7 @@ inline event_list task::acquire(backend_ctx_untyped& ctx)
   {
     logical_data_untyped d = e.get_data();
 
-    if (e.get_access_mode() == access_mode::redux)
+    if (e.get_access_mode() == access_mode::relaxed)
     {
       // Save the last task accessing the instance in with a relaxed coherency mode
       d.get_data_instance(e.get_instance_id()).set_last_task_relaxed(*this);

cudax/include/cuda/experimental/__stf/internal/constants.cuh

@@ -35,11 +35,11 @@ namespace cuda::experimental::stf
  */
 enum class access_mode : unsigned int
 {
-  none  = 0,
-  read  = 1,
-  write = 2,
-  rw    = 3, // READ + WRITE
-  redux = 4, /* operator ? */
+  none    = 0,
+  read    = 1,
+  write   = 2,
+  rw      = 3, // READ + WRITE
+  relaxed = 4, /* operator ? */
 };
 
 /**
@@ -50,8 +50,8 @@ inline access_mode operator|(access_mode lhs, access_mode rhs)
 {
   assert(as_underlying(lhs) < 16);
   assert(as_underlying(rhs) < 16);
-  EXPECT(lhs != access_mode::redux);
-  EXPECT(rhs != access_mode::redux);
+  EXPECT(lhs != access_mode::relaxed);
+  EXPECT(rhs != access_mode::relaxed);
   return access_mode(as_underlying(lhs) | as_underlying(rhs));
 }
 
@@ -75,8 +75,8 @@ inline const char* access_mode_string(access_mode mode)
       return "rw";
     case access_mode::write:
       return "write";
-    case access_mode::redux:
-      return "redux"; // op ?
+    case access_mode::relaxed:
+      return "relaxed"; // op ?
     default:
       assert(false);
       abort();

cudax/include/cuda/experimental/__stf/internal/logical_data.cuh

@@ -242,7 +242,7 @@ public:
 
   reserved::logical_data_state state;
 
-  // For temporary or redux accesses, we need to be able to find an available entry
+  // For temporary or relaxed accesses, we need to be able to find an available entry
   ::std::vector<data_instance> used_instances;
 
   // A string useful for debugging purpose
@@ -917,7 +917,7 @@ public:
         break;
       }
 
-      case access_mode::redux:
+      case access_mode::relaxed:
         current_instance.set_msir(reserved::msir_state_id::reduction);
         break;
       default:
@@ -1238,9 +1238,9 @@ public:
     return task_dep_untyped(*this, access_mode::rw, mv(dp));
   }
 
-  task_dep_untyped redux(::std::shared_ptr<reduction_operator_base> op, data_place dp = data_place::affine)
+  task_dep_untyped relaxed(::std::shared_ptr<reduction_operator_base> op, data_place dp = data_place::affine)
   {
-    return task_dep_untyped(*this, access_mode::redux, mv(dp), op);
+    return task_dep_untyped(*this, access_mode::relaxed, mv(dp), op);
   }
 
   ///@}
@@ -1732,7 +1732,7 @@ inline void reserved::logical_data_untyped_impl::erase()
     assert(ref_id != instance_id_t::invalid);
 
     // Get the state in which we store previous writer, readers, ...
-    if (h_state.current_mode == access_mode::redux)
+    if (h_state.current_mode == access_mode::relaxed)
     {
       // Reconstruction of the data on the reference data place needed
 
@@ -1863,14 +1863,14 @@ inline event_list enforce_stf_deps_before(
   auto& dot                 = *bctx.get_dot();
   const bool dot_is_tracing = dot.is_tracing();
 
-  if (mode == access_mode::redux)
+  if (mode == access_mode::relaxed)
   {
     // A reduction only needs to wait for previous accesses on the data instance
-    ctx_.current_mode = access_mode::redux;
+    ctx_.current_mode = access_mode::relaxed;
 
     if (dot_is_tracing)
     {
-      // Add this task to the list of task accessing the logical data in redux mode
+      // Add this task to the list of task accessing the logical data in relaxed mode
       // We only store its id since this is used for dot
       ctx_.pending_redux_id.push_back(task.get_unique_id());
     }
@@ -1882,13 +1882,13 @@ inline event_list enforce_stf_deps_before(
   }
 
   // This is not a reduction, but perhaps we need to reconstruct the data first?
-  if (ctx_.current_mode == access_mode::redux)
+  if (ctx_.current_mode == access_mode::relaxed)
   {
     assert(eplace.has_value());
     if (dot_is_tracing)
     {
       // Add a dependency between previous tasks accessing the handle
-      // in redux mode, and this task which forces its
+      // in relaxed mode, and this task which forces its
       // reconstruction.
       for (const int redux_task_id : ctx_.pending_redux_id)
       {
@@ -1998,7 +1998,7 @@ inline event_list enforce_stf_deps_before(
 template <typename task_type>
 inline void enforce_stf_deps_after(logical_data_untyped& handle, const task_type& task, const access_mode mode)
 {
-  if (mode == access_mode::redux)
+  if (mode == access_mode::relaxed)
   {
     // no further action is required
     return;
@@ -2295,9 +2295,9 @@ public:
   }
 
   template <typename... Pack>
-  task_dep<T> redux(Pack&&... pack)
+  task_dep<T> relaxed(Pack&&... pack)
   {
-    return task_dep<T>(*this, access_mode::redux, ::std::forward<Pack>(pack)...);
+    return task_dep<T>(*this, access_mode::relaxed, ::std::forward<Pack>(pack)...);
   }
   ///@}
 };

cudax/include/cuda/experimental/__stf/internal/task.cuh

@@ -448,7 +448,7 @@ void dep_allocate(
     }
 
     // After allocating a reduction instance, we need to initialize it
-    if (mode == access_mode::redux)
+    if (mode == access_mode::relaxed)
     {
       assert(eplace.has_value());
       // We have just allocated a new piece of data to perform

cudax/test/stf/examples/07-cholesky-redux.cu

@@ -263,7 +263,7 @@ void DGEMM(
 
   // If beta == 1.0 (we assume this is exactly 1.0), then this operation is
   // an accumulation with the add operator
-  auto dep_c = (beta == 1.0) ? C.handle(C_row, C_col).redux(redux_op) : C.handle(C_row, C_col).rw();
+  auto dep_c = (beta == 1.0) ? C.handle(C_row, C_col).relaxed(redux_op) : C.handle(C_row, C_col).rw();
   auto t     = ctx.task(exec_place::device(A.get_preferred_devid(C_row, C_col)),
                     A.handle(A_row, A_col).read(),
                     B.handle(B_row, B_col).read(),

cudax/test/stf/hashtable/fusion_reduction.cu

@@ -98,7 +98,7 @@ int main()
 
   for (size_t dev_id = 0; dev_id < 4; dev_id++)
   {
-    ctx.task(h_handle.redux(fusion_op))->*[&](auto stream, auto h) {
+    ctx.task(h_handle.relaxed(fusion_op))->*[&](auto stream, auto h) {
       EXPECT(h.get_capacity() == 2048);
       fill_table<<<32, 32, 0, stream>>>(dev_id, 10, h);
     };

cudax/test/stf/reductions/many_inc.cu

@@ -35,7 +35,7 @@ int main()
   for (int i = 0; i < K; i++)
   {
     // Increment the variable by 1
-    ctx.task(exec_place::device(i % ndevs), handle.redux(redux_op))->*[](auto stream, auto s) {
+    ctx.task(exec_place::device(i % ndevs), handle.relaxed(redux_op))->*[](auto stream, auto s) {
       add<<<1, 1, 0, stream>>>(s.data_handle());
     };
   }

cudax/test/stf/reductions/redux_test.cu

@@ -25,7 +25,7 @@ int main()
   int a         = 17;
   auto handle   = ctx.logical_data(make_slice(&a, 1));
   auto redux_op = std::make_shared<slice_reduction_op_sum<int>>();
-  ctx.task(handle.redux(redux_op))->*[](auto stream, auto s) {
+  ctx.task(handle.relaxed(redux_op))->*[](auto stream, auto s) {
     add<<<1, 1, 0, stream>>>(s.data_handle(), 42);
   };
 

cudax/test/stf/reductions/redux_test2.cu

@@ -58,7 +58,7 @@ int main()
   };
 
   // REDUX dev1 (18 + 42)
-  ctx.task(exec_place::device(1), handle.redux(redux_op))->*[](auto stream, auto s) {
+  ctx.task(exec_place::device(1), handle.relaxed(redux_op))->*[](auto stream, auto s) {
     add<<<1, 1, 0, stream>>>(s.data_handle(), 42);
   };
 

cudax/test/stf/reductions/slice2d_reduction.cu

@@ -38,7 +38,7 @@ int main()
 
   auto redux_op = std::make_shared<slice_reduction_op_sum<int, 2>>();
 
-  ctx.task(handle.redux(redux_op))->*[](auto stream, auto s) {
+  ctx.task(handle.relaxed(redux_op))->*[](auto stream, auto s) {
     add<<<32, 32, 0, stream>>>(s, 42);
   };
 

cudax/test/stf/reductions/slice_custom_op.cu

@@ -51,12 +51,12 @@ int main()
   auto op = std::make_shared<slice_reduction_op<bool, 1, OR_op>>();
 
   // C |= A
-  ctx.task(lC.redux(op), lA.read())->*[](auto stream, auto sC, auto sA) {
+  ctx.task(lC.relaxed(op), lA.read())->*[](auto stream, auto sC, auto sA) {
     cudaMemcpyAsync(sC.data_handle(), sA.data_handle(), sA.extent(0) * sizeof(bool), cudaMemcpyDeviceToDevice, stream);
   };
 
   // C |= B
-  ctx.task(lC.redux(op), lB.read())->*[](auto stream, auto sC, auto sB) {
+  ctx.task(lC.relaxed(op), lB.read())->*[](auto stream, auto sC, auto sB) {
     cudaMemcpyAsync(sC.data_handle(), sB.data_handle(), sB.extent(0) * sizeof(bool), cudaMemcpyDeviceToDevice, stream);
   };
 

cudax/test/stf/reductions/successive_reductions.cu

@@ -50,7 +50,7 @@ int main()
     // We add i (total = N(N-1)/2 + initial_value)
     for (int i = 0; i < N; i++)
     {
-      ctx.task(var_handle.redux(redux_op))->*[=](cudaStream_t stream, auto d_var) {
+      ctx.task(var_handle.relaxed(redux_op))->*[=](cudaStream_t stream, auto d_var) {
         add_val<<<1, 1, 0, stream>>>(d_var.data_handle(), i);
         cuda_safe_call(cudaGetLastError());
       };

cudax/test/stf/reductions/successive_reductions_pfor.cu

@@ -29,7 +29,7 @@ int main()
     // We add i (total = N(N-1)/2 + initial_value)
     for (int i = 0; i < N; i++)
     {
-      ctx.parallel_for(var_handle.shape(), var_handle.redux(op))->*[=] _CCCL_DEVICE(size_t ind, auto d_var) {
+      ctx.parallel_for(var_handle.shape(), var_handle.relaxed(op))->*[=] _CCCL_DEVICE(size_t ind, auto d_var) {
         atomicAdd(d_var.data_handle(), i);
       };
     }

cudax/test/stf/reductions/sum.cu

@@ -92,7 +92,7 @@ int main()
   // We add i (total = N(N-1)/2 + initial_value)
   for (int i = 0; i < N; i++)
   {
-    ctx.task(var_handle.redux(redux_op))->*[&](cudaStream_t stream, auto d_var) {
+    ctx.task(var_handle.relaxed(redux_op))->*[&](cudaStream_t stream, auto d_var) {
       add_val<<<1, 1, 0, stream>>>(d_var.data_handle(), i);
     };
   }

cudax/test/stf/reductions/sum_array.cu

@@ -98,9 +98,10 @@ int main()
 
   for (int i = 0; i < N; i++)
   {
-    ctx.task(var_handle.redux(redux_op), array_handles[i].read())->*[](cudaStream_t stream, auto d_var, auto d_array_i) {
-      add<<<1, 1, 0, stream>>>(d_array_i.data_handle(), d_var.data_handle());
-    };
+    ctx.task(var_handle.relaxed(redux_op), array_handles[i].read())
+        ->*[](cudaStream_t stream, auto d_var, auto d_array_i) {
+              add<<<1, 1, 0, stream>>>(d_array_i.data_handle(), d_var.data_handle());
+            };
   }
 
   // Force the reconstruction of data on the device, so that no transfers are

cudax/test/stf/reductions/sum_multiple_places.cu

@@ -41,13 +41,13 @@ int main()
     // device
     for (int d = 0; d < ndevs; d++)
     {
-      ctx.task(exec_place::device(d), var_handle.redux(redux_op))->*[=](cudaStream_t s, auto var) {
+      ctx.task(exec_place::device(d), var_handle.relaxed(redux_op))->*[=](cudaStream_t s, auto var) {
         add_val<int><<<1, 1, 0, s>>>(var, i);
       };
     }
 
     // host
-    ctx.host_launch(var_handle.redux(redux_op))->*[=](auto var) {
+    ctx.host_launch(var_handle.relaxed(redux_op))->*[=](auto var) {
       var(0) += i;
     };
   }

cudax/test/stf/reductions/sum_multiple_places_no_refvalue.cu

@@ -94,13 +94,13 @@ int main()
     // device
     for (int d = 0; d < ndevs; d++)
     {
-      ctx.task(exec_place::device(d), var_handle.redux(redux_op))->*[&](cudaStream_t s, auto var) {
+      ctx.task(exec_place::device(d), var_handle.relaxed(redux_op))->*[&](cudaStream_t s, auto var) {
         add_val<int><<<1, 1, 0, s>>>(var.data_handle(), i);
       };
     }
 
     // host
-    ctx.task(exec_place::host, var_handle.redux(redux_op))->*[&](cudaStream_t s, auto var) {
+    ctx.task(exec_place::host, var_handle.relaxed(redux_op))->*[&](cudaStream_t s, auto var) {
       cuda_safe_call(cudaStreamSynchronize(s));
       *var.data_handle() += i;
     };

cudax/test/stf/reductions/write_back_after_redux.cu

@@ -49,13 +49,13 @@ int main()
     // device
     for (int d = 0; d < ndevs; d++)
     {
-      ctx.task(exec_place::device(d), var_handle.redux(redux_op))->*[=](cudaStream_t s, auto var) {
+      ctx.task(exec_place::device(d), var_handle.relaxed(redux_op))->*[=](cudaStream_t s, auto var) {
         add_val<int><<<1, 1, 0, s>>>(var, i);
       };
     }
 
     // host
-    ctx.host_launch(var_handle.redux(redux_op))->*[=](auto var) {
+    ctx.host_launch(var_handle.relaxed(redux_op))->*[=](auto var) {
       var(0) += i;
     };
   }

docs/cudax/stf.rst

@@ -558,7 +558,7 @@ write-only access (using the ``write()`` member of ``lX``). A write-only
 access will indeed allocate ``lX`` at the appropriate location, but it
 will not try to load a valid copy of it prior to executing the task.
 
-Using other access modes such as ``read()``, ``redux()`` or ``rw()``
+Using other access modes such as ``read()``, ``relaxed()`` or ``rw()``
 that attempt to provide a valid instance will result in an error.
 
 Similarly, it is possible to define a logical data from a slice shapes