Implement frozen_logical_data::get_access_mode to query the access mo…

…de used to (#3646)

freeze a logical data (and avoid storing this information separately)

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

@@ -136,6 +136,11 @@ private:
       ld.set_automatic_unfreeze(fake_task, flag);
     }
 
+    access_mode get_access_mode() const
+    {
+      return m;
+    }
+
   private:
     logical_data<T> ld;
     access_mode m;
@@ -199,6 +204,12 @@ public:
     return *this;
   }
 
+  access_mode get_access_mode() const
+  {
+    assert(pimpl);
+    return pimpl->get_access_mode();
+  }
+
 private:
   ::std::shared_ptr<impl> pimpl = nullptr;
 };

cudax/test/stf/freeze/freeze_rw.cu

@@ -56,6 +56,9 @@ int main()
   for (int k = 0; k < 4; k++)
   {
     auto fx = ctx.freeze(lX, access_mode::rw, data_place::current_device());
+
+    _CCCL_ASSERT(fx.get_access_mode() == access_mode::rw, "invalid access mode");
+
     auto dX = fx.get(data_place::current_device(), stream);
     mult<<<8, 4, 0, stream>>>(dX, 4);
     fx.unfreeze(stream);

docs/cudax/stf.rst

@@ -1853,6 +1853,9 @@ stream passed to ``freeze`` (for example, by using a blocking call such as
 depend on the completion of the work in the streams used for any preceding
 ``freeze`` and ``get`` calls.
 
+It is possible to retrieve the access mode used to freeze a logical data with
+the ``get_access_mode()`` method of the ``frozen_logical_data`` object.
+
 Logical token
 ^^^^^^^^^^^^^