unified accessor subscript operator

celerity · Jun 5, 2023 · cc51408 · cc51408
1 parent 1aa701e
commit cc51408
Showing 1 changed file with 9 additions and 26 deletions.
diff --git a/include/accessor.h b/include/accessor.h
@@ -228,7 +228,7 @@ class accessor<DataT, Dims, Mode, target::device> : public detail::accessor_base
 	accessor& operator=(accessor&&) noexcept = default;
 
 
-#if !defined(SYCL_DEVICE_ONLY)
+#if !defined(__SYCL_DEVICE_ONLY__)
 	accessor(const accessor& other) { copy_and_hydrate(other); }
 
 	accessor& operator=(const accessor& other) {
@@ -255,35 +255,18 @@ class accessor<DataT, Dims, Mode, target::device> : public detail::accessor_base
 	}
 
 	template <access_mode M = Mode>
-	inline std::enable_if_t<detail::access::mode_traits::is_producer(M), DataT&> operator[](const id<Dims>& index) const {
-		auto linear_index = get_linear_offset(index);
+	inline std::conditional_t<detail::access::mode_traits::is_producer(M), DataT&, const DataT&> operator[](const id<Dims>& index) const {
 #if CELERITY_DETAIL_ACCESSOR_BOUNDARY_CHECK
-		auto actual_index = index - m_backing_buffer_offset;
-		if((actual_index >= m_backing_buffer_range) != detail::id_cast<Dims>(id<3>(false, false, false))) {
+		if(((index - m_backing_buffer_offset) >= m_backing_buffer_range) != detail::id_cast<Dims>(id<3>(false, false, false))) {
 			for(int d = 0; d < Dims; ++d) {
 				sycl::atomic_ref<size_t, sycl::memory_order::relaxed, sycl::memory_scope::device>{m_attempted_index[0][d]}.fetch_min(index[d]);
 				sycl::atomic_ref<size_t, sycl::memory_order::relaxed, sycl::memory_scope::device>{m_attempted_index[1][d]}.fetch_max(index[d] + 1);
 			}
-			linear_index = 0;
+			// TODO: make sure we always allocate at least one element in the backing buffer, if CELERITY_DETAIL_ACCESSOR_BOUNDARY_CHECK is enabled
+			return m_device_ptr[0];
 		}
 #endif
-		return m_device_ptr[linear_index];
-	}
-
-	template <access_mode M = Mode>
-	inline std::enable_if_t<detail::access::mode_traits::is_pure_consumer(M), const DataT&> operator[](const id<Dims>& index) const {
-		auto linear_index = get_linear_offset(index);
-#if CELERITY_DETAIL_ACCESSOR_BOUNDARY_CHECK
-		auto actual_index = index - m_backing_buffer_offset;
-		if((actual_index >= m_backing_buffer_range) != detail::id_cast<Dims>(id<3>(false, false, false))) {
-			for(int d = 0; d < Dims; ++d) {
-				sycl::atomic_ref<size_t, sycl::memory_order::relaxed, sycl::memory_scope::device>{m_attempted_index[0][d]}.fetch_min(index[d]);
-				sycl::atomic_ref<size_t, sycl::memory_order::relaxed, sycl::memory_scope::device>{m_attempted_index[1][d]}.fetch_max(index[d] + 1);
-			}
-			linear_index = 0;
-		}
-#endif
-		return m_device_ptr[linear_index];
+		return m_device_ptr[get_linear_offset(index)];
 	}
 
 	template <int D = Dims>
@@ -353,7 +336,7 @@ class accessor<DataT, Dims, Mode, target::device> : public detail::accessor_base
 	// Constructor for tests, called through accessor_testspy.
 	accessor(DataT* const ptr, const id<Dims>& backing_buffer_offset, const range<Dims>& backing_buffer_range)
 	    : m_device_ptr(ptr), m_backing_buffer_offset(backing_buffer_offset), m_backing_buffer_range(backing_buffer_range) {
-#if defined(SYCL_DEVICE_ONLY)
+#if defined(__SYCL_DEVICE_ONLY__)
 #if CELERITY_WORKAROUND_HIPSYCL // hipSYCL does not yet implement is_device_copyable_v
 		static_assert(std::is_trivially_copyable_v<accessor>);
 #else
@@ -374,7 +357,7 @@ class accessor<DataT, Dims, Mode, target::device> : public detail::accessor_base
 		m_attempted_index = other.m_attempted_index;
 #endif
 
-#if !defined(__SYCL_DEVICE_ONLY__)
+#if !defined(____SYCL_DEVICE_ONLY____)
 		if(detail::is_embedded_hydration_id(m_device_ptr)) {
 			if(detail::cgf_diagnostics::is_available() && detail::cgf_diagnostics::get_instance().is_checking()) {
 				detail::cgf_diagnostics::get_instance().register_accessor(detail::extract_hydration_id(m_device_ptr), target::device);
@@ -732,7 +715,7 @@ class local_accessor {
 	template <int D = Dims, typename = std::enable_if_t<D == 0>>
 	local_accessor(handler& cgh) : local_accessor(range<0>(), cgh) {}
 
-#if !defined(__SYCL_DEVICE_ONLY__)
+#if !defined(____SYCL_DEVICE_ONLY____)
 	local_accessor(const range<Dims>& allocation_size, handler& cgh) : m_sycl_acc{}, m_allocation_size(allocation_size) {}
 
 	local_accessor(const local_accessor& other)