From 0d77bb7b3a00f30322e7024b48fb7f353d00e0d3 Mon Sep 17 00:00:00 2001
From: Oliver Rietmann <oliver.rietmann@cern.ch>
Date: Wed, 4 Feb 2026 19:21:17 +0100
Subject: [PATCH 1/2] correct result

---
 GPU/Common/GPUCommonAlgorithm.h               |  18 +-
 GPU/Common/GPUCommonAlgorithmThrust.h         |  14 +-
 GPU/Common/MemLayout.h                        | 341 ++++++++----------
 GPU/GPUTracking/DataTypes/GPUDataTypes.h      |   5 +-
 GPU/GPUTracking/Global/GPUChainTracking.cxx   |  13 +-
 GPU/GPUTracking/Global/GPUChainTracking.h     |   7 +-
 GPU/GPUTracking/Global/GPUChainTrackingIO.cxx |   8 +-
 GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx     |  54 +--
 GPU/GPUTracking/Merger/GPUTPCGMMerger.h       |   7 +-
 .../Merger/GPUTPCGMSectorTrack.cxx            |  16 +-
 GPU/GPUTracking/Merger/GPUTPCGMSectorTrack.h  |   9 +-
 .../SectorTracker/GPUTPCBaseTrackParam.h      |  52 +--
 .../GPUTPCExtrapolationTracking.cxx           |   6 +-
 .../GPUTPCSectorDebugSortKernels.cxx          |   6 +-
 GPU/GPUTracking/SectorTracker/GPUTPCTrack.h   |  63 ++--
 .../SectorTracker/GPUTPCTrackLinearisation.h  |   4 +-
 .../SectorTracker/GPUTPCTrackParam.cxx        |  10 +-
 .../SectorTracker/GPUTPCTrackParam.h          |  35 +-
 .../SectorTracker/GPUTPCTracker.cxx           |   9 +-
 GPU/GPUTracking/SectorTracker/GPUTPCTracker.h |  13 +-
 .../SectorTracker/GPUTPCTrackerDump.cxx       |   2 +-
 .../SectorTracker/GPUTPCTracklet.h            |  31 +-
 .../GPUTPCTrackletConstructor.cxx             |  14 +-
 .../SectorTracker/GPUTPCTrackletConstructor.h |  11 +-
 .../SectorTracker/GPUTPCTrackletSelector.cxx  |   2 +-
 .../display/render/GPUDisplayDraw.cxx         |   2 +-
 26 files changed, 344 insertions(+), 408 deletions(-)

diff --git a/GPU/Common/GPUCommonAlgorithm.h b/GPU/Common/GPUCommonAlgorithm.h
index db57e7ec06d4b..655c0291129ff 100644
--- a/GPU/Common/GPUCommonAlgorithm.h
+++ b/GPU/Common/GPUCommonAlgorithm.h
@@ -16,6 +16,7 @@
 #define GPUCOMMONALGORITHM_H
 
 #include "GPUCommonDef.h"
+#include "MemLayout.h"
 
 #if !defined(GPUCA_GPUCODE) // Could also enable custom search on the CPU, but it is not always faster, so we stick to std::sort
 #include <algorithm>
@@ -28,19 +29,20 @@ namespace o2::gpu
 {
 class GPUCommonAlgorithm
 {
+
  public:
   template <class T>
-  GPUd() static void sort(T* begin, T* end);
+  GPUd() static void sort(T begin, T end);
   template <class T>
   GPUd() static void sortInBlock(T* begin, T* end);
   template <class T>
-  GPUd() static void sortDeviceDynamic(T* begin, T* end);
+  GPUd() static void sortDeviceDynamic(T begin, T end);
   template <class T, class S>
-  GPUd() static void sort(T* begin, T* end, const S& comp);
+  GPUd() static void sort(T begin, T end, const S& comp);
   template <class T, class S>
   GPUd() static void sortInBlock(T* begin, T* end, const S& comp);
   template <class T, class S>
-  GPUd() static void sortDeviceDynamic(T* begin, T* end, const S& comp);
+  GPUd() static void sortDeviceDynamic(T begin, T end, const S& comp);
 #ifndef __OPENCL__
   template <class T, class S>
   GPUh() static void sortOnDevice(auto* rec, int32_t stream, T* begin, size_t N, const S& comp);
@@ -224,7 +226,7 @@ namespace o2::gpu
 {
 
 template <class T>
-GPUdi() void GPUCommonAlgorithm::sortDeviceDynamic(T* begin, T* end)
+GPUdi() void GPUCommonAlgorithm::sortDeviceDynamic(T begin, T end)
 {
 #ifndef GPUCA_GPUCODE
   GPUCommonAlgorithm::sort(begin, end);
@@ -234,7 +236,7 @@ GPUdi() void GPUCommonAlgorithm::sortDeviceDynamic(T* begin, T* end)
 }
 
 template <class T, class S>
-GPUdi() void GPUCommonAlgorithm::sortDeviceDynamic(T* begin, T* end, const S& comp)
+GPUdi() void GPUCommonAlgorithm::sortDeviceDynamic(T begin, T end, const S& comp)
 {
   GPUCommonAlgorithm::sort(begin, end, comp);
 }
@@ -248,7 +250,7 @@ namespace o2::gpu
 {
 
 template <class T>
-GPUdi() void GPUCommonAlgorithm::sort(T* begin, T* end)
+GPUdi() void GPUCommonAlgorithm::sort(T begin, T end)
 {
 #ifdef GPUCA_ALGORITHM_STD
   std::sort(begin, end);
@@ -258,7 +260,7 @@ GPUdi() void GPUCommonAlgorithm::sort(T* begin, T* end)
 }
 
 template <class T, class S>
-GPUdi() void GPUCommonAlgorithm::sort(T* begin, T* end, const S& comp)
+GPUdi() void GPUCommonAlgorithm::sort(T begin, T end, const S& comp)
 {
 #ifdef GPUCA_ALGORITHM_STD
   std::sort(begin, end, comp);
diff --git a/GPU/Common/GPUCommonAlgorithmThrust.h b/GPU/Common/GPUCommonAlgorithmThrust.h
index 7af3138d45490..fa842b485d746 100644
--- a/GPU/Common/GPUCommonAlgorithmThrust.h
+++ b/GPU/Common/GPUCommonAlgorithmThrust.h
@@ -19,12 +19,14 @@
 #pragma GCC diagnostic push
 #pragma GCC diagnostic ignored "-Wshadow"
 #include <thrust/sort.h>
+#include <thrust/iterator/iterator_traits.h>
 #include <thrust/execution_policy.h>
 #include <thrust/device_ptr.h>
 #pragma GCC diagnostic pop
 
 #include "GPUCommonDef.h"
 #include "GPUCommonHelpers.h"
+#include "GPUTPCTrack.h"
 
 #ifndef __HIPCC__ // CUDA
 #include <cub/cub.cuh>
@@ -81,19 +83,15 @@ GPUdi() void GPUCommonAlgorithm::sortInBlock(T* begin, T* end, const S& comp)
 */
 
 template <class T>
-GPUdi() void GPUCommonAlgorithm::sortDeviceDynamic(T* begin, T* end)
+GPUdi() void GPUCommonAlgorithm::sortDeviceDynamic(T begin, T end)
 {
-  thrust::device_ptr<T> thrustBegin(begin);
-  thrust::device_ptr<T> thrustEnd(end);
-  thrust::sort(GPUCA_THRUST_NAMESPACE::par, thrustBegin, thrustEnd);
+  thrust::sort(GPUCA_THRUST_NAMESPACE::par, begin, end);
 }
 
 template <class T, class S>
-GPUdi() void GPUCommonAlgorithm::sortDeviceDynamic(T* begin, T* end, const S& comp)
+GPUdi() void GPUCommonAlgorithm::sortDeviceDynamic(T begin, T end, const S& comp)
 {
-  thrust::device_ptr<T> thrustBegin(begin);
-  thrust::device_ptr<T> thrustEnd(end);
-  thrust::sort(GPUCA_THRUST_NAMESPACE::par, thrustBegin, thrustEnd, comp);
+  thrust::sort(GPUCA_THRUST_NAMESPACE::par, begin, end, comp);
 }
 
 #ifndef GPUCA_GPUCODE_COMPILEKERNELS
diff --git a/GPU/Common/MemLayout.h b/GPU/Common/MemLayout.h
index 53ca642192d65..1d4952904cdb1 100644
--- a/GPU/Common/MemLayout.h
+++ b/GPU/Common/MemLayout.h
@@ -1,257 +1,212 @@
 #ifndef MEMLAYOUT_H
 #define MEMLAYOUT_H
 
-#include "GPUCommonDefAPI.h"
-
 namespace MemLayout {
 
-template <class T> using value = T;
+using size_t = decltype(sizeof(0));
+using ptrdiff_t = decltype(static_cast<int*>(nullptr) - static_cast<int*>(nullptr));
 
+template <class T> using value = T;
 template <class T> using reference = T&;
-template <class T> using reference_restrict = T& GPUrestrict();
-
 template <class T> using const_reference = const T&;
-template <class T> using const_reference_restrict = const T& GPUrestrict();
-
 template <class T> using pointer = T*;
-template <class T> using pointer_restrtict = T* GPUrestrict();
-
 template <class T> using const_pointer = const T*;
-template <class T> using const_pointer_restrict = const T* GPUrestrict();
-
-using size_t = decltype(sizeof 0);
-using ptrdiff_t = decltype(static_cast<int*>(nullptr) - static_cast<int*>(nullptr));
-
-enum Flag { soa, aos };
-
-// The types S<value>, S<reference>, and S<const_reference> need to be aggregate constructible
-template <template <template <class> class> class S, template <class> class F, Flag L>
-struct wrapper;
-
-template <template <template <class> class> class S, template <class> class F>
-struct wrapper<S, F, Flag::aos> { using type = F<S<value>>; };
-
-template <template <template <class> class> class S, template <class> class F>
-struct wrapper<S, F, Flag::soa> { using type = S<F>; };
-
-namespace type_traits {
-
-template<class T> struct remove_reference { using type = T; };
-template<class T> struct remove_reference<T&> { using type = T; };
-template<class T> struct remove_reference<T&&> { using type = T; };
-
-}  // namespace type_traits
-
-template< class T >
-constexpr type_traits::remove_reference<T>::type&& move(T&& t) noexcept {
-    return static_cast<typename type_traits::remove_reference<T>::type&&>(t);
-}
 
 template <class SF>
 struct RandomAccessAt {
-    size_t i;
+    MemLayout::size_t i;
     template <class... Args>
-    constexpr SF operator()(Args& ...args) const { return {{}, args[i]...}; }
+    constexpr SF operator()(Args& ...args) const { return {{args[i]...}}; }
+    template <class... Args>
+    constexpr SF operator()(const Args& ...args) const { return {{args[i]...}}; }
 };
 
 template <class SF>
 struct GetPointer {
     template <class... Args>
-    constexpr SF operator()(Args& ...args) const { return {{}, &args...}; }
+    constexpr SF operator()(Args& ...args) const { return {{&args...}}; }
+    template <class... Args>
+    constexpr SF operator()(const Args& ...args) const { return {{&args...}}; }
 };
 
 template <class SF>
 struct AggregateConstructor {
     template <class... Args>
-    constexpr SF operator()(Args& ...args) const { return {{}, args...}; } 
+    constexpr SF operator()(Args& ...args) const { return {{args...}}; }
+    template <class... Args>
+    constexpr SF operator()(const Args& ...args) const { return {{args...}}; } 
 };
 
-template <
-    template <class> class F_left,
-    template <class> class F_right
->
-struct CopyAssignment {
-    template <class T>
-    constexpr void operator()(F_left<T>& left, F_right<T>& right) const { left = right; }
+struct FirstMember {
+    template <class T, class... Args>
+    constexpr T& operator()(T& t, Args& ...args) const { return t; }
+    template <class T, class... Args>
+    constexpr const T& operator()(const T& t, const Args& ...args) const { return t; }
 };
 
-template <
-    template <class> class F_left,
-    template <class> class F_right
->
-struct MoveAssignment {
-    template <class T>
-    constexpr void operator()(F_left<T>& left, F_right<T>& right) const { left = move(right); }
+template <class SF>
+struct PreIncrement {
+    template <class... Args>
+    constexpr SF operator()(Args& ...args) const { return {{++args...}}; }
+};
+
+template <class SF>
+struct PreDecrement {
+    template <class... Args>
+    constexpr SF operator()(Args& ...args) const { return {{--args...}}; }
+};
+
+template <class SF>
+struct Advance {
+    ptrdiff_t i;
+    template <class... Args>
+    constexpr SF operator()(const Args& ...args) const { return {{(args + i)...}}; }
+};
+
+struct CopyAssignment {
+    template <class Left, class Right>
+    constexpr Left& operator()(Left& left, const Right& right) const { return left = right; }
 };
 
 template <
     template <template <class> class> class S,
     template <class> class F
 >
-struct CRTP {
-    using Derived = S<F>;
-    template <template <class> class F_out>
-    constexpr operator S<F_out>() {
-        return static_cast<Derived*>(this)->apply(AggregateConstructor<S<F_out>>{});
-    }
-    template <template <class> class F_out>
-    constexpr operator S<F_out>() const {
-        return static_cast<const Derived*>(this)->apply(AggregateConstructor<S<F_out>>{});
-    }
-    constexpr S<reference> operator[] (size_t i) {
-        return static_cast<Derived*>(this)->apply(RandomAccessAt<S<reference>>{i});
-    }
-    constexpr S<const_reference> operator[] (size_t i) const {
-        return static_cast<const Derived*>(this)->apply(RandomAccessAt<S<const_reference>>{i});
-    }
-    constexpr S<pointer> operator& () {
-        return static_cast<Derived*>(this)->apply(GetPointer<S<pointer>>{});
-    }
-    constexpr S<const_pointer> operator& () const {
-        return static_cast<const Derived*>(this)->apply(GetPointer<S<const_pointer>>{});
-    }
-    constexpr S<reference> operator*() {
-        return static_cast<Derived*>(this)->operator[](0);
-    }
-    constexpr S<const_reference> operator*() const {
-        return static_cast<const Derived*>(this)->operator[](0);
-    }
+struct wrapper : public S<F> {
+    using Base = S<F>;
+
+    constexpr wrapper() = default;
+    constexpr wrapper(Base b) : Base{static_cast<Base&&>(b)} {}
+    template <template <class> class F_other>
+    constexpr wrapper(wrapper<S, F_other>& other) : Base{other.apply(AggregateConstructor<wrapper>{})} {}
+
+    constexpr wrapper<S, reference> operator[] (size_t i) { return Base::apply(RandomAccessAt<wrapper<S, reference>>{i}); }
+    constexpr wrapper<S, const_reference> operator[] (size_t i) const { return Base::apply(RandomAccessAt<wrapper<S, const_reference>>{i}); }
+
+    constexpr wrapper<S, reference> operator*() { return operator[](0); }
+    constexpr wrapper<S, const_reference> operator*(ptrdiff_t) const { return operator[](0); }
 };
 
 template <template <template <class> class> class S>
-struct CRTP<S, value> {
-    using Derived = S<value>;
-    template <template <class> class F_out>
-    constexpr operator S<F_out>() {
-        return static_cast<Derived*>(this)->apply(AggregateConstructor<S<F_out>>{});
-    }
-    template <template <class> class F_out>
-    constexpr operator S<F_out>() const {
-        return static_cast<const Derived*>(this)->apply(AggregateConstructor<S<F_out>>{});
-    }
+struct wrapper<S, value> : public S<value> {
+    using Base = S<value>;
+
+    constexpr wrapper() = default;
+    constexpr wrapper(Base b) : Base{static_cast<Base&&>(b)} {}
+    constexpr wrapper(wrapper<S, reference> other) : Base(other.apply(AggregateConstructor<wrapper>{})) {}
+    constexpr wrapper(wrapper<S, const_reference> other) : Base(other.apply(AggregateConstructor<wrapper>{})) {}
+
+    //constexpr S<MemLayout::pointer> operator& () { return apply(GetPointer<S<MemLayout::pointer>>{}); }
+    //constexpr S<MemLayout::const_pointer> operator& () const { return apply(GetPointer<S<MemLayout::const_pointer>>{}); }
 };
 
 template <template <template <class> class> class S>
-struct CRTP<S, reference> {
-    using Derived = S<reference>;
-    template <template <class> class F_out>
-    constexpr operator S<F_out>() {
-        return static_cast<Derived*>(this)->apply(AggregateConstructor<S<F_out>>{});
-    }
-    template <template <class> class F_out>
-    constexpr operator S<F_out>() const {
-        return static_cast<const Derived*>(this)->apply(AggregateConstructor<S<F_out>>{});
-    }
-    template <template <class> class F_other>
-    constexpr Derived& operator=(S<F_other>& other) {
-        memberwise(*static_cast<Derived*>(this), other, CopyAssignment<reference, F_other>{});
-        return *static_cast<Derived*>(this);
-    }
-    template <template <class> class F_other>
-    constexpr Derived& operator=(S<F_other>&& other) {
-        memberwise(*static_cast<Derived*>(this), other, MoveAssignment<reference, F_other>{});
-        return *static_cast<Derived*>(this);
+struct wrapper<S, reference> : public S<reference> {
+    using Base = S<reference>;
+
+    constexpr wrapper() = delete;
+    constexpr wrapper(Base b) : Base{static_cast<Base&&>(b)} {}
+    constexpr wrapper(wrapper<S, value>& other) : Base(other.apply(AggregateConstructor<wrapper>{})) {}
+    
+    constexpr wrapper(const wrapper& other) = default;
+
+    constexpr wrapper& operator=(const wrapper& other) {
+        Base::apply(other, CopyAssignment{});
+        return *this;
     }
-    constexpr S<pointer> operator& () {
-        return static_cast<Derived*>(this)->apply(GetPointer<S<pointer>>{});
+    constexpr wrapper& operator=(const wrapper<S, const_reference>& other) {
+        Base::apply(other, CopyAssignment{});
+        return *this;
     }
-    constexpr S<const_pointer> operator& () const {
-        return static_cast<const Derived*>(this)->apply(GetPointer<S<const_pointer>>{});
+    constexpr wrapper& operator=(const wrapper<S, value>& other) {
+        Base::apply(other, CopyAssignment{});
+        return *this;
     }
+
+    constexpr wrapper(wrapper&& other) = default;
+    
+    constexpr wrapper& operator=(wrapper&& other) { return operator=(other); }
+
+    constexpr wrapper<S, pointer> operator& () { return Base::apply(GetPointer<wrapper<S, pointer>>{}); }
+    //constexpr wrapper<S, const_pointer> operator& () const { return Base::apply(GetPointer<wrapper<S, const_pointer>>{}); }
 };
 
 template <template <template <class> class> class S>
-struct CRTP<S, const_reference> {
-    using Derived = S<const_reference>;
-    template <template <class> class F_out>
-    constexpr operator S<F_out>() const {
-        return static_cast<const Derived*>(this)->apply(AggregateConstructor<S<F_out>>{});
-    }
-    constexpr S<const_pointer> operator& () const {
-        return static_cast<const Derived*>(this)->apply(GetPointer<S<const_pointer>>{});
-    }
+struct wrapper<S, const_reference> : public S<const_reference> {
+    using Base = S<const_reference>;
+
+    constexpr wrapper() = delete;
+    constexpr wrapper(Base b) : Base{static_cast<Base&&>(b)} {}
+    constexpr wrapper(wrapper<S, value> other) : Base(other.apply(AggregateConstructor<wrapper>{})) {}
+    constexpr wrapper(wrapper<S, reference> other) : Base(other.apply(AggregateConstructor<wrapper>{})) {}
+
+    constexpr wrapper<S, const_pointer> operator&() const { return Base::apply(GetPointer<wrapper<S, const_pointer>>{}); }
 };
 
 template <template <template <class> class> class S>
-struct CRTP<S, reference_restrict> {
-    using Derived = S<reference_restrict>;
-    template <template <class> class F_out>
-    constexpr operator S<F_out>() {
-        return static_cast<Derived*>(this)->apply(AggregateConstructor<S<F_out>>{});
-    }
-    template <template <class> class F_out>
-    constexpr operator S<F_out>() const {
-        return static_cast<const Derived*>(this)->apply(AggregateConstructor<S<F_out>>{});
-    }
-    template <template <class> class F_other>
-    constexpr Derived& operator=(S<F_other>& other) {
-        memberwise(*static_cast<Derived*>(this), other, CopyAssignment<reference_restrict, F_other>{});
-        return *static_cast<Derived*>(this);
-    }
-    template <template <class> class F_other>
-    constexpr Derived& operator=(S<F_other>&& other) {
-        memberwise(*static_cast<Derived*>(this), other, MoveAssignment<reference_restrict, F_other>{});
-        return *static_cast<Derived*>(this);
-    }
-    constexpr S<pointer_restrtict> operator& () {
-        return static_cast<Derived*>(this)->apply(GetPointer<S<pointer_restrtict>>{});
-    }
-    constexpr S<const_pointer_restrict> operator& () const {
-        return static_cast<const Derived*>(this)->apply(GetPointer<S<const_pointer_restrict>>{});
-    }
+struct wrapper<S, pointer> : public S<pointer> {
+    using Base = S<pointer>;
+
+    constexpr wrapper() = default;
+    constexpr wrapper(Base b) : Base{static_cast<Base&&>(b)} {}
+
+    constexpr wrapper<S, reference> operator[] (size_t i) { return Base::apply(RandomAccessAt<wrapper<S, reference>>{i}); }
+    constexpr wrapper<S, const_reference> operator[] (size_t i) const { return Base::apply(RandomAccessAt<wrapper<S, const_reference>>{i}); }
+
+    constexpr wrapper<S, reference> operator*() { return operator[](0); }
+    constexpr wrapper<S, const_reference> operator*() const { return operator[](0); }
+
+    constexpr bool operator==(const wrapper& other) const { return Base::apply(FirstMember{}) == other.apply(FirstMember{}); }
+    constexpr bool operator!=(const wrapper& other) const { return !this->operator==(other); }
+    constexpr bool operator<(const wrapper& other) const { return Base::apply(FirstMember{}) < other.apply(FirstMember{}); }
+
+    constexpr wrapper operator+(ptrdiff_t i) const { return Base::apply(Advance<wrapper>{i}); }
+    constexpr wrapper operator-(ptrdiff_t i) const { return operator+(-i); }
+    constexpr ptrdiff_t operator-(const wrapper& other) const { return Base::apply(FirstMember{}) - other.apply(FirstMember{}); }
+
+    constexpr wrapper& operator++() { Base::apply(PreIncrement<wrapper>{}); return *this; }
+    constexpr wrapper& operator+=(ptrdiff_t i) { return *this = *this + i; }
+    constexpr wrapper& operator--() { Base::apply(PreDecrement<wrapper>{}); return *this; }
+    constexpr wrapper& operator-=(ptrdiff_t i) { return *this = *this - i; }
 };
 
 template <template <template <class> class> class S>
-struct CRTP<S, const_reference_restrict> {
-    using Derived = S<const_reference_restrict>;
-    template <template <class> class F_out>
-    constexpr operator S<F_out>() const {
-        return static_cast<const Derived*>(this)->apply(AggregateConstructor<S<F_out>>{});
-    }
-    constexpr S<const_pointer_restrict> operator& () const {
-        return static_cast<const Derived*>(this)->apply(GetPointer<S<const_pointer_restrict>>{});
-    }
-};
+struct wrapper<S, const_pointer> : public S<const_pointer> {
+    using Base = S<const_pointer>;
 
-template <
-    template <template <class> class> class S,
-    template <class> class F
->
-struct iterator {
-    //using iterator_category = std::random_access_iterator_tag;
-    using difference_type = ptrdiff_t;
-    using value_type = S<value>;
-    using pointer = S<F>;
-    using reference = S<reference>;
-
-    difference_type index;
-    pointer handle;
-
-    constexpr bool operator==(iterator const& other) const { return index == other.index; }
-    constexpr bool operator!=(iterator const& other) const { return index != other.index; }
-    constexpr bool operator<(iterator const& other) const { return index < other.index; }
-
-    constexpr iterator operator+(difference_type i) const { return {index + i, handle}; }
-    constexpr iterator operator-(difference_type i) const { return {index - i, handle}; }
-    
-    constexpr difference_type operator-(iterator const& other) const { 
-        return difference_type(index) - difference_type(other.index); 
-    }
-    
-    constexpr iterator& operator++() { ++index; return *this; }
-    constexpr iterator& operator--() { --index; return *this; }
+    constexpr wrapper() = default;
+    constexpr wrapper(Base b) : Base{static_cast<Base&&>(b)} {}
+    constexpr wrapper(wrapper<S, pointer> other) : Base(other.apply(AggregateConstructor<wrapper>{})) {}
 
-    constexpr reference operator*() { return handle[index]; }
+    constexpr wrapper<S, const_reference> operator[] (size_t i) const { return Base::apply(RandomAccessAt<wrapper<S, const_reference>>{i}); }
+    constexpr wrapper<S, const_reference> operator*() const { return operator[](0); }
 };
 
+enum Flag { soa, aos };
+
+template <template <template <class> class> class S, template <class> class F, Flag L>
+struct interface;
+
+template <template <template <class> class> class S, template <class> class F>
+struct interface<S, F, Flag::aos> { using type = F<S<value>>; };
+
+template <template <template <class> class> class S, template <class> class F>
+struct interface<S, F, Flag::soa> { using type = wrapper<S, F>; };
+
 }  // namespace MemLayout
 
-#define MEMLAYOUT_MEMBERFUNCTIONS(STRUCT, CONTAINER, ...)\
+#define MEMLAYOUT_APPLY_UNARY(...)\
     template <class Function>\
     constexpr auto apply(Function&& f) { return f(__VA_ARGS__); }\
     template <class Function>\
     constexpr auto apply(Function&& f) const { return f(__VA_ARGS__); }\
-    template <template <class> class F_left, template <class> class F_right, class FunctionObject>\
-    constexpr friend void memberwise(STRUCT<F_left>& left, STRUCT<F_right>& right, FunctionObject&& f);\
+
+#define MEMLAYOUT_EXPAND(m) f(m, other.m)
+
+#define MEMLAYOUT_APPLY_BINARY(STRUCT_NAME, ...)\
+    template <template <class> class F_other, class Function>\
+    constexpr STRUCT_NAME apply(const STRUCT_NAME<F_other>& other, Function&& f) { return {__VA_ARGS__}; }\
+    template <template <class> class F_other, class Function>\
+    constexpr STRUCT_NAME apply(const STRUCT_NAME<F_other>& other, Function&& f) const { return {__VA_ARGS__}; }\
 
 #endif // MEMLAYOUT_H
diff --git a/GPU/GPUTracking/DataTypes/GPUDataTypes.h b/GPU/GPUTracking/DataTypes/GPUDataTypes.h
index 6713b0a906bf1..4e08fbe12e100 100644
--- a/GPU/GPUTracking/DataTypes/GPUDataTypes.h
+++ b/GPU/GPUTracking/DataTypes/GPUDataTypes.h
@@ -24,6 +24,7 @@
 #include <cstddef>
 #endif
 #include "GPUTRDDef.h"
+#include "GPUTPCTrack.h"
 #include "MemLayout.h"
 
 struct AliHLTTPCClusterMCLabel;
@@ -98,8 +99,6 @@ struct TPCZSLinkMapping;
 
 #include "utils/bitfield.h"
 
-template <template <class> class F>
-class GPUTPCTrackSkeleton;
 class GPUTPCHitId;
 class GPUTPCGMMergedTrack;
 struct GPUTPCGMMergedTrackHit;
@@ -226,7 +225,7 @@ struct GPUTrackingInOutPointers {
   const AliHLTTPCRawCluster* rawClusters[NSECTORS] = {nullptr};
   uint32_t nRawClusters[NSECTORS] = {0};
   const o2::tpc::ClusterNativeAccess* clustersNative = nullptr;
-  const GPUTPCTrackSkeleton<MemLayout::value>* sectorTracks[NSECTORS] = {nullptr};
+  MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_pointer> sectorTracks[NSECTORS];// = {{nullptr, nullptr, nullptr, {nullptr, nullptr, nullptr, nullptr}}};
   uint32_t nSectorTracks[NSECTORS] = {0};
   const GPUTPCHitId* sectorClusters[NSECTORS] = {nullptr};
   uint32_t nSectorClusters[NSECTORS] = {0};
diff --git a/GPU/GPUTracking/Global/GPUChainTracking.cxx b/GPU/GPUTracking/Global/GPUChainTracking.cxx
index 9d2578731a30c..b9749ff5ea609 100644
--- a/GPU/GPUTracking/Global/GPUChainTracking.cxx
+++ b/GPU/GPUTracking/Global/GPUChainTracking.cxx
@@ -554,7 +554,18 @@ void GPUChainTracking::AllocateIOMemory()
   for (uint32_t i = 0; i < NSECTORS; i++) {
     AllocateIOMemoryHelper(mIOPtrs.nClusterData[i], mIOPtrs.clusterData[i], mIOMem.clusterData[i]);
     AllocateIOMemoryHelper(mIOPtrs.nRawClusters[i], mIOPtrs.rawClusters[i], mIOMem.rawClusters[i]);
-    AllocateIOMemoryHelper(mIOPtrs.nSectorTracks[i], mIOPtrs.sectorTracks[i], mIOMem.sectorTracks[i]);
+    
+    //AllocateIOMemoryHelper(mIOPtrs.nSectorTracks[i], mIOPtrs.sectorTracks[i], mIOMem.sectorTracks[i]);
+
+    AllocateIOMemoryHelper(mIOPtrs.nSectorTracks[i], mIOPtrs.sectorTracks[i].mFirstHitID, mIOMem.sectorTracks[i].mFirstHitID);
+    AllocateIOMemoryHelper(mIOPtrs.nSectorTracks[i], mIOPtrs.sectorTracks[i].mNHits, mIOMem.sectorTracks[i].mNHits);
+    AllocateIOMemoryHelper(mIOPtrs.nSectorTracks[i], mIOPtrs.sectorTracks[i].mLocalTrackId, mIOMem.sectorTracks[i].mLocalTrackId);
+
+    AllocateIOMemoryHelper(mIOPtrs.nSectorTracks[i], mIOPtrs.sectorTracks[i].mParam.mX, mIOMem.sectorTracks[i].mParam.mX);
+    AllocateIOMemoryHelper(mIOPtrs.nSectorTracks[i], mIOPtrs.sectorTracks[i].mParam.mC, mIOMem.sectorTracks[i].mParam.mC);
+    AllocateIOMemoryHelper(mIOPtrs.nSectorTracks[i], mIOPtrs.sectorTracks[i].mParam.mZOffset, mIOMem.sectorTracks[i].mParam.mZOffset);
+    AllocateIOMemoryHelper(mIOPtrs.nSectorTracks[i], mIOPtrs.sectorTracks[i].mParam.mP, mIOMem.sectorTracks[i].mParam.mP);
+
     AllocateIOMemoryHelper(mIOPtrs.nSectorClusters[i], mIOPtrs.sectorClusters[i], mIOMem.sectorClusters[i]);
   }
   mIOMem.clusterNativeAccess.reset(new ClusterNativeAccess);
diff --git a/GPU/GPUTracking/Global/GPUChainTracking.h b/GPU/GPUTracking/Global/GPUChainTracking.h
index 2a9ce862d86f7..251311aebed4b 100644
--- a/GPU/GPUTracking/Global/GPUChainTracking.h
+++ b/GPU/GPUTracking/Global/GPUChainTracking.h
@@ -108,7 +108,12 @@ class GPUChainTracking : public GPUChain
     std::unique_ptr<AliHLTTPCRawCluster[]> rawClusters[NSECTORS];
     std::unique_ptr<o2::tpc::ClusterNative[]> clustersNative;
     std::unique_ptr<o2::tpc::ClusterNativeAccess> clusterNativeAccess;
-    std::unique_ptr<GPUTPCTrackSkeleton<MemLayout::value>[]> sectorTracks[NSECTORS];
+
+    template <class T>
+    using unique_ptr_array = std::unique_ptr<T[]>;
+
+    MemLayout::wrapper<GPUTPCTrackSkeleton, unique_ptr_array> sectorTracks[NSECTORS];
+
     std::unique_ptr<GPUTPCHitId[]> sectorClusters[NSECTORS];
     std::unique_ptr<AliHLTTPCClusterMCLabel[]> mcLabelsTPC;
     std::unique_ptr<GPUTPCMCInfo[]> mcInfosTPC;
diff --git a/GPU/GPUTracking/Global/GPUChainTrackingIO.cxx b/GPU/GPUTracking/Global/GPUChainTrackingIO.cxx
index dd11e9989f684..d47713bdb3edd 100644
--- a/GPU/GPUTracking/Global/GPUChainTrackingIO.cxx
+++ b/GPU/GPUTracking/Global/GPUChainTrackingIO.cxx
@@ -141,7 +141,9 @@ void GPUChainTracking::DumpData(const char* filename, const GPUTrackingInOutPoin
     uint32_t n = 1;
     DumpData(fp, &ioPtrs->settingsTF, &n, InOutPointerType::TF_SETTINGS);
   }
-  DumpData(fp, ioPtrs->sectorTracks, ioPtrs->nSectorTracks, InOutPointerType::SECTOR_OUT_TRACK);
+
+  //DumpData(fp, ioPtrs->sectorTracks, ioPtrs->nSectorTracks, InOutPointerType::SECTOR_OUT_TRACK);
+
   DumpData(fp, ioPtrs->sectorClusters, ioPtrs->nSectorClusters, InOutPointerType::SECTOR_OUT_CLUSTER);
   DumpData(fp, &ioPtrs->mcLabelsTPC, &ioPtrs->nMCLabelsTPC, InOutPointerType::MC_LABEL_TPC);
   DumpData(fp, &ioPtrs->mcInfosTPC, &ioPtrs->nMCInfosTPC, InOutPointerType::MC_INFO_TPC);
@@ -242,7 +244,9 @@ int32_t GPUChainTracking::ReadData(const char* filename)
   }
   uint32_t n;
   ReadData(fp, &mIOPtrs.settingsTF, &n, &mIOMem.settingsTF, InOutPointerType::TF_SETTINGS);
-  ReadData(fp, mIOPtrs.sectorTracks, mIOPtrs.nSectorTracks, mIOMem.sectorTracks, InOutPointerType::SECTOR_OUT_TRACK);
+
+  //ReadData(fp, mIOPtrs.sectorTracks, mIOPtrs.nSectorTracks, mIOMem.sectorTracks, InOutPointerType::SECTOR_OUT_TRACK);
+  
   ReadData(fp, mIOPtrs.sectorClusters, mIOPtrs.nSectorClusters, mIOMem.sectorClusters, InOutPointerType::SECTOR_OUT_CLUSTER);
   ReadData(fp, &mIOPtrs.mcLabelsTPC, &mIOPtrs.nMCLabelsTPC, &mIOMem.mcLabelsTPC, InOutPointerType::MC_LABEL_TPC);
   ReadData(fp, &mIOPtrs.mcInfosTPC, &mIOPtrs.nMCInfosTPC, &mIOMem.mcInfosTPC, InOutPointerType::MC_INFO_TPC);
diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx
index 338ecae4f9b95..d2fd297f552ff 100644
--- a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx
+++ b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx
@@ -237,7 +237,7 @@ int64_t GPUTPCGMMerger::GetTrackLabelA(const S& trk) const
   int32_t nClusters = 0;
   if constexpr (std::is_same_v<S, GPUTPCGMBorderTrack&>) {
     sectorTrack = &mSectorTrackInfos[trk.TrackID()];
-    nClusters = sectorTrack->OrigTrack()->NHits();
+    nClusters = (*sectorTrack->OrigTrack()).NHits();
   } else {
     nClusters = trk.NClusters();
   }
@@ -246,7 +246,7 @@ int64_t GPUTPCGMMerger::GetTrackLabelA(const S& trk) const
     int32_t id;
     if constexpr (std::is_same_v<S, GPUTPCGMBorderTrack&>) {
       const GPUTPCTracker& tracker = GetConstantMem()->tpcTrackers[sectorTrack->Sector()];
-      const GPUTPCHitId& ic = tracker.TrackHits()[sectorTrack->OrigTrack()->FirstHitID() + i];
+      const GPUTPCHitId& ic = tracker.TrackHits()[(*sectorTrack->OrigTrack()).FirstHitID() + i];
       id = tracker.Data().ClusterDataIndex(tracker.Data().Row(ic.RowIndex()), ic.HitIndex()) + GetConstantMem()->ioPtrs.clustersNative->clusterOffset[sectorTrack->Sector()][0];
     } else {
       id = mClusters[trk.FirstClusterRef() + i].num;
@@ -488,7 +488,7 @@ GPUd() void GPUTPCGMMerger::ClearTrackLinks(int32_t nBlocks, int32_t nThreads, i
   }
 }
 
-GPUd() int32_t GPUTPCGMMerger::RefitSectorTrack(GPUTPCGMSectorTrack& sectorTrack, const GPUTPCTrack* inTrack, float alpha, int32_t sector)
+GPUd() int32_t GPUTPCGMMerger::RefitSectorTrack(GPUTPCGMSectorTrack& sectorTrack, MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_pointer> inTrack, float alpha, int32_t sector)
 {
   GPUTPCGMPropagator prop;
   prop.SetMaterialTPC();
@@ -497,14 +497,15 @@ GPUd() int32_t GPUTPCGMMerger::RefitSectorTrack(GPUTPCGMSectorTrack& sectorTrack
   prop.SetFitInProjections(false);
   prop.SetPolynomialField(&Param().polynomialField);
   GPUTPCGMTrackParam trk;
-  trk.X() = inTrack->Param().GetX();
-  trk.Y() = inTrack->Param().GetY();
-  trk.Z() = inTrack->Param().GetZ();
-  trk.SinPhi() = inTrack->Param().GetSinPhi();
-  trk.DzDs() = inTrack->Param().GetDzDs();
-  trk.QPt() = inTrack->Param().GetQPt();
-  trk.TZOffset() = Param().par.earlyTpcTransform ? inTrack->Param().GetZOffset() : GetConstantMem()->calibObjects.fastTransformHelper->getCorrMap()->convZOffsetToVertexTime(sector, inTrack->Param().GetZOffset(), Param().continuousMaxTimeBin);
-  trk.ShiftZ(this, sector, sectorTrack.ClusterZT0(), sectorTrack.ClusterZTN(), inTrack->Param().GetX(), inTrack->Param().GetX()); // We do not store the inner / outer cluster X, so we just use the track X instead
+  MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_reference> inTrack_cref = *inTrack;
+  trk.X() = inTrack_cref.Param().GetX();
+  trk.Y() = inTrack_cref.Param().GetY();
+  trk.Z() = inTrack_cref.Param().GetZ();
+  trk.SinPhi() = inTrack_cref.Param().GetSinPhi();
+  trk.DzDs() = inTrack_cref.Param().GetDzDs();
+  trk.QPt() = inTrack_cref.Param().GetQPt();
+  trk.TZOffset() = Param().par.earlyTpcTransform ? inTrack_cref.Param().GetZOffset() : GetConstantMem()->calibObjects.fastTransformHelper->getCorrMap()->convZOffsetToVertexTime(sector, inTrack_cref.Param().GetZOffset(), Param().continuousMaxTimeBin);
+  trk.ShiftZ(this, sector, sectorTrack.ClusterZT0(), sectorTrack.ClusterZTN(), inTrack_cref.Param().GetX(), inTrack_cref.Param().GetX()); // We do not store the inner / outer cluster X, so we just use the track X instead
   sectorTrack.SetX2(0.f);
   for (int32_t way = 0; way < 2; way++) {
     if (way) {
@@ -513,14 +514,14 @@ GPUd() int32_t GPUTPCGMMerger::RefitSectorTrack(GPUTPCGMSectorTrack& sectorTrack
     }
     trk.ResetCovariance();
     prop.SetTrack(&trk, alpha);
-    int32_t start = way ? inTrack->NHits() - 1 : 0;
-    int32_t end = way ? 0 : (inTrack->NHits() - 1);
+    int32_t start = way ? inTrack_cref.NHits() - 1 : 0;
+    int32_t end = way ? 0 : (inTrack_cref.NHits() - 1);
     int32_t incr = way ? -1 : 1;
     for (int32_t i = start; i != end; i += incr) {
       float x, y, z;
       int32_t row, flags;
       const GPUTPCTracker& tracker = GetConstantMem()->tpcTrackers[sector];
-      const GPUTPCHitId& ic = tracker.TrackHits()[inTrack->FirstHitID() + i];
+      const GPUTPCHitId& ic = tracker.TrackHits()[inTrack_cref.FirstHitID() + i];
       int32_t clusterIndex = tracker.Data().ClusterDataIndex(tracker.Data().Row(ic.RowIndex()), ic.HitIndex());
       row = ic.RowIndex();
       const ClusterNative& cl = GetConstantMem()->ioPtrs.clustersNative->clustersLinear[GetConstantMem()->ioPtrs.clustersNative->clusterOffset[sector][0] + clusterIndex];
@@ -550,11 +551,12 @@ GPUd() int32_t GPUTPCGMMerger::RefitSectorTrack(GPUTPCGMSectorTrack& sectorTrack
   return 0;
 }
 
-GPUd() void GPUTPCGMMerger::SetTrackClusterZT(GPUTPCGMSectorTrack& track, int32_t iSector, const GPUTPCTrack* sectorTr)
+GPUd() void GPUTPCGMMerger::SetTrackClusterZT(GPUTPCGMSectorTrack& track, int32_t iSector, MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_pointer> sectorTr)
 {
+  MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_reference> sectorTr_cref = *sectorTr;
   const GPUTPCTracker& trk = GetConstantMem()->tpcTrackers[iSector];
-  const GPUTPCHitId& ic1 = trk.TrackHits()[sectorTr->FirstHitID()];
-  const GPUTPCHitId& ic2 = trk.TrackHits()[sectorTr->FirstHitID() + sectorTr->NHits() - 1];
+  const GPUTPCHitId& ic1 = trk.TrackHits()[sectorTr_cref.FirstHitID()];
+  const GPUTPCHitId& ic2 = trk.TrackHits()[sectorTr_cref.FirstHitID() + sectorTr_cref.NHits() - 1];
   int32_t clusterIndex1 = trk.Data().ClusterDataIndex(trk.Data().Row(ic1.RowIndex()), ic1.HitIndex());
   int32_t clusterIndex2 = trk.Data().ClusterDataIndex(trk.Data().Row(ic2.RowIndex()), ic2.HitIndex());
   if (Param().par.earlyTpcTransform) {
@@ -574,12 +576,13 @@ GPUd() void GPUTPCGMMerger::UnpackSectorGlobal(int32_t nBlocks, int32_t nThreads
 {
   const GPUTPCTracker& trk = GetConstantMem()->tpcTrackers[iSector];
   float alpha = Param().Alpha(iSector);
-  const GPUTPCTrack* sectorTr = mMemory->firstExtrapolatedTracks[iSector];
+  // GPUTPCTrackSkeleton<MemLayout::const_pointer> sectorTr = mMemory->firstExtrapolatedTracks[iSector];
+  MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_pointer> sectorTr = mMemory->firstExtrapolatedTracks[iSector];
   uint32_t nLocalTracks = trk.CommonMemory()->nLocalTracks;
   uint32_t nTracks = *trk.NTracks();
   for (uint32_t itr = nLocalTracks + iBlock * nThreads + iThread; itr < nTracks; itr += nBlocks * nThreads) {
     sectorTr = &trk.Tracks()[itr];
-    int32_t localId = mTrackIDs[(sectorTr->LocalTrackId() >> 24) * mNMaxSingleSectorTracks + (sectorTr->LocalTrackId() & 0xFFFFFF)];
+    int32_t localId = mTrackIDs[((*sectorTr).LocalTrackId() >> 24) * mNMaxSingleSectorTracks + ((*sectorTr).LocalTrackId() & 0xFFFFFF)];
     if (localId == -1) {
       continue;
     }
@@ -611,10 +614,9 @@ GPUd() void GPUTPCGMMerger::RefitSectorTracks(int32_t nBlocks, int32_t nThreads,
   uint32_t nLocalTracks = trk.CommonMemory()->nLocalTracks;
 
   float alpha = Param().Alpha(iSector);
-  const GPUTPCTrack* sectorTr = nullptr;
 
   for (uint32_t itr = iBlock * nThreads + iThread; itr < nLocalTracks; itr += nBlocks * nThreads) {
-    sectorTr = &trk.Tracks()[itr];
+    MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_pointer> sectorTr = &trk.Tracks()[itr];
     GPUTPCGMSectorTrack track;
     SetTrackClusterZT(track, iSector, sectorTr);
     if (Param().rec.tpc.mergerCovSource == 0) {
@@ -642,7 +644,7 @@ GPUd() void GPUTPCGMMerger::RefitSectorTracks(int32_t nBlocks, int32_t nThreads,
     track.SetExtrapolatedTrackId(0, -1);
     track.SetExtrapolatedTrackId(1, -1);
     uint32_t myTrack = CAMath::AtomicAdd(&mMemory->nUnpackedTracks, 1u);
-    mTrackIDs[iSector * mNMaxSingleSectorTracks + sectorTr->LocalTrackId()] = myTrack;
+    mTrackIDs[iSector * mNMaxSingleSectorTracks + (*sectorTr).LocalTrackId()] = myTrack;
     mSectorTrackInfos[myTrack] = track;
   }
 }
@@ -701,7 +703,7 @@ GPUd() void GPUTPCGMMerger::MergeSectorsPrepareStep2(int32_t nBlocks, int32_t nT
       const GPUTPCGMSectorTrack* trackMin = track;
       while (track->NextSegmentNeighbour() >= 0 && track->Sector() == mSectorTrackInfos[track->NextSegmentNeighbour()].Sector()) {
         track = &mSectorTrackInfos[track->NextSegmentNeighbour()];
-        if (track->OrigTrack()->Param().X() < trackMin->OrigTrack()->Param().X()) {
+        if ((*track->OrigTrack()).Param().X() < (*trackMin->OrigTrack()).Param().X()) {
           trackMin = track;
         }
       }
@@ -1603,7 +1605,7 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread
       trackCluster* c2 = trackClusters + nHits + nTrackHits - 1;
       for (int32_t i = 0; i < nTrackHits; i++, c2--) {
         const GPUTPCTracker& trk = GetConstantMem()->tpcTrackers[t->Sector()];
-        const GPUTPCHitId& ic = trk.TrackHits()[t->OrigTrack()->FirstHitID() + i];
+        const GPUTPCHitId& ic = trk.TrackHits()[(*t->OrigTrack()).FirstHitID() + i];
         uint32_t id = trk.Data().ClusterDataIndex(trk.Data().Row(ic.RowIndex()), ic.HitIndex()) + GetConstantMem()->ioPtrs.clustersNative->clusterOffset[t->Sector()][0];
         *c2 = trackCluster{id, (uint8_t)ic.RowIndex(), t->Sector(), t->Leg()};
       }
@@ -1657,9 +1659,9 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread
           if (trackParts[i]->Leg() != baseLeg) {
             break;
           }
-          if (trackParts[i]->OrigTrack()->NHits() > length) {
+          if ((*trackParts[i]->OrigTrack()).NHits() > length) {
             iLongest = i;
-            length = trackParts[i]->OrigTrack()->NHits();
+            length = (*trackParts[i]->OrigTrack()).NHits();
           }
         }
         bool outwards;
diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMerger.h b/GPU/GPUTracking/Merger/GPUTPCGMMerger.h
index 54a541ebe0fd6..1adc6c377a0ba 100644
--- a/GPU/GPUTracking/Merger/GPUTPCGMMerger.h
+++ b/GPU/GPUTracking/Merger/GPUTPCGMMerger.h
@@ -24,6 +24,7 @@
 #include "GPUProcessor.h"
 #include "GPUTPCGMMergerTypes.h"
 #include "GPUGeneralKernels.h"
+#include "MemLayout.h"
 
 #if !defined(GPUCA_GPUCODE)
 #include <cmath>
@@ -73,7 +74,7 @@ class GPUTPCGMMerger : public GPUProcessor
     GPUAtomic(uint32_t) nMergedTrackClusters;
     GPUAtomic(uint32_t) nO2Tracks;
     GPUAtomic(uint32_t) nO2ClusRefs;
-    const GPUTPCTrack* firstExtrapolatedTracks[NSECTORS];
+    MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_pointer> firstExtrapolatedTracks[NSECTORS];
     GPUAtomic(uint32_t) tmpCounter[2 * NSECTORS];
     GPUAtomic(uint32_t) nLooperMatchCandidates;
   };
@@ -153,8 +154,8 @@ class GPUTPCGMMerger : public GPUProcessor
   GPUd() uint16_t MemoryResOutputO2MC() const { return mMemoryResOutputO2MC; }
   GPUd() uint16_t MemoryResOutputO2Scratch() const { return mMemoryResOutputO2Scratch; }
 
-  GPUd() int32_t RefitSectorTrack(GPUTPCGMSectorTrack& sectorTrack, const GPUTPCTrack* inTrack, float alpha, int32_t sector);
-  GPUd() void SetTrackClusterZT(GPUTPCGMSectorTrack& track, int32_t iSector, const GPUTPCTrack* sectorTr);
+  GPUd() int32_t RefitSectorTrack(GPUTPCGMSectorTrack& sectorTrack, MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_pointer> inTrack, float alpha, int32_t sector);
+  GPUd() void SetTrackClusterZT(GPUTPCGMSectorTrack& track, int32_t iSector, MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_pointer> sectorTr);
 
   int32_t CheckSectors();
   GPUd() void RefitSectorTracks(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, int32_t iSector);
diff --git a/GPU/GPUTracking/Merger/GPUTPCGMSectorTrack.cxx b/GPU/GPUTracking/Merger/GPUTPCGMSectorTrack.cxx
index d86a63ec61c53..d7acdc56b87d8 100644
--- a/GPU/GPUTracking/Merger/GPUTPCGMSectorTrack.cxx
+++ b/GPU/GPUTracking/Merger/GPUTPCGMSectorTrack.cxx
@@ -24,9 +24,9 @@
 using namespace o2::gpu;
 using namespace o2::tpc;
 
-GPUd() void GPUTPCGMSectorTrack::Set(const GPUTPCGMMerger* merger, const GPUTPCTrack* sectorTr, float alpha, int32_t sector)
+GPUd() void GPUTPCGMSectorTrack::Set(const GPUTPCGMMerger* merger, MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_pointer> sectorTr, float alpha, int32_t sector)
 {
-  GPUTPCBaseTrackParamSkeleton<MemLayout::const_reference> t = sectorTr->Param();
+  MemLayout::wrapper<GPUTPCBaseTrackParamSkeleton, MemLayout::const_reference> t = (*sectorTr).Param();
   mOrigTrack = sectorTr;
   mParam.mX = t.GetX();
   mParam.mY = t.GetY();
@@ -43,10 +43,10 @@ GPUd() void GPUTPCGMSectorTrack::Set(const GPUTPCGMMerger* merger, const GPUTPCT
   } else {
     mTZOffset = merger->GetConstantMem()->calibObjects.fastTransformHelper->getCorrMap()->convZOffsetToVertexTime(sector, t.GetZOffset(), merger->Param().continuousMaxTimeBin);
   }
-  mNClusters = sectorTr->NHits();
+  mNClusters = (*sectorTr).NHits();
 }
 
-GPUd() void GPUTPCGMSectorTrack::Set(const GPUTPCGMTrackParam& trk, const GPUTPCTrack* sectorTr, float alpha, int32_t sector)
+GPUd() void GPUTPCGMSectorTrack::Set(const GPUTPCGMTrackParam& trk, MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_pointer> sectorTr, float alpha, int32_t sector)
 {
   mOrigTrack = sectorTr;
   mParam.mX = trk.GetX();
@@ -60,7 +60,7 @@ GPUd() void GPUTPCGMSectorTrack::Set(const GPUTPCGMTrackParam& trk, const GPUTPC
   mAlpha = alpha;
   mSector = sector;
   mTZOffset = trk.GetTZOffset();
-  mNClusters = sectorTr->NHits();
+  mNClusters = (*sectorTr).NHits();
   mParam.mC0 = trk.GetCov(0);
   mParam.mC2 = trk.GetCov(2);
   mParam.mC3 = trk.GetCov(3);
@@ -96,11 +96,11 @@ GPUd() void GPUTPCGMSectorTrack::SetParam2(const GPUTPCGMTrackParam& trk)
 GPUd() bool GPUTPCGMSectorTrack::FilterErrors(const GPUTPCGMMerger* merger, int32_t iSector, float maxSinPhi, float sinPhiMargin)
 {
   float lastX;
-  // float lastX = GPUTPCGeometry::Row2X(mOrigTrack->Cluster(mOrigTrack->NClusters() - 1).GetRow()); // TODO: Why is this needed to be set below, Row2X should work, but looses some tracks
+  // float lastX = GPUTPCGeometry::Row2X((*mOrigTrack).Cluster((*mOrigTrack).NClusters() - 1).GetRow()); // TODO: Why is this needed to be set below, Row2X should work, but looses some tracks
   float y, z;
   int32_t row, index;
   const GPUTPCTracker& trk = merger->GetConstantMem()->tpcTrackers[iSector];
-  const GPUTPCHitId& ic = trk.TrackHits()[mOrigTrack->FirstHitID() + mOrigTrack->NHits() - 1];
+  const GPUTPCHitId& ic = trk.TrackHits()[(*mOrigTrack).FirstHitID() + (*mOrigTrack).NHits() - 1];
   index = trk.Data().ClusterDataIndex(trk.Data().Row(ic.RowIndex()), ic.HitIndex()) + merger->GetConstantMem()->ioPtrs.clustersNative->clusterOffset[iSector][0];
   row = ic.RowIndex();
   const ClusterNative& cl = merger->GetConstantMem()->ioPtrs.clustersNative->clustersLinear[index];
@@ -493,7 +493,7 @@ GPUd() bool GPUTPCGMSectorTrack::TransportToXAlpha(GPUTPCGMMerger* merger, float
 
 GPUd() void GPUTPCGMSectorTrack::CopyBaseTrackCov()
 {
-  const float* GPUrestrict() cov = mOrigTrack -> Param().mC;
+  const float* GPUrestrict() cov = (*mOrigTrack).Param().mC;
   mParam.mC0 = cov[0];
   mParam.mC2 = cov[2];
   mParam.mC3 = cov[3];
diff --git a/GPU/GPUTracking/Merger/GPUTPCGMSectorTrack.h b/GPU/GPUTracking/Merger/GPUTPCGMSectorTrack.h
index 27e4a89300ca4..412e5fe16d748 100644
--- a/GPU/GPUTracking/Merger/GPUTPCGMSectorTrack.h
+++ b/GPU/GPUTracking/Merger/GPUTPCGMSectorTrack.h
@@ -19,6 +19,7 @@
 #include "GPUTPCGMTrackParam.h"
 #include "GPUCommonMath.h"
 #include "GPUO2DataTypes.h"
+#include "MemLayout.h"
 
 namespace o2::gpu
 {
@@ -45,7 +46,7 @@ class GPUTPCGMSectorTrack
   {
     return (i < 2) ? mSegmentNeighbour[i] : mNeighbour[i - 2];
   }
-  GPUd() const GPUTPCTrack* OrigTrack() const { return mOrigTrack; }
+  GPUd() MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_pointer> OrigTrack() const { return mOrigTrack; }
   GPUd() float X() const { return mParam.mX; }
   GPUd() float Y() const { return mParam.mY; }
   GPUd() float Z() const { return mParam.mZ; }
@@ -73,9 +74,9 @@ class GPUTPCGMSectorTrack
     mClusterZT[1] = v2;
   }
 
-  GPUd() void Set(const GPUTPCGMTrackParam& trk, const GPUTPCTrack* sectorTr, float alpha, int32_t sector);
+  GPUd() void Set(const GPUTPCGMTrackParam& trk, MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_pointer> sectorTr, float alpha, int32_t sector);
   GPUd() void SetParam2(const GPUTPCGMTrackParam& trk);
-  GPUd() void Set(const GPUTPCGMMerger* merger, const GPUTPCTrack* sectorTr, float alpha, int32_t sector);
+  GPUd() void Set(const GPUTPCGMMerger* merger, MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_pointer> sectorTr, float alpha, int32_t sector);
   GPUd() void UseParam2() { mParam = mParam2; }
   GPUd() void SetX2(float v) { mParam2.mX = v; }
   GPUd() float X2() const { return mParam2.mX; }
@@ -124,7 +125,7 @@ class GPUTPCGMSectorTrack
   };
 
  private:
-  const GPUTPCTrack* mOrigTrack;    // pointer to original sector track
+  MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_pointer> mOrigTrack;    // pointer to original sector track
   sectorTrackParam mParam;          // Track parameters
   sectorTrackParam mParam2;         // Parameters at other side
   float mTZOffset;                  // Z offset with early transform, T offset otherwise
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCBaseTrackParam.h b/GPU/GPUTracking/SectorTracker/GPUTPCBaseTrackParam.h
index 56b34896b9000..e8db8f131a7be 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCBaseTrackParam.h
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCBaseTrackParam.h
@@ -45,11 +45,10 @@ struct Parameters {
  * This class is used for transfer between tracker and merger and does not contain the covariance matrice
  */
 template <template <class> class F>
-struct GPUTPCBaseTrackParamSkeleton : public MemLayout::CRTP<GPUTPCBaseTrackParamSkeleton, F>
+struct GPUTPCBaseTrackParamSkeleton
 {
-  using Base = MemLayout::CRTP<GPUTPCBaseTrackParamSkeleton, F>;
-  using Base::operator=;
-  MEMLAYOUT_MEMBERFUNCTIONS(GPUTPCBaseTrackParamSkeleton, F, mX, mC, mZOffset, mP)
+  MEMLAYOUT_APPLY_UNARY(mX, mC, mZOffset, mP)
+  MEMLAYOUT_APPLY_BINARY(GPUTPCBaseTrackParamSkeleton, MEMLAYOUT_EXPAND(mX), MEMLAYOUT_EXPAND(mC), MEMLAYOUT_EXPAND(mZOffset), MEMLAYOUT_EXPAND(mP))
 
   GPUd() float X() const { return mX; }
   GPUd() float Y() const { return mP[0]; }
@@ -92,32 +91,6 @@ struct GPUTPCBaseTrackParamSkeleton : public MemLayout::CRTP<GPUTPCBaseTrackPara
   GPUd() void SetQPt(float v) { mP[4] = v; }
   GPUd() void SetZOffset(float v) { mZOffset = v; }
 
-  // Needed for iterators and std::sort
-  /*template <template <class> class F_in>
-  constexpr GPUTPCBaseTrackParamSkeleton& operator=(GPUTPCBaseTrackParamSkeleton<F_in> other) {
-    mX = other.mX;
-    mC = other.mC;
-    mZOffset = other.mZOffset;
-    mP = other.mP;
-    return *this;
-  }*/
-  //GPUTPCBaseTrackParamSkeleton() = default;
-  //GPUTPCBaseTrackParamSkeleton(F<float> X, F<Covariance> C, F<float> ZOffset, F<Parameters> P) : mX(X), mC(C), mZOffset(ZOffset), mP(P) { }
-  /*GPUTPCBaseTrackParamSkeleton(const GPUTPCBaseTrackParamSkeleton& other) = default;
-  GPUTPCBaseTrackParamSkeleton& operator=(const GPUTPCBaseTrackParamSkeleton& other) {
-    mX = other.mX;
-    mC = other.mC;
-    mZOffset = other.mZOffset;
-    mP = other.mP;
-    return *this;
-  }
-  friend void swap(GPUTPCBaseTrackParamSkeleton& a, GPUTPCBaseTrackParamSkeleton& b) {
-    std::swap(a.mX, b.mX);
-    std::swap(a.mC, b.mC);
-    std::swap(a.mZOffset, b.mZOffset);
-    std::swap(a.mP, b.mP);
-  }*/
-
   // WARNING, Track Param Data is copied in the GPU Tracklet Constructor element by element instead of using copy constructor!!!
   // This is neccessary for performance reasons!!!
   // Changes to Elements of this class therefore must also be applied to TrackletConstructor!!!
@@ -127,19 +100,16 @@ struct GPUTPCBaseTrackParamSkeleton : public MemLayout::CRTP<GPUTPCBaseTrackPara
   F<Parameters> mP;    // 'active' track parameters: Y, Z, SinPhi, DzDs, q/Pt
 };
 
-template <
-    template <class> class F_left,
-    template <class> class F_right,
-    class FunctionObject
->
-constexpr void memberwise(GPUTPCBaseTrackParamSkeleton<F_left>& left, GPUTPCBaseTrackParamSkeleton<F_right>& right, FunctionObject&& f) {
-    f(left.mX, right.mX);
-    f(left.mC, right.mC);
-    f(left.mZOffset, right.mZOffset);
-    f(left.mP, right.mP);
+// Needed for sorting
+constexpr void swap(GPUTPCBaseTrackParamSkeleton<MemLayout::reference> a, GPUTPCBaseTrackParamSkeleton<MemLayout::reference> b) {
+    using std::swap;
+    swap(a.mX, b.mX);
+    swap(a.mC, b.mC);
+    swap(a.mZOffset, b.mZOffset);
+    swap(a.mP, b.mP);
 }
 
-using GPUTPCBaseTrackParam = GPUTPCBaseTrackParamSkeleton<MemLayout::value>;
+using GPUTPCBaseTrackParam = MemLayout::wrapper<GPUTPCBaseTrackParamSkeleton, MemLayout::value>;
 
 } // namespace o2::gpu
 
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCExtrapolationTracking.cxx b/GPU/GPUTracking/SectorTracker/GPUTPCExtrapolationTracking.cxx
index 58325877fdd2a..81e91e70d0737 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCExtrapolationTracking.cxx
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCExtrapolationTracking.cxx
@@ -20,6 +20,7 @@
 #include "GPUTPCTracker.h"
 #include "GPUCommonMath.h"
 #include "GPUParam.inc"
+#include "MemLayout.h"
 
 using namespace o2::gpu;
 
@@ -31,7 +32,7 @@ GPUd() int32_t GPUTPCExtrapolationTracking::PerformExtrapolationTrackingRun(GPUT
     (float) Data().HitDataY(Row(mTrackHits[Tracks()[iTrack].FirstHitID() + j].RowIndex()), mTrackHits[Tracks()[iTrack].FirstHitID() + j].HitIndex()) * Row(mTrackHits[Tracks()[iTrack].FirstHitID() + j].RowIndex()).HstepY() + Row(mTrackHits[Tracks()[iTrack].FirstHitID() + j].RowIndex()).Grid().YMin());
   }*/
 
-  GPUTPCTrackParamSkeleton<MemLayout::value> tParam;
+  MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::value> tParam;
   tParam.InitParam();
   tParam.SetCov(0, 0.05f);
   tParam.SetCov(2, 0.05f);
@@ -109,7 +110,8 @@ GPUd() int32_t GPUTPCExtrapolationTracking::PerformExtrapolationTrackingRun(GPUT
         rowIndex--;
       }
     }
-    GPUTPCTrack& GPUrestrict() track = tracker.Tracks()[trackId];
+    //GPUTPCTrack& GPUrestrict() track = tracker.Tracks()[trackId];
+    MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::reference> track = tracker.Tracks()[trackId]; // reference_restrict
     track.SetParam(tParam.GetParam());
     track.SetNHits(nHits);
     track.SetFirstHitID(hitId);
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCSectorDebugSortKernels.cxx b/GPU/GPUTracking/SectorTracker/GPUTPCSectorDebugSortKernels.cxx
index 7981ef5af26d8..d2833442956be 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCSectorDebugSortKernels.cxx
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCSectorDebugSortKernels.cxx
@@ -91,7 +91,7 @@ GPUdii() void GPUTPCSectorDebugSortKernels::Thread<GPUTPCSectorDebugSortKernels:
   if (iThread || iBlock) {
     return;
   }
-  auto sorter = [&tracker](const GPUTPCTrack& trk1, const GPUTPCTrack& trk2) {
+  auto sorter = [&tracker](MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_reference> trk1, MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_reference> trk2) {
     if (trk1.NHits() != trk2.NHits()) {
       return trk1.NHits() > trk2.NHits();
     }
@@ -106,6 +106,6 @@ GPUdii() void GPUTPCSectorDebugSortKernels::Thread<GPUTPCSectorDebugSortKernels:
     }
     return tracker.TrackHits()[trk1.FirstHitID()].HitIndex() > tracker.TrackHits()[trk2.FirstHitID()].HitIndex();
   };
-  GPUCommonAlgorithm::sortDeviceDynamic(tracker.Tracks(), tracker.Tracks() + tracker.CommonMemory()->nLocalTracks, sorter);
-  GPUCommonAlgorithm::sortDeviceDynamic(tracker.Tracks() + tracker.CommonMemory()->nLocalTracks, tracker.Tracks() + *tracker.NTracks(), sorter);
+  GPUCommonAlgorithm::sortDeviceDynamic(tracker.Tracks(), &tracker.Tracks()[tracker.CommonMemory()->nLocalTracks], sorter);
+  GPUCommonAlgorithm::sortDeviceDynamic(&tracker.Tracks()[tracker.CommonMemory()->nLocalTracks], &tracker.Tracks()[*tracker.NTracks()], sorter);
 }
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTrack.h b/GPU/GPUTracking/SectorTracker/GPUTPCTrack.h
index de4b12b6f0d4a..66cc17f2311e5 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCTrack.h
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCTrack.h
@@ -29,55 +29,60 @@ namespace o2::gpu
  * The track parameters at both ends are stored separately in the GPUTPCEndPoint class
  */
 template <template <class> class F>
-class GPUTPCTrackSkeleton : public MemLayout::CRTP<GPUTPCTrackSkeleton, F>
+class GPUTPCTrackSkeleton
 {
  public:
-    using Base = MemLayout::CRTP<GPUTPCTrackSkeleton, F>;
-    using Base::operator=;
-  MEMLAYOUT_MEMBERFUNCTIONS(GPUTPCTrackSkeleton, F, mFirstHitID, mNHits, mLocalTrackId, mParam)
-    
-#if !defined(GPUCA_GPUCODE)
-  constexpr GPUTPCTrackSkeleton() : mFirstHitID(0), mNHits(0), mLocalTrackId(-1), mParam() { }
-  constexpr GPUTPCTrackSkeleton(
-    F<int32_t> FirstHitID,
-    F<int32_t> NHits,
-    F<int32_t> LocalTrackId,
-    GPUTPCBaseTrackParamSkeleton<F> Param
-  ) : mFirstHitID(FirstHitID), mNHits(NHits), mLocalTrackId(LocalTrackId), mParam(Param) { }
-  ~GPUTPCTrackSkeleton() = default;
-#endif //! GPUCA_GPUCODE
+    MEMLAYOUT_APPLY_UNARY(mFirstHitID, mNHits, mLocalTrackId, mParam)
+    MEMLAYOUT_APPLY_BINARY(GPUTPCTrackSkeleton, MEMLAYOUT_EXPAND(mFirstHitID), MEMLAYOUT_EXPAND(mNHits), MEMLAYOUT_EXPAND(mLocalTrackId), MEMLAYOUT_EXPAND(mParam))
+
+//#if !defined(GPUCA_GPUCODE)
+//  ~GPUTPCTrackSkeleton() = default;
+//#endif //! GPUCA_GPUCODE
 
   GPUhd() int32_t NHits() const { return mNHits; }
   GPUhd() int32_t LocalTrackId() const { return mLocalTrackId; }
   GPUhd() int32_t FirstHitID() const { return mFirstHitID; }
-  GPUhd() GPUTPCBaseTrackParamSkeleton<MemLayout::const_reference> Param() const { return mParam; }
+  GPUhd() MemLayout::wrapper<GPUTPCBaseTrackParamSkeleton, MemLayout::const_reference> Param() const { return mParam; }
 
   GPUhd() void SetNHits(int32_t v) { mNHits = v; }
   GPUhd() void SetLocalTrackId(int32_t v) { mLocalTrackId = v; }
   GPUhd() void SetFirstHitID(int32_t v) { mFirstHitID = v; }
-  GPUhd() void SetParam(GPUTPCBaseTrackParamSkeleton<MemLayout::const_reference> v) { mParam = v; }
+  GPUhd() void SetParam(MemLayout::wrapper<GPUTPCBaseTrackParamSkeleton, MemLayout::const_reference> v) { mParam = v; }
 
  //private:
   F<int32_t> mFirstHitID;         // index of the first track cell in the track->cell pointer array
   F<int32_t> mNHits;              // number of track cells
   F<int32_t> mLocalTrackId;       // Id of local track this extrapolated track belongs to, index of this track itself if it is a local track
-  GPUTPCBaseTrackParamSkeleton<F> mParam; // track parameters
+  MemLayout::wrapper<GPUTPCBaseTrackParamSkeleton, F> mParam; // track parameters
 };
 
-template <
-    template <class> class F_left,
-    template <class> class F_right,
-    class FunctionObject
->
-constexpr void memberwise(GPUTPCTrackSkeleton<F_left>& left, GPUTPCTrackSkeleton<F_right>& right, FunctionObject&& f) {
-    f(left.mFirstHitID, right.mFirstHitID);
-    f(left.mNHits, right.mNHits);
-    f(left.mLocalTrackId, right.mLocalTrackId);
-    f(left.mParam, right.mParam);
+// Needed for sorting
+constexpr void swap(GPUTPCTrackSkeleton<MemLayout::reference> a, GPUTPCTrackSkeleton<MemLayout::reference> b) {
+    using std::swap;
+    swap(a.mFirstHitID, b.mFirstHitID);
+    swap(a.mNHits, b.mNHits);
+    swap(a.mLocalTrackId, b.mLocalTrackId);
+    swap(a.mParam, b.mParam);
 }
 
-using GPUTPCTrack = GPUTPCTrackSkeleton<MemLayout::value>;
+using GPUTPCTrack = MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::value>;
 
 } // namespace o2::gpu
 
+namespace std {
+
+template<class T> class iterator_traits;
+class random_access_iterator_tag;
+
+template<>
+struct iterator_traits<MemLayout::wrapper<o2::gpu::GPUTPCTrackSkeleton, MemLayout::pointer>> {
+    using iterator_category = random_access_iterator_tag;
+    using difference_type = MemLayout::ptrdiff_t;
+    using value_type = MemLayout::wrapper<o2::gpu::GPUTPCTrackSkeleton, MemLayout::value>;
+    using pointer = void; //MemLayout::wrapper<o2::gpu::GPUTPCTrackSkeleton, MemLayout::pointer>;
+    using reference = MemLayout::wrapper<o2::gpu::GPUTPCTrackSkeleton, MemLayout::reference>;
+};
+
+}
+
 #endif // GPUTPCTRACK_H
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTrackLinearisation.h b/GPU/GPUTracking/SectorTracker/GPUTPCTrackLinearisation.h
index 775083f4a2ec6..b13033632bdd7 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCTrackLinearisation.h
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCTrackLinearisation.h
@@ -40,7 +40,7 @@ class GPUTPCTrackLinearisation
   GPUTPCTrackLinearisation() : mSinPhi(0), mCosPhi(1), mDzDs(0), mQPt(0) {}
   GPUTPCTrackLinearisation(float SinPhi1, float CosPhi1, float DzDs1, float QPt1) : mSinPhi(SinPhi1), mCosPhi(CosPhi1), mDzDs(DzDs1), mQPt(QPt1) {}
 
-  GPUd() GPUTPCTrackLinearisation(GPUTPCTrackParamSkeleton<MemLayout::const_reference_restrict> t);
+  GPUd() GPUTPCTrackLinearisation(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::const_reference> t); // const_reference_restrict
 
   GPUd() void Set(float SinPhi1, float CosPhi1, float DzDs1, float QPt1);
 
@@ -66,7 +66,7 @@ class GPUTPCTrackLinearisation
   float mQPt;    // QPt
 };
 
-GPUdi() GPUTPCTrackLinearisation::GPUTPCTrackLinearisation(GPUTPCTrackParamSkeleton<MemLayout::const_reference_restrict> t) : mSinPhi(t.SinPhi()), mCosPhi(0), mDzDs(t.DzDs()), mQPt(t.QPt())
+GPUdi() GPUTPCTrackLinearisation::GPUTPCTrackLinearisation(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::const_reference> t) : mSinPhi(t.SinPhi()), mCosPhi(0), mDzDs(t.DzDs()), mQPt(t.QPt()) // const_reference_restrict
 {
   if (mSinPhi > GPUCA_MAX_SIN_PHI) {
     mSinPhi = GPUCA_MAX_SIN_PHI;
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTrackParam.cxx b/GPU/GPUTracking/SectorTracker/GPUTPCTrackParam.cxx
index d141ce7e4a7f5..42de67eb6b02b 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCTrackParam.cxx
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCTrackParam.cxx
@@ -31,7 +31,7 @@ using namespace o2::gpu;
 //
 
 template <template <class> class F>
-GPUd() float GPUTPCTrackParamSkeleton<F>::GetDist2(GPUTPCTrackParamSkeleton<MemLayout::const_reference_restrict> t) const
+GPUd() float GPUTPCTrackParamSkeleton<F>::GetDist2(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::const_reference> t) const // const_reference_restrict
 {
   // get squared distance between tracks
 
@@ -42,7 +42,7 @@ GPUd() float GPUTPCTrackParamSkeleton<F>::GetDist2(GPUTPCTrackParamSkeleton<MemL
 }
 
 template <template <class> class F>
-GPUd() float GPUTPCTrackParamSkeleton<F>::GetDistXZ2(GPUTPCTrackParamSkeleton<MemLayout::const_reference_restrict> t) const
+GPUd() float GPUTPCTrackParamSkeleton<F>::GetDistXZ2(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::const_reference> t) const // const_reference_restrict
 {
   // get squared distance between tracks in X&Z
 
@@ -305,7 +305,7 @@ template <template <class> class F>
 GPUd() bool GPUTPCTrackParamSkeleton<F>::TransportToX(float x, float Bz, float maxSinPhi)
 {
   //* Transport the track parameters to X=x
-  GPUTPCTrackLinearisation t0(*this);
+  GPUTPCTrackLinearisation t0(*static_cast<MemLayout::wrapper<GPUTPCTrackParamSkeleton, F>*>(this));
   return TransportToX(x, t0, Bz, maxSinPhi);
 }
 
@@ -332,7 +332,7 @@ template <template <class> class F>
 GPUd() bool GPUTPCTrackParamSkeleton<F>::TransportToXWithMaterial(float x, GPUTPCTrackFitParam& GPUrestrict() par, float Bz, float maxSinPhi)
 {
   //* Transport the track parameters to X=x  taking into account material budget
-  GPUTPCTrackLinearisation t0(*this);
+  GPUTPCTrackLinearisation t0(*static_cast<MemLayout::wrapper<GPUTPCTrackParamSkeleton, F>*>(this));
   return TransportToXWithMaterial(x, t0, par, Bz, maxSinPhi);
 }
 
@@ -879,5 +879,5 @@ GPUd() int32_t GPUTPCTrackParamSkeleton<F>::GetPropagatedYZ(float bz, float x, f
 namespace o2::gpu {
   template class GPUTPCTrackParamSkeleton<MemLayout::value>;
   template class GPUTPCTrackParamSkeleton<MemLayout::reference>;
-  template class GPUTPCTrackParamSkeleton<MemLayout::reference_restrict>;
+  //template class GPUTPCTrackParamSkeleton<MemLayout::reference_restrict>;
 } // namespace o2::gpu
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTrackParam.h b/GPU/GPUTracking/SectorTracker/GPUTPCTrackParam.h
index 5536114c88752..e7b6b02efb722 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCTrackParam.h
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCTrackParam.h
@@ -33,19 +33,18 @@ class GPUTPCTrackLinearisation;
  *
  */
 template <template <class> class F>
-class GPUTPCTrackParamSkeleton : public MemLayout::CRTP<GPUTPCTrackParamSkeleton, F>
+class GPUTPCTrackParamSkeleton
 {
  public:
-    using Base = MemLayout::CRTP<GPUTPCTrackParamSkeleton, F>;
-    using Base::operator=;
-  MEMLAYOUT_MEMBERFUNCTIONS(GPUTPCTrackParamSkeleton, F, mParam, mSignCosPhi, mChi2, mNDF)
+  MEMLAYOUT_APPLY_UNARY(mParam, mSignCosPhi, mChi2, mNDF)
+  MEMLAYOUT_APPLY_BINARY(GPUTPCTrackParamSkeleton, MEMLAYOUT_EXPAND(mParam), MEMLAYOUT_EXPAND(mSignCosPhi), MEMLAYOUT_EXPAND(mChi2), MEMLAYOUT_EXPAND(mNDF))
 
   struct GPUTPCTrackFitParam {
     float bethe, e, theta2, EP2, sigmadE2, k22, k33, k43, k44; // parameters
   };
 
-  GPUd() GPUTPCBaseTrackParamSkeleton<MemLayout::const_reference> GetParam() const { return mParam; }
-  GPUd() void SetParam(GPUTPCBaseTrackParamSkeleton<MemLayout::const_reference> v);
+  GPUd() MemLayout::wrapper<GPUTPCBaseTrackParamSkeleton, MemLayout::const_reference> GetParam() const { return mParam; }
+  GPUd() void SetParam(MemLayout::wrapper<GPUTPCBaseTrackParamSkeleton, MemLayout::const_reference> v);
   GPUd() void InitParam();
 
   GPUd() float X() const { return mParam.X(); }
@@ -99,8 +98,8 @@ class GPUTPCTrackParamSkeleton : public MemLayout::CRTP<GPUTPCTrackParamSkeleton
   GPUd() void SetChi2(float v) { mChi2 = v; }
   GPUd() void SetNDF(int32_t v) { mNDF = v; }
 
-  GPUd() float GetDist2(GPUTPCTrackParamSkeleton<MemLayout::const_reference_restrict> t) const;
-  GPUd() float GetDistXZ2(GPUTPCTrackParamSkeleton<MemLayout::const_reference_restrict> t) const;
+  GPUd() float GetDist2(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::const_reference> t) const; // const_reference_restrict
+  GPUd() float GetDistXZ2(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::const_reference> t) const; // const_reference_restrict
 
   GPUd() float GetS(float x, float y, float Bz) const;
 
@@ -146,10 +145,10 @@ class GPUTPCTrackParamSkeleton : public MemLayout::CRTP<GPUTPCTrackParamSkeleton
 
   GPUd() void Print() const;
 
-#ifndef GPUCA_GPUCODE
+//#ifndef GPUCA_GPUCODE
  //private:
-#endif                         //! GPUCA_GPUCODE
-  GPUTPCBaseTrackParamSkeleton<F> mParam;
+//#endif                         //! GPUCA_GPUCODE
+  MemLayout::wrapper<GPUTPCBaseTrackParamSkeleton, F> mParam;
 
  //private:
   // WARNING, Track Param Data is copied in the GPU Tracklet Constructor element by element instead of using copy constructor!!!
@@ -160,20 +159,8 @@ class GPUTPCTrackParamSkeleton : public MemLayout::CRTP<GPUTPCTrackParamSkeleton
   F<int32_t> mNDF;      // the Number of Degrees of Freedom
 };
 
-template <
-    template <class> class F_left,
-    template <class> class F_right,
-    class FunctionObject
->
-constexpr void memberwise(GPUTPCTrackParamSkeleton<F_left>& left, GPUTPCTrackParamSkeleton<F_right>& right, FunctionObject&& f) {
-    f(left.mParam, right.mParam);
-    f(left.mSignCosPhi, right.mSignCosPhi);
-    f(left.mChi2, right.mChi2);
-    f(left.mNDF, right.mNDF);
-}
-
 template <template <class> class F>
-GPUd() void GPUTPCTrackParamSkeleton<F>::SetParam(GPUTPCBaseTrackParamSkeleton<MemLayout::const_reference> v) { mParam = v; }
+GPUd() void GPUTPCTrackParamSkeleton<F>::SetParam(MemLayout::wrapper<GPUTPCBaseTrackParamSkeleton, MemLayout::const_reference> v) { mParam = v; }
 
 template <template <class> class F>
 GPUd() void GPUTPCTrackParamSkeleton<F>::InitParam()
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTracker.cxx b/GPU/GPUTracking/SectorTracker/GPUTPCTracker.cxx
index a1c1f0bb8da6b..bb3f09aa6f93d 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCTracker.cxx
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCTracker.cxx
@@ -134,7 +134,14 @@ GPUhd() void* GPUTPCTracker::SetPointersTracklets(void* mem)
 
 GPUhd() void* GPUTPCTracker::SetPointersOutput(void* mem)
 {
-  computePointerWithAlignment(mem, mTracks, mNMaxTracks);
+  computePointerWithAlignment(mem, mTracks.mFirstHitID, mNMaxTracklets);
+  computePointerWithAlignment(mem, mTracks.mNHits, mNMaxTracklets);
+  computePointerWithAlignment(mem, mTracks.mLocalTrackId, mNMaxTracklets);
+  computePointerWithAlignment(mem, mTracks.mParam.mX, mNMaxTracklets);
+  computePointerWithAlignment(mem, mTracks.mParam.mC, mNMaxTracklets);
+  computePointerWithAlignment(mem, mTracks.mParam.mZOffset, mNMaxTracklets);
+  computePointerWithAlignment(mem, mTracks.mParam.mP, mNMaxTracklets);
+  //computePointerWithAlignment(mem, mTracks, mNMaxTracks);
   computePointerWithAlignment(mem, mTrackHits, mNMaxTrackHits);
   return mem;
 }
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTracker.h b/GPU/GPUTracking/SectorTracker/GPUTPCTracker.h
index 5af0939188391..8fe7b0832f046 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCTracker.h
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCTracker.h
@@ -25,6 +25,7 @@
 #include "GPUTPCTrackingData.h"
 #include "GPUTPCTrackParam.h"
 #include "GPUTPCTracklet.h"
+#include "GPUTPCTrack.h"
 #include "GPUProcessor.h"
 #include "MemLayout.h"
 
@@ -62,7 +63,7 @@ class GPUTPCTracker : public GPUProcessor
 #endif
 
   template <MemLayout::Flag layout>
-  using TrackletArrayType = MemLayout::wrapper<GPUTPCTrackletSkeleton, MemLayout::pointer, layout>::type;
+  using TrackletArrayType = MemLayout::interface<GPUTPCTrackletSkeleton, MemLayout::pointer, layout>::type;
 
   struct commonMemoryStruct {
     GPUAtomic(uint32_t) nStartHits = 0; // number of start hits
@@ -84,13 +85,13 @@ class GPUTPCTracker : public GPUProcessor
     return (mCommonMem);
   }
 
-  GPUdi() static void GetErrors2Seeding(const GPUParam& param, char sector, int32_t iRow, GPUTPCTrackParamSkeleton<MemLayout::const_reference> t, float time, float& ErrY2, float& ErrZ2)
+  GPUdi() static void GetErrors2Seeding(const GPUParam& param, char sector, int32_t iRow, MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::const_reference> t, float time, float& ErrY2, float& ErrZ2)
   {
     // param.GetClusterErrors2(sector, iRow, param.GetContinuousTracking() != 0. ? 125.f : t.Z(), t.SinPhi(), t.DzDs(), time, 0.f, 0.f, ErrY2, ErrZ2);
     param.GetClusterErrorsSeeding2(sector, iRow, param.par.continuousTracking != 0.f ? 125.f : t.Z(), t.SinPhi(), t.DzDs(), time, ErrY2, ErrZ2);
   }
 
-  GPUdi() void GetErrors2Seeding(int32_t iRow, GPUTPCTrackParamSkeleton<MemLayout::const_reference> t, float time, float& ErrY2, float& ErrZ2) const
+  GPUdi() void GetErrors2Seeding(int32_t iRow, MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::const_reference> t, float time, float& ErrY2, float& ErrZ2) const
   {
     // Param().GetClusterErrors2(mISector, iRow, Param().GetContinuousTracking() != 0. ? 125.f : t.Z(), t.SinPhi(), t.DzDs(), time, 0.f, 0.f, ErrY2, ErrZ2);
     Param().GetClusterErrorsSeeding2(mISector, iRow, Param().par.continuousTracking != 0.f ? 125.f : t.Z(), t.SinPhi(), t.DzDs(), time, ErrY2, ErrZ2);
@@ -192,14 +193,14 @@ class GPUTPCTracker : public GPUProcessor
   GPUhd() GPUglobalref() GPUTPCHitId* TrackletStartHits() { return mTrackletStartHits; }
   GPUhd() GPUglobalref() GPUTPCHitId* TrackletTmpStartHits() const { return mTrackletTmpStartHits; }
 
-  GPUhd() GPUglobalref() GPUTPCTrackletSkeleton<MemLayout::reference_restrict> Tracklet(int32_t i) { return mTracklets[i]; }
+  GPUhd() GPUglobalref() MemLayout::wrapper<GPUTPCTrackletSkeleton, MemLayout::reference> Tracklet(int32_t i) { return mTracklets[i]; } // reference_restrict
 
   template <MemLayout::Flag layout>
   GPUhd() GPUglobalref() TrackletArrayType<layout> Tracklets() const { return mTracklets; }
   GPUhd() GPUglobalref() calink* TrackletRowHits() const { return mTrackletRowHits; }
 
   GPUhd() GPUglobalref() GPUAtomic(uint32_t) * NTracks() const { return &mCommonMem->nTracks; }
-  GPUhd() GPUglobalref() GPUTPCTrackSkeleton<MemLayout::value>* Tracks() const { return mTracks; }
+  GPUhd() GPUglobalref() MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::pointer> Tracks() const { return mTracks; }
   GPUhd() GPUglobalref() GPUAtomic(uint32_t) * NTrackHits() const { return &mCommonMem->nTrackHits; }
   GPUhd() GPUglobalref() GPUTPCHitId* TrackHits() const { return mTrackHits; }
 
@@ -251,7 +252,7 @@ class GPUTPCTracker : public GPUProcessor
   GPUglobalref() GPUTPCHitId* mTrackletStartHits = nullptr; // start hits for the tracklets
   GPUglobalref() TrackletArrayType<MemLayout::Flag::soa> mTracklets; // tracklets
   GPUglobalref() calink* mTrackletRowHits = nullptr;        // Hits for each Tracklet in each row
-  GPUglobalref() GPUTPCTrackSkeleton<MemLayout::value>* mTracks = nullptr; // reconstructed tracks
+  GPUglobalref() MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::pointer> mTracks; // reconstructed tracks
   GPUglobalref() GPUTPCHitId* mTrackHits = nullptr;         // array of track hit numbers
 
   static int32_t StarthitSortComparison(const void* a, const void* b);
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTrackerDump.cxx b/GPU/GPUTracking/SectorTracker/GPUTPCTrackerDump.cxx
index 3c3f113aa8c21..f291c9cc2fb9a 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCTrackerDump.cxx
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCTrackerDump.cxx
@@ -102,7 +102,7 @@ void GPUTPCTracker::DumpTrackHits(std::ostream& out)
     if (Tracks()[j].NHits() == 0) {
       continue;
     }
-    GPUTPCBaseTrackParamSkeleton<MemLayout::const_reference> p = Tracks()[j].Param();
+    MemLayout::wrapper<GPUTPCBaseTrackParamSkeleton, MemLayout::const_reference> p = Tracks()[j].Param();
     out << "  " << j << " x " << p.GetX() << " offset " << p.GetZOffset() << " y " << p.GetY() << " z " << p.GetZ() << " snp " << p.GetSinPhi() << " tgl " << p.GetDzDs() << " qpt " << p.GetQPt() << " - ";
     for (int32_t k = 0; k < 15; k++) {
       out << p.GetCov(k) << " ";
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTracklet.h b/GPU/GPUTracking/SectorTracker/GPUTPCTracklet.h
index 95bc83e8ee65f..df069cdb35d60 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCTracklet.h
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCTracklet.h
@@ -28,49 +28,36 @@ namespace o2::gpu
  * The class is dedicated for internal use by the GPUTPCTracker algorithm.
  */
 template <template <class> class F>
-class GPUTPCTrackletSkeleton : public MemLayout::CRTP<GPUTPCTrackletSkeleton, F>
+class GPUTPCTrackletSkeleton
 {
  public:
-  using Base = MemLayout::CRTP<GPUTPCTrackletSkeleton, F>;
-  using Base::operator=;
-  MEMLAYOUT_MEMBERFUNCTIONS(GPUTPCTrackletSkeleton, F, mFirstRow, mLastRow, mParam, mHitWeight, mFirstHit)
-#if !defined(GPUCA_GPUCODE)
+  MEMLAYOUT_APPLY_UNARY(mFirstRow, mLastRow, mParam, mHitWeight, mFirstHit)
+  MEMLAYOUT_APPLY_BINARY(GPUTPCTrackletSkeleton, MEMLAYOUT_EXPAND(mFirstRow), MEMLAYOUT_EXPAND(mLastRow), MEMLAYOUT_EXPAND(mParam), MEMLAYOUT_EXPAND(mHitWeight), MEMLAYOUT_EXPAND(mFirstHit))
+
+//#if !defined(GPUCA_GPUCODE)
   //GPUTPCTrackletSkeleton() : mFirstRow(0), mLastRow(0), mParam(), mHitWeight(0), mFirstHit(0) {};
-#endif //! GPUCA_GPUCODE
+//#endif //! GPUCA_GPUCODE
 
   GPUhd() int32_t FirstRow() const { return mFirstRow; }
   GPUhd() int32_t LastRow() const { return mLastRow; }
   GPUhd() int32_t HitWeight() const { return mHitWeight; }
   GPUhd() uint32_t FirstHit() const { return mFirstHit; }
-  GPUhd() GPUTPCBaseTrackParamSkeleton<MemLayout::const_reference> Param() const { return mParam; }
+  GPUhd() MemLayout::wrapper<GPUTPCBaseTrackParamSkeleton, MemLayout::const_reference> Param() const { return mParam; }
 
   GPUhd() void SetFirstRow(int32_t v) { mFirstRow = v; }
   GPUhd() void SetLastRow(int32_t v) { mLastRow = v; }
   GPUhd() void SetFirstHit(uint32_t v) { mFirstHit = v; }
-  GPUhd() void SetParam(GPUTPCBaseTrackParamSkeleton<MemLayout::const_reference> v) { mParam = v; }
+  GPUhd() void SetParam(MemLayout::wrapper<GPUTPCBaseTrackParamSkeleton, MemLayout::const_reference> v) { mParam = v; }
   GPUhd() void SetHitWeight(const int32_t w) { mHitWeight = w; }
 
 // private:
   F<int32_t> mFirstRow;           // first TPC row // TODO: We can use smaller data format here!
   F<int32_t> mLastRow;            // last TPC row
-  GPUTPCBaseTrackParamSkeleton<F> mParam; // tracklet parameters
+  MemLayout::wrapper<GPUTPCBaseTrackParamSkeleton, F> mParam; // tracklet parameters
   F<int32_t> mHitWeight;          // Hit Weight of Tracklet
   F<uint32_t> mFirstHit;          // first hit in row hit array
 };
 
-template <
-    template <class> class F_left,
-    template <class> class F_right,
-    class FunctionObject
->
-constexpr void memberwise(GPUTPCTrackletSkeleton<F_left>& left, GPUTPCTrackletSkeleton<F_right>& right, FunctionObject&& f) {
-    f(left.mFirstRow, right.mFirstRow);
-    f(left.mLastRow, right.mLastRow);
-    f(left.mParam, right.mParam);
-    f(left.mHitWeight, right.mHitWeight);
-    f(left.mFirstHit, right.mFirstHit);
-}
-
 } // namespace o2::gpu
 
 #endif // GPUTPCTRACKLET_H
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.cxx b/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.cxx
index b5c815a0ee9d6..777b55bbe5594 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.cxx
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.cxx
@@ -30,13 +30,13 @@
 
 using namespace o2::gpu;
 
-GPUdii() void GPUTPCTrackletConstructor::InitTracklet(GPUTPCTrackParamSkeleton<MemLayout::reference_restrict> tParam)
+GPUdii() void GPUTPCTrackletConstructor::InitTracklet(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference> tParam) // reference_restrict
 {
   // Initialize Tracklet Parameters using default values
   tParam.InitParam();
 }
 
-GPUd() bool GPUTPCTrackletConstructor::CheckCov(GPUTPCTrackParamSkeleton<MemLayout::reference_restrict> tParam)
+GPUd() bool GPUTPCTrackletConstructor::CheckCov(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference> tParam) // reference_restrict
 {
   bool ok = 1;
   const float* c = tParam.Cov();
@@ -53,7 +53,7 @@ GPUd() bool GPUTPCTrackletConstructor::CheckCov(GPUTPCTrackParamSkeleton<MemLayo
   return (ok);
 }
 
-GPUd() void GPUTPCTrackletConstructor::StoreTracklet(int32_t /*nBlocks*/, int32_t /*nThreads*/, int32_t /*iBlock*/, int32_t /*iThread*/, GPUsharedref() GPUSharedMemory& s, GPUTPCThreadMemory& GPUrestrict() r, GPUconstantref() GPUTPCTracker& GPUrestrict() tracker, GPUTPCTrackParamSkeleton<MemLayout::reference_restrict> tParam, calink* rowHits)
+GPUd() void GPUTPCTrackletConstructor::StoreTracklet(int32_t /*nBlocks*/, int32_t /*nThreads*/, int32_t /*iBlock*/, int32_t /*iThread*/, GPUsharedref() GPUSharedMemory& s, GPUTPCThreadMemory& GPUrestrict() r, GPUconstantref() GPUTPCTracker& GPUrestrict() tracker, MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference> tParam, calink* rowHits) // reference_restrict
 {
   // reconstruction of tracklets, tracklet store step
   const uint32_t nHits = r.mLastRow + 1 - r.mFirstRow;
@@ -79,7 +79,7 @@ GPUd() void GPUTPCTrackletConstructor::StoreTracklet(int32_t /*nBlocks*/, int32_
     return;
   }
 
-  GPUglobalref() GPUTPCTrackletSkeleton<MemLayout::reference_restrict> tracklet = tracker.Tracklet(itrout);
+  GPUglobalref() MemLayout::wrapper<GPUTPCTrackletSkeleton, MemLayout::reference> tracklet = tracker.Tracklet(itrout); // reference_restrict
 
   CADEBUG(printf("    Storing tracklet: %d rows\n", nHits));
 
@@ -104,7 +104,7 @@ GPUd() void GPUTPCTrackletConstructor::StoreTracklet(int32_t /*nBlocks*/, int32_
 }
 
 template <class T>
-GPUdic(2, 1) void GPUTPCTrackletConstructor::UpdateTracklet(int32_t /*nBlocks*/, int32_t /*nThreads*/, int32_t /*iBlock*/, int32_t /*iThread*/, GPUsharedref() T& s, GPUTPCThreadMemory& GPUrestrict() r, GPUconstantref() GPUTPCTracker& GPUrestrict() tracker, GPUTPCTrackParamSkeleton<MemLayout::reference_restrict> tParam, int32_t iRow, calink& rowHit, calink* rowHits)
+GPUdic(2, 1) void GPUTPCTrackletConstructor::UpdateTracklet(int32_t /*nBlocks*/, int32_t /*nThreads*/, int32_t /*iBlock*/, int32_t /*iThread*/, GPUsharedref() T& s, GPUTPCThreadMemory& GPUrestrict() r, GPUconstantref() GPUTPCTracker& GPUrestrict() tracker, MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference> tParam, int32_t iRow, calink& rowHit, calink* rowHits) // reference_restrict
 {
   // reconstruction of tracklets, tracklets update step
   CA_MAKE_SHARED_REF(GPUTPCRow, row, tracker.Row(iRow), s.mRows[iRow]);
@@ -403,7 +403,7 @@ GPUdic(2, 1) void GPUTPCTrackletConstructor::UpdateTracklet(int32_t /*nBlocks*/,
 GPUdic(2, 1) void GPUTPCTrackletConstructor::DoTracklet(GPUconstantref() GPUTPCTracker& GPUrestrict() tracker, GPUsharedref() GPUTPCTrackletConstructor::GPUSharedMemory& s, GPUTPCThreadMemory& GPUrestrict() r)
 {
   int32_t iRow = 0, iRowEnd = GPUCA_ROW_COUNT;
-  GPUTPCTrackParamSkeleton<MemLayout::value> tParam;
+  MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::value> tParam;
   calink rowHits[GPUCA_ROW_COUNT];
   if (r.mGo) {
     GPUTPCHitId id = tracker.TrackletStartHits()[r.mISH];
@@ -491,7 +491,7 @@ GPUdii() void GPUTPCTrackletConstructor::Thread(int32_t nBlocks, int32_t nThread
 }
 
 template <> // FIXME: GPUgeneric() needed to make the clang spirv output link correctly
-GPUd() int32_t GPUTPCTrackletConstructor::GPUTPCTrackletConstructorExtrapolationTracking<GPUgeneric() GPUTPCExtrapolationTracking::GPUSharedMemory>(GPUconstantref() GPUTPCTracker& GPUrestrict() tracker, GPUsharedref() GPUTPCExtrapolationTracking::GPUSharedMemory& sMem, GPUTPCTrackParamSkeleton<MemLayout::reference_restrict> tParam, int32_t row, int32_t increment, int32_t iTracklet, calink* rowHits)
+GPUd() int32_t GPUTPCTrackletConstructor::GPUTPCTrackletConstructorExtrapolationTracking<GPUgeneric() GPUTPCExtrapolationTracking::GPUSharedMemory>(GPUconstantref() GPUTPCTracker& GPUrestrict() tracker, GPUsharedref() GPUTPCExtrapolationTracking::GPUSharedMemory& sMem, MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference> tParam, int32_t row, int32_t increment, int32_t iTracklet, calink* rowHits) // reference_restrict
 {
   GPUTPCThreadMemory rMem;
   rMem.mISH = iTracklet;
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.h b/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.h
index c4792b9e00ec5..d6c315a104082 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.h
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.h
@@ -72,19 +72,18 @@ class GPUTPCTrackletConstructor : public GPUKernelTemplate
 #endif                // GPUCA_TRACKLET_CONSTRUCTOR_DO_PROFILE
   };
 
-  GPUd() static void InitTracklet(GPUTPCTrackParamSkeleton<MemLayout::reference_restrict> tParam);
+  GPUd() static void InitTracklet(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference> tParam); // reference_restrict
 
   template <class T>
-  GPUd() static void UpdateTracklet(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() T& s, GPUTPCThreadMemory& r, GPUconstantref() GPUTPCTracker& tracker, GPUTPCTrackParamSkeleton<MemLayout::reference_restrict> tParam, int32_t iRow, calink& rowHit, calink* rowHits);
+  GPUd() static void UpdateTracklet(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() T& s, GPUTPCThreadMemory& r, GPUconstantref() GPUTPCTracker& tracker, MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference> tParam, int32_t iRow, calink& rowHit, calink* rowHits); // reference_restrict
 
-  GPUd() static void StoreTracklet(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& s, GPUTPCThreadMemory& r, GPUconstantref() GPUTPCTracker& tracker, GPUTPCTrackParamSkeleton<MemLayout::reference_restrict> tParam, calink* rowHits);
-
-  GPUd() static bool CheckCov(GPUTPCTrackParamSkeleton<MemLayout::reference_restrict> tParam);
+  GPUd() static void StoreTracklet(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& s, GPUTPCThreadMemory& r, GPUconstantref() GPUTPCTracker& tracker, MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference> tParam, calink* rowHits); // reference_restrict
 
+  GPUd() static bool CheckCov(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference> tParam); // reference_restrict
   GPUd() static void DoTracklet(GPUconstantref() GPUTPCTracker& tracker, GPUsharedref() GPUTPCTrackletConstructor::GPUSharedMemory& sMem, GPUTPCThreadMemory& rMem);
 
   template <class T>
-  GPUd() static int32_t GPUTPCTrackletConstructorExtrapolationTracking(GPUconstantref() GPUTPCTracker& tracker, GPUsharedref() T& sMem, GPUTPCTrackParamSkeleton<MemLayout::reference_restrict> tParam, int32_t startrow, int32_t increment, int32_t iTracklet, calink* rowHits);
+  GPUd() static int32_t GPUTPCTrackletConstructorExtrapolationTracking(GPUconstantref() GPUTPCTracker& tracker, GPUsharedref() T& sMem, MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference> tParam, int32_t startrow, int32_t increment, int32_t iTracklet, calink* rowHits); // reference_restrict
 
   typedef GPUconstantref() GPUTPCTracker processorType;
   GPUhdi() constexpr static GPUDataTypes::RecoStep GetRecoStep() { return GPUDataTypes::RecoStep::TPCSectorTracking; }
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTrackletSelector.cxx b/GPU/GPUTracking/SectorTracker/GPUTPCTrackletSelector.cxx
index 01823a641d92d..592998af560e2 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCTrackletSelector.cxx
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCTrackletSelector.cxx
@@ -40,7 +40,7 @@ GPUdii() void GPUTPCTrackletSelector::Thread<0>(int32_t nBlocks, int32_t nThread
   for (int32_t itr = s.mItr0 + iThread; itr < s.mNTracklets; itr += s.mNThreadsTotal) {
     GPUbarrierWarp();
 
-    GPUglobalref() GPUTPCTrackletSkeleton<MemLayout::reference_restrict> tracklet = tracker.Tracklet(itr);
+    GPUglobalref() MemLayout::wrapper<GPUTPCTrackletSkeleton, MemLayout::reference> tracklet = tracker.Tracklet(itr); // reference_restrict
 
     int32_t firstRow = tracklet.FirstRow();
     int32_t lastRow = tracklet.LastRow();
diff --git a/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx b/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx
index c25efb21e2279..5a46e07c8ee2c 100644
--- a/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx
+++ b/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx
@@ -262,7 +262,7 @@ GPUDisplay::vboList GPUDisplay::DrawTracklets(const GPUTPCTracker& tracker)
   }
   size_t startCount = mVertexBufferStart[iSector].size();
   for (uint32_t i = 0; i < *tracker.NTracklets(); i++) {
-    GPUTPCTrackletSkeleton<MemLayout::const_reference> tracklet = tracker.Tracklet(i);
+    MemLayout::wrapper<GPUTPCTrackletSkeleton, MemLayout::const_reference> tracklet = tracker.Tracklet(i);
     size_t startCountInner = mVertexBuffer[iSector].size();
     float4 oldpos;
     for (int32_t j = tracklet.FirstRow(); j <= tracklet.LastRow(); j++) {

From d94784a5c5101c2c74e653318c0b41b3cc32c2a4 Mon Sep 17 00:00:00 2001
From: Oliver Rietmann <oliver.rietmann@cern.ch>
Date: Fri, 6 Feb 2026 14:02:22 +0100
Subject: [PATCH 2/2] Add iterators and specializations for pointer wrapper

---
 .github/workflows/standalone-benchmark.yml    |   1 +
 GPU/Common/GPUCommonAlgorithm.h               |   1 -
 GPU/Common/GPUCommonAlgorithmThrust.h         |   1 -
 GPU/Common/MemLayout.h                        | 150 ++++++++++++++----
 .../GPUTPCCompressionTrackModel.h             |   3 +-
 GPU/GPUTracking/DataTypes/GPUDataTypes.h      |   2 +-
 GPU/GPUTracking/Global/GPUChainTracking.cxx   |  34 ++--
 GPU/GPUTracking/Global/GPUChainTracking.h     |  12 +-
 GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx     |  56 +++----
 GPU/GPUTracking/Merger/GPUTPCGMMerger.h       |   5 +-
 GPU/GPUTracking/Merger/GPUTPCGMMergerDump.cxx |   5 +-
 .../Merger/GPUTPCGMSectorTrack.cxx            |  20 +--
 GPU/GPUTracking/Merger/GPUTPCGMSectorTrack.h  |   4 +-
 .../SectorTracker/GPUTPCBaseTrackParam.h      |  35 ++--
 .../GPUTPCExtrapolationTracking.cxx           |   3 +-
 .../GPUTPCSectorDebugSortKernels.cxx          |   4 +-
 GPU/GPUTracking/SectorTracker/GPUTPCTrack.h   |   5 +-
 .../SectorTracker/GPUTPCTrackLinearisation.h  |   4 +-
 .../SectorTracker/GPUTPCTrackParam.cxx        |  21 ++-
 .../SectorTracker/GPUTPCTrackParam.h          |  25 +--
 .../SectorTracker/GPUTPCTracker.cxx           |  48 +++---
 GPU/GPUTracking/SectorTracker/GPUTPCTracker.h |  27 ++--
 .../SectorTracker/GPUTPCTrackerDump.cxx       |  19 +--
 .../SectorTracker/GPUTPCTracklet.h            |   2 +
 .../GPUTPCTrackletConstructor.cxx             |  12 +-
 .../SectorTracker/GPUTPCTrackletConstructor.h |  10 +-
 .../SectorTracker/GPUTPCTrackletSelector.cxx  |   2 +-
 .../display/render/GPUDisplayDraw.cxx         |   2 +-
 28 files changed, 314 insertions(+), 199 deletions(-)

diff --git a/.github/workflows/standalone-benchmark.yml b/.github/workflows/standalone-benchmark.yml
index 0713ffc6a1ed7..9527e6a5a4403 100644
--- a/.github/workflows/standalone-benchmark.yml
+++ b/.github/workflows/standalone-benchmark.yml
@@ -73,6 +73,7 @@ jobs:
           STANDALONE_DIR: /root/standalone
           BUILD_DIR: /root/standalone/build
           ARTIFACT_FILE: /root/artifact.txt
+          ALIBUILD_O2_FORCE_GPU: 1
 
       - name: Upload Artifact
         uses: actions/upload-artifact@v4
diff --git a/GPU/Common/GPUCommonAlgorithm.h b/GPU/Common/GPUCommonAlgorithm.h
index 655c0291129ff..a1bbdd930b4cb 100644
--- a/GPU/Common/GPUCommonAlgorithm.h
+++ b/GPU/Common/GPUCommonAlgorithm.h
@@ -16,7 +16,6 @@
 #define GPUCOMMONALGORITHM_H
 
 #include "GPUCommonDef.h"
-#include "MemLayout.h"
 
 #if !defined(GPUCA_GPUCODE) // Could also enable custom search on the CPU, but it is not always faster, so we stick to std::sort
 #include <algorithm>
diff --git a/GPU/Common/GPUCommonAlgorithmThrust.h b/GPU/Common/GPUCommonAlgorithmThrust.h
index fa842b485d746..be74e2143ae37 100644
--- a/GPU/Common/GPUCommonAlgorithmThrust.h
+++ b/GPU/Common/GPUCommonAlgorithmThrust.h
@@ -19,7 +19,6 @@
 #pragma GCC diagnostic push
 #pragma GCC diagnostic ignored "-Wshadow"
 #include <thrust/sort.h>
-#include <thrust/iterator/iterator_traits.h>
 #include <thrust/execution_policy.h>
 #include <thrust/device_ptr.h>
 #pragma GCC diagnostic pop
diff --git a/GPU/Common/MemLayout.h b/GPU/Common/MemLayout.h
index 1d4952904cdb1..250ad36435474 100644
--- a/GPU/Common/MemLayout.h
+++ b/GPU/Common/MemLayout.h
@@ -12,29 +12,34 @@ template <class T> using const_reference = const T&;
 template <class T> using pointer = T*;
 template <class T> using const_pointer = const T*;
 
+template <class T> using reference_restrict = T& __restrict__;
+template <class T> using const_reference_restrict = const T& __restrict__;
+template <class T> using pointer_restrict = T* __restrict__;
+template <class T> using const_pointer_restrict = const T* __restrict__;
+
 template <class SF>
 struct RandomAccessAt {
     MemLayout::size_t i;
     template <class... Args>
-    constexpr SF operator()(Args& ...args) const { return {{args[i]...}}; }
+    constexpr SF operator()(Args& ...args) const { return {args[i]...}; }
     template <class... Args>
-    constexpr SF operator()(const Args& ...args) const { return {{args[i]...}}; }
+    constexpr SF operator()(const Args& ...args) const { return {args[i]...}; }
 };
 
 template <class SF>
 struct GetPointer {
     template <class... Args>
-    constexpr SF operator()(Args& ...args) const { return {{&args...}}; }
+    constexpr SF operator()(Args& ...args) const { return {&args...}; }
     template <class... Args>
-    constexpr SF operator()(const Args& ...args) const { return {{&args...}}; }
+    constexpr SF operator()(const Args& ...args) const { return {&args...}; }
 };
 
 template <class SF>
 struct AggregateConstructor {
     template <class... Args>
-    constexpr SF operator()(Args& ...args) const { return {{args...}}; }
+    constexpr SF operator()(Args& ...args) const { return {args...}; }
     template <class... Args>
-    constexpr SF operator()(const Args& ...args) const { return {{args...}}; } 
+    constexpr SF operator()(const Args& ...args) const { return {args...}; }
 };
 
 struct FirstMember {
@@ -47,20 +52,20 @@ struct FirstMember {
 template <class SF>
 struct PreIncrement {
     template <class... Args>
-    constexpr SF operator()(Args& ...args) const { return {{++args...}}; }
+    constexpr SF operator()(Args& ...args) const { return {++args...}; }
 };
 
 template <class SF>
 struct PreDecrement {
     template <class... Args>
-    constexpr SF operator()(Args& ...args) const { return {{--args...}}; }
+    constexpr SF operator()(Args& ...args) const { return {--args...}; }
 };
 
 template <class SF>
 struct Advance {
     ptrdiff_t i;
     template <class... Args>
-    constexpr SF operator()(const Args& ...args) const { return {{(args + i)...}}; }
+    constexpr SF operator()(const Args& ...args) const { return {(args + i)...}; }
 };
 
 struct CopyAssignment {
@@ -78,10 +83,12 @@ struct wrapper : public S<F> {
     constexpr wrapper() = default;
     constexpr wrapper(Base b) : Base{static_cast<Base&&>(b)} {}
     template <template <class> class F_other>
-    constexpr wrapper(wrapper<S, F_other>& other) : Base{other.apply(AggregateConstructor<wrapper>{})} {}
+    constexpr wrapper(S<F_other>& other) : Base{other.apply(AggregateConstructor<Base>{})} {}
+    template <template <class> class F_other>
+    constexpr wrapper(const S<F_other>& other) : Base{other.apply(AggregateConstructor<Base>{})} {}
 
-    constexpr wrapper<S, reference> operator[] (size_t i) { return Base::apply(RandomAccessAt<wrapper<S, reference>>{i}); }
-    constexpr wrapper<S, const_reference> operator[] (size_t i) const { return Base::apply(RandomAccessAt<wrapper<S, const_reference>>{i}); }
+    constexpr wrapper<S, reference> operator[] (size_t i) { return Base::apply(RandomAccessAt<S<reference>>{i}); }
+    constexpr wrapper<S, const_reference> operator[] (size_t i) const { return Base::apply(RandomAccessAt<S<const_reference>>{i}); }
 
     constexpr wrapper<S, reference> operator*() { return operator[](0); }
     constexpr wrapper<S, const_reference> operator*(ptrdiff_t) const { return operator[](0); }
@@ -93,8 +100,8 @@ struct wrapper<S, value> : public S<value> {
 
     constexpr wrapper() = default;
     constexpr wrapper(Base b) : Base{static_cast<Base&&>(b)} {}
-    constexpr wrapper(wrapper<S, reference> other) : Base(other.apply(AggregateConstructor<wrapper>{})) {}
-    constexpr wrapper(wrapper<S, const_reference> other) : Base(other.apply(AggregateConstructor<wrapper>{})) {}
+    constexpr wrapper(const S<reference>& other) : Base(other.apply(AggregateConstructor<Base>{})) {}
+    constexpr wrapper(const S<const_reference>& other) : Base(other.apply(AggregateConstructor<Base>{})) {}
 
     //constexpr S<MemLayout::pointer> operator& () { return apply(GetPointer<S<MemLayout::pointer>>{}); }
     //constexpr S<MemLayout::const_pointer> operator& () const { return apply(GetPointer<S<MemLayout::const_pointer>>{}); }
@@ -106,10 +113,15 @@ struct wrapper<S, reference> : public S<reference> {
 
     constexpr wrapper() = delete;
     constexpr wrapper(Base b) : Base{static_cast<Base&&>(b)} {}
-    constexpr wrapper(wrapper<S, value>& other) : Base(other.apply(AggregateConstructor<wrapper>{})) {}
+    constexpr wrapper(S<value>& other) : Base(other.apply(AggregateConstructor<Base>{})) {}
+    constexpr wrapper(S<reference_restrict> other) : Base(other.apply(AggregateConstructor<Base>{})) {}
     
     constexpr wrapper(const wrapper& other) = default;
 
+    constexpr wrapper& operator=(const wrapper<S, value>& other) {
+        Base::apply(other, CopyAssignment{});
+        return *this;
+    }
     constexpr wrapper& operator=(const wrapper& other) {
         Base::apply(other, CopyAssignment{});
         return *this;
@@ -118,17 +130,63 @@ struct wrapper<S, reference> : public S<reference> {
         Base::apply(other, CopyAssignment{});
         return *this;
     }
+    constexpr wrapper& operator=(const wrapper<S, reference_restrict>& other) {
+        Base::apply(other, CopyAssignment{});
+        return *this;
+    }
+    constexpr wrapper& operator=(const wrapper<S, const_reference_restrict>& other) {
+        Base::apply(other, CopyAssignment{});
+        return *this;
+    }
+
+    constexpr wrapper(wrapper&& other) = default;
+    
+    constexpr wrapper& operator=(wrapper&& other) { return operator=(other); }
+
+    constexpr wrapper<S, pointer> operator&() { return Base::apply(GetPointer<S<pointer>>{}); }
+    //constexpr wrapper<S, const_pointer> operator&() const { return Base::apply(GetPointer<S<const_pointer>>{}); }
+    constexpr pointer<wrapper<S, reference>> operator->() { return this; }
+};
+
+template <template <template <class> class> class S>
+struct wrapper<S, reference_restrict> : public S<reference_restrict> {
+    using Base = S<reference_restrict>;
+
+    constexpr wrapper() = delete;
+    constexpr wrapper(Base b) : Base{static_cast<Base&&>(b)} {}
+    constexpr wrapper(S<value>& other) : Base(other.apply(AggregateConstructor<Base>{})) {}
+    constexpr wrapper(S<reference> other) : Base(other.apply(AggregateConstructor<Base>{})) {}
+    
+    constexpr wrapper(const wrapper& other) = default;
+
     constexpr wrapper& operator=(const wrapper<S, value>& other) {
         Base::apply(other, CopyAssignment{});
         return *this;
     }
+    constexpr wrapper& operator=(const wrapper& other) {
+        Base::apply(other, CopyAssignment{});
+        return *this;
+    }
+    constexpr wrapper& operator=(const wrapper<S, reference>& other) {
+        Base::apply(other, CopyAssignment{});
+        return *this;
+    }
+    constexpr wrapper& operator=(const wrapper<S, const_reference>& other) {
+        Base::apply(other, CopyAssignment{});
+        return *this;
+    }
+    constexpr wrapper& operator=(const wrapper<S, const_reference_restrict>& other) {
+        Base::apply(other, CopyAssignment{});
+        return *this;
+    }
 
     constexpr wrapper(wrapper&& other) = default;
     
     constexpr wrapper& operator=(wrapper&& other) { return operator=(other); }
 
-    constexpr wrapper<S, pointer> operator& () { return Base::apply(GetPointer<wrapper<S, pointer>>{}); }
-    //constexpr wrapper<S, const_pointer> operator& () const { return Base::apply(GetPointer<wrapper<S, const_pointer>>{}); }
+    constexpr wrapper<S, pointer> operator&() { return Base::apply(GetPointer<S<pointer>>{}); }
+    //constexpr wrapper<S, const_pointer> operator&() const { return Base::apply(GetPointer<S<const_pointer>>{}); }
+    constexpr pointer<wrapper<S, reference>> operator->() { return this; }
 };
 
 template <template <template <class> class> class S>
@@ -137,10 +195,28 @@ struct wrapper<S, const_reference> : public S<const_reference> {
 
     constexpr wrapper() = delete;
     constexpr wrapper(Base b) : Base{static_cast<Base&&>(b)} {}
-    constexpr wrapper(wrapper<S, value> other) : Base(other.apply(AggregateConstructor<wrapper>{})) {}
-    constexpr wrapper(wrapper<S, reference> other) : Base(other.apply(AggregateConstructor<wrapper>{})) {}
+    constexpr wrapper(const S<value>& other) : Base(other.apply(AggregateConstructor<Base>{})) {}
+    constexpr wrapper(const S<reference>& other) : Base(other.apply(AggregateConstructor<Base>{})) {}
+    constexpr wrapper(const S<reference_restrict>& other) : Base(other.apply(AggregateConstructor<Base>{})) {}
+    constexpr wrapper(const S<const_reference_restrict>& other) : Base(other.apply(AggregateConstructor<Base>{})) {}
 
-    constexpr wrapper<S, const_pointer> operator&() const { return Base::apply(GetPointer<wrapper<S, const_pointer>>{}); }
+    constexpr wrapper<S, const_pointer> operator&() const { return Base::apply(GetPointer<S<const_pointer>>{}); }
+    constexpr const_pointer<wrapper<S, const_reference>> operator->() const { return this; }
+};
+
+template <template <template <class> class> class S>
+struct wrapper<S, const_reference_restrict> : public S<const_reference_restrict> {
+    using Base = S<const_reference_restrict>;
+
+    constexpr wrapper() = delete;
+    constexpr wrapper(Base b) : Base{static_cast<Base&&>(b)} {}
+    constexpr wrapper(const S<value>& other) : Base(other.apply(AggregateConstructor<Base>{})) {}
+    constexpr wrapper(const S<reference>& other) : Base(other.apply(AggregateConstructor<Base>{})) {}
+    constexpr wrapper(const S<reference_restrict>& other) : Base(other.apply(AggregateConstructor<Base>{})) {}
+    constexpr wrapper(const S<const_reference>& other) : Base(other.apply(AggregateConstructor<Base>{})) {}
+
+    constexpr wrapper<S, const_pointer> operator&() const { return Base::apply(GetPointer<S<const_pointer>>{}); }
+    constexpr const_pointer<wrapper<S, const_reference>> operator->() const { return this; }
 };
 
 template <template <template <class> class> class S>
@@ -150,23 +226,25 @@ struct wrapper<S, pointer> : public S<pointer> {
     constexpr wrapper() = default;
     constexpr wrapper(Base b) : Base{static_cast<Base&&>(b)} {}
 
-    constexpr wrapper<S, reference> operator[] (size_t i) { return Base::apply(RandomAccessAt<wrapper<S, reference>>{i}); }
-    constexpr wrapper<S, const_reference> operator[] (size_t i) const { return Base::apply(RandomAccessAt<wrapper<S, const_reference>>{i}); }
+    constexpr wrapper<S, reference> operator[] (size_t i) { return Base::apply(RandomAccessAt<S<reference>>{i}); }
+    constexpr const wrapper<S, const_reference> operator[] (size_t i) const { return Base::apply(RandomAccessAt<S<const_reference>>{i}); }
 
     constexpr wrapper<S, reference> operator*() { return operator[](0); }
     constexpr wrapper<S, const_reference> operator*() const { return operator[](0); }
+    constexpr wrapper<S, reference> operator->() { return operator[](0); }
+    constexpr wrapper<S, const_reference> operator->() const { return operator[](0); }
 
     constexpr bool operator==(const wrapper& other) const { return Base::apply(FirstMember{}) == other.apply(FirstMember{}); }
     constexpr bool operator!=(const wrapper& other) const { return !this->operator==(other); }
     constexpr bool operator<(const wrapper& other) const { return Base::apply(FirstMember{}) < other.apply(FirstMember{}); }
 
-    constexpr wrapper operator+(ptrdiff_t i) const { return Base::apply(Advance<wrapper>{i}); }
+    constexpr wrapper operator+(ptrdiff_t i) const { return Base::apply(Advance<Base>{i}); }
     constexpr wrapper operator-(ptrdiff_t i) const { return operator+(-i); }
     constexpr ptrdiff_t operator-(const wrapper& other) const { return Base::apply(FirstMember{}) - other.apply(FirstMember{}); }
 
-    constexpr wrapper& operator++() { Base::apply(PreIncrement<wrapper>{}); return *this; }
+    constexpr wrapper& operator++() { Base::apply(PreIncrement<Base>{}); return *this; }
     constexpr wrapper& operator+=(ptrdiff_t i) { return *this = *this + i; }
-    constexpr wrapper& operator--() { Base::apply(PreDecrement<wrapper>{}); return *this; }
+    constexpr wrapper& operator--() { Base::apply(PreDecrement<Base>{}); return *this; }
     constexpr wrapper& operator-=(ptrdiff_t i) { return *this = *this - i; }
 };
 
@@ -176,10 +254,24 @@ struct wrapper<S, const_pointer> : public S<const_pointer> {
 
     constexpr wrapper() = default;
     constexpr wrapper(Base b) : Base{static_cast<Base&&>(b)} {}
-    constexpr wrapper(wrapper<S, pointer> other) : Base(other.apply(AggregateConstructor<wrapper>{})) {}
+    constexpr wrapper(const S<pointer>& other) : Base(other.apply(AggregateConstructor<Base>{})) {}
 
-    constexpr wrapper<S, const_reference> operator[] (size_t i) const { return Base::apply(RandomAccessAt<wrapper<S, const_reference>>{i}); }
+    constexpr wrapper<S, const_reference> operator[] (size_t i) const { return Base::apply(RandomAccessAt<S<const_reference>>{i}); }
     constexpr wrapper<S, const_reference> operator*() const { return operator[](0); }
+    constexpr wrapper<S, const_reference> operator->() const { return operator[](0); }
+
+    constexpr bool operator==(const wrapper& other) const { return Base::apply(FirstMember{}) == other.apply(FirstMember{}); }
+    constexpr bool operator!=(const wrapper& other) const { return !this->operator==(other); }
+    constexpr bool operator<(const wrapper& other) const { return Base::apply(FirstMember{}) < other.apply(FirstMember{}); }
+
+    constexpr wrapper operator+(ptrdiff_t i) const { return Base::apply(Advance<Base>{i}); }
+    constexpr wrapper operator-(ptrdiff_t i) const { return operator+(-i); }
+    constexpr ptrdiff_t operator-(const wrapper& other) const { return Base::apply(FirstMember{}) - other.apply(FirstMember{}); }
+
+    constexpr wrapper& operator++() { Base::apply(PreIncrement<Base>{}); return *this; }
+    constexpr wrapper& operator+=(ptrdiff_t i) { return *this = *this + i; }
+    constexpr wrapper& operator--() { Base::apply(PreDecrement<Base>{}); return *this; }
+    constexpr wrapper& operator-=(ptrdiff_t i) { return *this = *this - i; }
 };
 
 enum Flag { soa, aos };
@@ -204,6 +296,10 @@ struct interface<S, F, Flag::soa> { using type = wrapper<S, F>; };
 #define MEMLAYOUT_EXPAND(m) f(m, other.m)
 
 #define MEMLAYOUT_APPLY_BINARY(STRUCT_NAME, ...)\
+    template <template <class> class F_other, class Function>\
+    constexpr STRUCT_NAME apply(STRUCT_NAME<F_other>& other, Function&& f) { return {__VA_ARGS__}; }\
+    template <template <class> class F_other, class Function>\
+    constexpr STRUCT_NAME apply(STRUCT_NAME<F_other>& other, Function&& f) const { return {__VA_ARGS__}; }\
     template <template <class> class F_other, class Function>\
     constexpr STRUCT_NAME apply(const STRUCT_NAME<F_other>& other, Function&& f) { return {__VA_ARGS__}; }\
     template <template <class> class F_other, class Function>\
diff --git a/GPU/GPUTracking/DataCompression/GPUTPCCompressionTrackModel.h b/GPU/GPUTracking/DataCompression/GPUTPCCompressionTrackModel.h
index b3b4da27e625b..ea43459e7315f 100644
--- a/GPU/GPUTracking/DataCompression/GPUTPCCompressionTrackModel.h
+++ b/GPU/GPUTracking/DataCompression/GPUTPCCompressionTrackModel.h
@@ -27,6 +27,7 @@
 
 #elif defined(GPUCA_COMPRESSION_TRACK_MODEL_SECTORTRACKER)
 #include "GPUTPCTrackParam.h"
+#include "MemLayout.h"
 
 #else // Default internal track model for compression
 #endif
@@ -121,7 +122,7 @@ class GPUTPCCompressionTrackModel
   const GPUParam* mParam;
 
 #elif defined(GPUCA_COMPRESSION_TRACK_MODEL_SECTORTRACKER)
-  GPUTPCTrackParam mTrk;
+  GPUTPCTrackParamSkeleton<MemLayout::value> mTrk;
   float mAlpha;
   const GPUParam* mParam;
 
diff --git a/GPU/GPUTracking/DataTypes/GPUDataTypes.h b/GPU/GPUTracking/DataTypes/GPUDataTypes.h
index 4e08fbe12e100..32f87719ab2d8 100644
--- a/GPU/GPUTracking/DataTypes/GPUDataTypes.h
+++ b/GPU/GPUTracking/DataTypes/GPUDataTypes.h
@@ -225,7 +225,7 @@ struct GPUTrackingInOutPointers {
   const AliHLTTPCRawCluster* rawClusters[NSECTORS] = {nullptr};
   uint32_t nRawClusters[NSECTORS] = {0};
   const o2::tpc::ClusterNativeAccess* clustersNative = nullptr;
-  MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_pointer> sectorTracks[NSECTORS];// = {{nullptr, nullptr, nullptr, {nullptr, nullptr, nullptr, nullptr}}};
+  MemLayout::interface<GPUTPCTrackSkeleton, MemLayout::const_pointer, GPUTPCTrackLayout>::type sectorTracks[NSECTORS];
   uint32_t nSectorTracks[NSECTORS] = {0};
   const GPUTPCHitId* sectorClusters[NSECTORS] = {nullptr};
   uint32_t nSectorClusters[NSECTORS] = {0};
diff --git a/GPU/GPUTracking/Global/GPUChainTracking.cxx b/GPU/GPUTracking/Global/GPUChainTracking.cxx
index b9749ff5ea609..ea96587a5ff41 100644
--- a/GPU/GPUTracking/Global/GPUChainTracking.cxx
+++ b/GPU/GPUTracking/Global/GPUChainTracking.cxx
@@ -549,23 +549,33 @@ void GPUChainTracking::ClearIOPointers()
   new (&mIOMem) InOutMemory;
 }
 
+namespace {
+
+template <class Function>
+struct ApplyRecursive {
+    Function f;
+
+    template <class T>
+    const T * operator()(const T * & aosIOPtr, GPUChainTracking::unique_ptr_array<T>& aosIOMem) const { return f(aosIOPtr, aosIOMem); }
+
+    template <template <template <class> class> class S>
+    S<MemLayout::const_pointer> operator()(S<MemLayout::const_pointer>& soaIOPtr, S<GPUChainTracking::unique_ptr_array>& soaIOMem) const {
+        return soaIOPtr.apply(soaIOMem, ApplyRecursive{f});
+    }
+};
+
+}
+
 void GPUChainTracking::AllocateIOMemory()
 {
   for (uint32_t i = 0; i < NSECTORS; i++) {
     AllocateIOMemoryHelper(mIOPtrs.nClusterData[i], mIOPtrs.clusterData[i], mIOMem.clusterData[i]);
     AllocateIOMemoryHelper(mIOPtrs.nRawClusters[i], mIOPtrs.rawClusters[i], mIOMem.rawClusters[i]);
-    
-    //AllocateIOMemoryHelper(mIOPtrs.nSectorTracks[i], mIOPtrs.sectorTracks[i], mIOMem.sectorTracks[i]);
-
-    AllocateIOMemoryHelper(mIOPtrs.nSectorTracks[i], mIOPtrs.sectorTracks[i].mFirstHitID, mIOMem.sectorTracks[i].mFirstHitID);
-    AllocateIOMemoryHelper(mIOPtrs.nSectorTracks[i], mIOPtrs.sectorTracks[i].mNHits, mIOMem.sectorTracks[i].mNHits);
-    AllocateIOMemoryHelper(mIOPtrs.nSectorTracks[i], mIOPtrs.sectorTracks[i].mLocalTrackId, mIOMem.sectorTracks[i].mLocalTrackId);
-
-    AllocateIOMemoryHelper(mIOPtrs.nSectorTracks[i], mIOPtrs.sectorTracks[i].mParam.mX, mIOMem.sectorTracks[i].mParam.mX);
-    AllocateIOMemoryHelper(mIOPtrs.nSectorTracks[i], mIOPtrs.sectorTracks[i].mParam.mC, mIOMem.sectorTracks[i].mParam.mC);
-    AllocateIOMemoryHelper(mIOPtrs.nSectorTracks[i], mIOPtrs.sectorTracks[i].mParam.mZOffset, mIOMem.sectorTracks[i].mParam.mZOffset);
-    AllocateIOMemoryHelper(mIOPtrs.nSectorTracks[i], mIOPtrs.sectorTracks[i].mParam.mP, mIOMem.sectorTracks[i].mParam.mP);
-
+    auto sectorTrackAllocator = [this, nSectorTrack = this->mIOPtrs.nSectorTracks[i]](auto& IOPtrsTrack, auto& mIOMemTrack) {
+        AllocateIOMemoryHelper(nSectorTrack, IOPtrsTrack, mIOMemTrack);
+        return IOPtrsTrack;
+    };
+    ApplyRecursive{sectorTrackAllocator}(mIOPtrs.sectorTracks[i], mIOMem.sectorTracks[i]);
     AllocateIOMemoryHelper(mIOPtrs.nSectorClusters[i], mIOPtrs.sectorClusters[i], mIOMem.sectorClusters[i]);
   }
   mIOMem.clusterNativeAccess.reset(new ClusterNativeAccess);
diff --git a/GPU/GPUTracking/Global/GPUChainTracking.h b/GPU/GPUTracking/Global/GPUChainTracking.h
index 251311aebed4b..23b2c1266d6f2 100644
--- a/GPU/GPUTracking/Global/GPUChainTracking.h
+++ b/GPU/GPUTracking/Global/GPUChainTracking.h
@@ -17,7 +17,9 @@
 
 #include "GPUChain.h"
 #include "GPUDataTypes.h"
+#include "GPUTPCTrack.h"
 #include "MemLayout.h"
+
 #include <atomic>
 #include <mutex>
 #include <functional>
@@ -91,6 +93,9 @@ class GPUChainTracking : public GPUChain
   // Structures for input and output data
   GPUTrackingInOutPointers& mIOPtrs;
 
+  template <class T>
+  using unique_ptr_array = std::unique_ptr<T[]>;
+
   struct InOutMemory {
     InOutMemory();
     ~InOutMemory();
@@ -108,12 +113,7 @@ class GPUChainTracking : public GPUChain
     std::unique_ptr<AliHLTTPCRawCluster[]> rawClusters[NSECTORS];
     std::unique_ptr<o2::tpc::ClusterNative[]> clustersNative;
     std::unique_ptr<o2::tpc::ClusterNativeAccess> clusterNativeAccess;
-
-    template <class T>
-    using unique_ptr_array = std::unique_ptr<T[]>;
-
-    MemLayout::wrapper<GPUTPCTrackSkeleton, unique_ptr_array> sectorTracks[NSECTORS];
-
+    MemLayout::interface<GPUTPCTrackSkeleton, unique_ptr_array, GPUTPCTrackLayout>::type sectorTracks[NSECTORS];
     std::unique_ptr<GPUTPCHitId[]> sectorClusters[NSECTORS];
     std::unique_ptr<AliHLTTPCClusterMCLabel[]> mcLabelsTPC;
     std::unique_ptr<GPUTPCMCInfo[]> mcInfosTPC;
diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx
index d2fd297f552ff..83dc127e1f82e 100644
--- a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx
+++ b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx
@@ -237,7 +237,7 @@ int64_t GPUTPCGMMerger::GetTrackLabelA(const S& trk) const
   int32_t nClusters = 0;
   if constexpr (std::is_same_v<S, GPUTPCGMBorderTrack&>) {
     sectorTrack = &mSectorTrackInfos[trk.TrackID()];
-    nClusters = (*sectorTrack->OrigTrack()).NHits();
+    nClusters = sectorTrack->OrigTrack()->NHits();
   } else {
     nClusters = trk.NClusters();
   }
@@ -246,7 +246,7 @@ int64_t GPUTPCGMMerger::GetTrackLabelA(const S& trk) const
     int32_t id;
     if constexpr (std::is_same_v<S, GPUTPCGMBorderTrack&>) {
       const GPUTPCTracker& tracker = GetConstantMem()->tpcTrackers[sectorTrack->Sector()];
-      const GPUTPCHitId& ic = tracker.TrackHits()[(*sectorTrack->OrigTrack()).FirstHitID() + i];
+      const GPUTPCHitId& ic = tracker.TrackHits()[sectorTrack->OrigTrack()->FirstHitID() + i];
       id = tracker.Data().ClusterDataIndex(tracker.Data().Row(ic.RowIndex()), ic.HitIndex()) + GetConstantMem()->ioPtrs.clustersNative->clusterOffset[sectorTrack->Sector()][0];
     } else {
       id = mClusters[trk.FirstClusterRef() + i].num;
@@ -488,7 +488,7 @@ GPUd() void GPUTPCGMMerger::ClearTrackLinks(int32_t nBlocks, int32_t nThreads, i
   }
 }
 
-GPUd() int32_t GPUTPCGMMerger::RefitSectorTrack(GPUTPCGMSectorTrack& sectorTrack, MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_pointer> inTrack, float alpha, int32_t sector)
+GPUd() int32_t GPUTPCGMMerger::RefitSectorTrack(GPUTPCGMSectorTrack& sectorTrack, MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_reference> inTrack, float alpha, int32_t sector)
 {
   GPUTPCGMPropagator prop;
   prop.SetMaterialTPC();
@@ -497,15 +497,14 @@ GPUd() int32_t GPUTPCGMMerger::RefitSectorTrack(GPUTPCGMSectorTrack& sectorTrack
   prop.SetFitInProjections(false);
   prop.SetPolynomialField(&Param().polynomialField);
   GPUTPCGMTrackParam trk;
-  MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_reference> inTrack_cref = *inTrack;
-  trk.X() = inTrack_cref.Param().GetX();
-  trk.Y() = inTrack_cref.Param().GetY();
-  trk.Z() = inTrack_cref.Param().GetZ();
-  trk.SinPhi() = inTrack_cref.Param().GetSinPhi();
-  trk.DzDs() = inTrack_cref.Param().GetDzDs();
-  trk.QPt() = inTrack_cref.Param().GetQPt();
-  trk.TZOffset() = Param().par.earlyTpcTransform ? inTrack_cref.Param().GetZOffset() : GetConstantMem()->calibObjects.fastTransformHelper->getCorrMap()->convZOffsetToVertexTime(sector, inTrack_cref.Param().GetZOffset(), Param().continuousMaxTimeBin);
-  trk.ShiftZ(this, sector, sectorTrack.ClusterZT0(), sectorTrack.ClusterZTN(), inTrack_cref.Param().GetX(), inTrack_cref.Param().GetX()); // We do not store the inner / outer cluster X, so we just use the track X instead
+  trk.X() = inTrack.Param().GetX();
+  trk.Y() = inTrack.Param().GetY();
+  trk.Z() = inTrack.Param().GetZ();
+  trk.SinPhi() = inTrack.Param().GetSinPhi();
+  trk.DzDs() = inTrack.Param().GetDzDs();
+  trk.QPt() = inTrack.Param().GetQPt();
+  trk.TZOffset() = Param().par.earlyTpcTransform ? inTrack.Param().GetZOffset() : GetConstantMem()->calibObjects.fastTransformHelper->getCorrMap()->convZOffsetToVertexTime(sector, inTrack.Param().GetZOffset(), Param().continuousMaxTimeBin);
+  trk.ShiftZ(this, sector, sectorTrack.ClusterZT0(), sectorTrack.ClusterZTN(), inTrack.Param().GetX(), inTrack.Param().GetX()); // We do not store the inner / outer cluster X, so we just use the track X instead
   sectorTrack.SetX2(0.f);
   for (int32_t way = 0; way < 2; way++) {
     if (way) {
@@ -514,14 +513,14 @@ GPUd() int32_t GPUTPCGMMerger::RefitSectorTrack(GPUTPCGMSectorTrack& sectorTrack
     }
     trk.ResetCovariance();
     prop.SetTrack(&trk, alpha);
-    int32_t start = way ? inTrack_cref.NHits() - 1 : 0;
-    int32_t end = way ? 0 : (inTrack_cref.NHits() - 1);
+    int32_t start = way ? inTrack.NHits() - 1 : 0;
+    int32_t end = way ? 0 : (inTrack.NHits() - 1);
     int32_t incr = way ? -1 : 1;
     for (int32_t i = start; i != end; i += incr) {
       float x, y, z;
       int32_t row, flags;
       const GPUTPCTracker& tracker = GetConstantMem()->tpcTrackers[sector];
-      const GPUTPCHitId& ic = tracker.TrackHits()[inTrack_cref.FirstHitID() + i];
+      const GPUTPCHitId& ic = tracker.TrackHits()[inTrack.FirstHitID() + i];
       int32_t clusterIndex = tracker.Data().ClusterDataIndex(tracker.Data().Row(ic.RowIndex()), ic.HitIndex());
       row = ic.RowIndex();
       const ClusterNative& cl = GetConstantMem()->ioPtrs.clustersNative->clustersLinear[GetConstantMem()->ioPtrs.clustersNative->clusterOffset[sector][0] + clusterIndex];
@@ -551,12 +550,11 @@ GPUd() int32_t GPUTPCGMMerger::RefitSectorTrack(GPUTPCGMSectorTrack& sectorTrack
   return 0;
 }
 
-GPUd() void GPUTPCGMMerger::SetTrackClusterZT(GPUTPCGMSectorTrack& track, int32_t iSector, MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_pointer> sectorTr)
+GPUd() void GPUTPCGMMerger::SetTrackClusterZT(GPUTPCGMSectorTrack& track, int32_t iSector, MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_reference> sectorTr)
 {
-  MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_reference> sectorTr_cref = *sectorTr;
   const GPUTPCTracker& trk = GetConstantMem()->tpcTrackers[iSector];
-  const GPUTPCHitId& ic1 = trk.TrackHits()[sectorTr_cref.FirstHitID()];
-  const GPUTPCHitId& ic2 = trk.TrackHits()[sectorTr_cref.FirstHitID() + sectorTr_cref.NHits() - 1];
+  const GPUTPCHitId& ic1 = trk.TrackHits()[sectorTr.FirstHitID()];
+  const GPUTPCHitId& ic2 = trk.TrackHits()[sectorTr.FirstHitID() + sectorTr.NHits() - 1];
   int32_t clusterIndex1 = trk.Data().ClusterDataIndex(trk.Data().Row(ic1.RowIndex()), ic1.HitIndex());
   int32_t clusterIndex2 = trk.Data().ClusterDataIndex(trk.Data().Row(ic2.RowIndex()), ic2.HitIndex());
   if (Param().par.earlyTpcTransform) {
@@ -576,13 +574,11 @@ GPUd() void GPUTPCGMMerger::UnpackSectorGlobal(int32_t nBlocks, int32_t nThreads
 {
   const GPUTPCTracker& trk = GetConstantMem()->tpcTrackers[iSector];
   float alpha = Param().Alpha(iSector);
-  // GPUTPCTrackSkeleton<MemLayout::const_pointer> sectorTr = mMemory->firstExtrapolatedTracks[iSector];
-  MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_pointer> sectorTr = mMemory->firstExtrapolatedTracks[iSector];
   uint32_t nLocalTracks = trk.CommonMemory()->nLocalTracks;
   uint32_t nTracks = *trk.NTracks();
   for (uint32_t itr = nLocalTracks + iBlock * nThreads + iThread; itr < nTracks; itr += nBlocks * nThreads) {
-    sectorTr = &trk.Tracks()[itr];
-    int32_t localId = mTrackIDs[((*sectorTr).LocalTrackId() >> 24) * mNMaxSingleSectorTracks + ((*sectorTr).LocalTrackId() & 0xFFFFFF)];
+    MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_reference> sectorTr = trk.Tracks()[itr];
+    int32_t localId = mTrackIDs[(sectorTr.LocalTrackId() >> 24) * mNMaxSingleSectorTracks + (sectorTr.LocalTrackId() & 0xFFFFFF)];
     if (localId == -1) {
       continue;
     }
@@ -616,7 +612,7 @@ GPUd() void GPUTPCGMMerger::RefitSectorTracks(int32_t nBlocks, int32_t nThreads,
   float alpha = Param().Alpha(iSector);
 
   for (uint32_t itr = iBlock * nThreads + iThread; itr < nLocalTracks; itr += nBlocks * nThreads) {
-    MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_pointer> sectorTr = &trk.Tracks()[itr];
+    MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_reference> sectorTr = trk.Tracks()[itr];
     GPUTPCGMSectorTrack track;
     SetTrackClusterZT(track, iSector, sectorTr);
     if (Param().rec.tpc.mergerCovSource == 0) {
@@ -644,7 +640,7 @@ GPUd() void GPUTPCGMMerger::RefitSectorTracks(int32_t nBlocks, int32_t nThreads,
     track.SetExtrapolatedTrackId(0, -1);
     track.SetExtrapolatedTrackId(1, -1);
     uint32_t myTrack = CAMath::AtomicAdd(&mMemory->nUnpackedTracks, 1u);
-    mTrackIDs[iSector * mNMaxSingleSectorTracks + (*sectorTr).LocalTrackId()] = myTrack;
+    mTrackIDs[iSector * mNMaxSingleSectorTracks + sectorTr.LocalTrackId()] = myTrack;
     mSectorTrackInfos[myTrack] = track;
   }
 }
@@ -703,7 +699,7 @@ GPUd() void GPUTPCGMMerger::MergeSectorsPrepareStep2(int32_t nBlocks, int32_t nT
       const GPUTPCGMSectorTrack* trackMin = track;
       while (track->NextSegmentNeighbour() >= 0 && track->Sector() == mSectorTrackInfos[track->NextSegmentNeighbour()].Sector()) {
         track = &mSectorTrackInfos[track->NextSegmentNeighbour()];
-        if ((*track->OrigTrack()).Param().X() < (*trackMin->OrigTrack()).Param().X()) {
+        if (track->OrigTrack()->Param().X() < trackMin->OrigTrack()->Param().X()) {
           trackMin = track;
         }
       }
@@ -712,7 +708,7 @@ GPUd() void GPUTPCGMMerger::MergeSectorsPrepareStep2(int32_t nBlocks, int32_t nT
       if (Param().rec.tpc.mergerCovSource == 2 && trackTmp.X2() != 0.f) {
         trackTmp.UseParam2();
       } else {
-        trackTmp.Set(this, trackMin->OrigTrack(), trackMin->Alpha(), trackMin->Sector());
+        trackTmp.Set(this, *trackMin->OrigTrack(), trackMin->Alpha(), trackMin->Sector());
       }
     } else {
       if (CAMath::Abs(track->QPt()) * Param().qptB5Scaler < Param().rec.tpc.mergerLooperSecondHorizontalQPtB5Limit) {
@@ -1605,7 +1601,7 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread
       trackCluster* c2 = trackClusters + nHits + nTrackHits - 1;
       for (int32_t i = 0; i < nTrackHits; i++, c2--) {
         const GPUTPCTracker& trk = GetConstantMem()->tpcTrackers[t->Sector()];
-        const GPUTPCHitId& ic = trk.TrackHits()[(*t->OrigTrack()).FirstHitID() + i];
+        const GPUTPCHitId& ic = trk.TrackHits()[t->OrigTrack()->FirstHitID() + i];
         uint32_t id = trk.Data().ClusterDataIndex(trk.Data().Row(ic.RowIndex()), ic.HitIndex()) + GetConstantMem()->ioPtrs.clustersNative->clusterOffset[t->Sector()][0];
         *c2 = trackCluster{id, (uint8_t)ic.RowIndex(), t->Sector(), t->Leg()};
       }
@@ -1659,9 +1655,9 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread
           if (trackParts[i]->Leg() != baseLeg) {
             break;
           }
-          if ((*trackParts[i]->OrigTrack()).NHits() > length) {
+          if (trackParts[i]->OrigTrack()->NHits() > length) {
             iLongest = i;
-            length = (*trackParts[i]->OrigTrack()).NHits();
+            length = trackParts[i]->OrigTrack()->NHits();
           }
         }
         bool outwards;
diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMerger.h b/GPU/GPUTracking/Merger/GPUTPCGMMerger.h
index 1adc6c377a0ba..0bad9628d1ab6 100644
--- a/GPU/GPUTracking/Merger/GPUTPCGMMerger.h
+++ b/GPU/GPUTracking/Merger/GPUTPCGMMerger.h
@@ -74,7 +74,6 @@ class GPUTPCGMMerger : public GPUProcessor
     GPUAtomic(uint32_t) nMergedTrackClusters;
     GPUAtomic(uint32_t) nO2Tracks;
     GPUAtomic(uint32_t) nO2ClusRefs;
-    MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_pointer> firstExtrapolatedTracks[NSECTORS];
     GPUAtomic(uint32_t) tmpCounter[2 * NSECTORS];
     GPUAtomic(uint32_t) nLooperMatchCandidates;
   };
@@ -154,8 +153,8 @@ class GPUTPCGMMerger : public GPUProcessor
   GPUd() uint16_t MemoryResOutputO2MC() const { return mMemoryResOutputO2MC; }
   GPUd() uint16_t MemoryResOutputO2Scratch() const { return mMemoryResOutputO2Scratch; }
 
-  GPUd() int32_t RefitSectorTrack(GPUTPCGMSectorTrack& sectorTrack, MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_pointer> inTrack, float alpha, int32_t sector);
-  GPUd() void SetTrackClusterZT(GPUTPCGMSectorTrack& track, int32_t iSector, MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_pointer> sectorTr);
+  GPUd() int32_t RefitSectorTrack(GPUTPCGMSectorTrack& sectorTrack, MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_reference> inTrack, float alpha, int32_t sector);
+  GPUd() void SetTrackClusterZT(GPUTPCGMSectorTrack& track, int32_t iSector, MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_reference> sectorTr);
 
   int32_t CheckSectors();
   GPUd() void RefitSectorTracks(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, int32_t iSector);
diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMergerDump.cxx b/GPU/GPUTracking/Merger/GPUTPCGMMergerDump.cxx
index 9c924e74ec519..9dbca015f2204 100644
--- a/GPU/GPUTracking/Merger/GPUTPCGMMergerDump.cxx
+++ b/GPU/GPUTracking/Merger/GPUTPCGMMergerDump.cxx
@@ -37,6 +37,7 @@
 #include "GPUTrackingRefit.h"
 #include "CorrectionMapsHelper.h"
 #include "GPUConstantMem.h"
+#include "MemLayout.h"
 
 using namespace o2::gpu;
 using namespace gputpcgmmergertypes;
@@ -264,7 +265,7 @@ inline void GPUTPCGMMerger::MergedTrackStreamerInternal(const GPUTPCGMBorderTrac
   }
   const GPUTPCTracker& tracker1 = GetConstantMem()->tpcTrackers[sector1];
   const GPUTPCGMSectorTrack& sectorTrack1 = mSectorTrackInfos[b1.TrackID()];
-  const GPUTPCTrack& inTrack1 = *sectorTrack1.OrigTrack();
+  MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_reference> inTrack1 = *sectorTrack1.OrigTrack();
   for (int32_t i = 0; i < inTrack1.NHits(); i++) {
     const GPUTPCHitId& ic1 = tracker1.TrackHits()[inTrack1.FirstHitID() + i];
     int32_t clusterIndex = tracker1.Data().ClusterDataIndex(tracker1.Data().Row(ic1.RowIndex()), ic1.HitIndex());
@@ -272,7 +273,7 @@ inline void GPUTPCGMMerger::MergedTrackStreamerInternal(const GPUTPCGMBorderTrac
   }
   const GPUTPCTracker& tracker2 = GetConstantMem()->tpcTrackers[sector2];
   const GPUTPCGMSectorTrack& sectorTrack2 = mSectorTrackInfos[b2.TrackID()];
-  const GPUTPCTrack& inTrack2 = *sectorTrack2.OrigTrack();
+  MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_reference> inTrack2 = *sectorTrack2.OrigTrack();
   for (int32_t i = 0; i < inTrack2.NHits(); i++) {
     const GPUTPCHitId& ic2 = tracker2.TrackHits()[inTrack2.FirstHitID() + i];
     int32_t clusterIndex = tracker2.Data().ClusterDataIndex(tracker2.Data().Row(ic2.RowIndex()), ic2.HitIndex());
diff --git a/GPU/GPUTracking/Merger/GPUTPCGMSectorTrack.cxx b/GPU/GPUTracking/Merger/GPUTPCGMSectorTrack.cxx
index d7acdc56b87d8..ebe513fe51bdd 100644
--- a/GPU/GPUTracking/Merger/GPUTPCGMSectorTrack.cxx
+++ b/GPU/GPUTracking/Merger/GPUTPCGMSectorTrack.cxx
@@ -24,10 +24,10 @@
 using namespace o2::gpu;
 using namespace o2::tpc;
 
-GPUd() void GPUTPCGMSectorTrack::Set(const GPUTPCGMMerger* merger, MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_pointer> sectorTr, float alpha, int32_t sector)
+GPUd() void GPUTPCGMSectorTrack::Set(const GPUTPCGMMerger* merger, MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_reference> sectorTr, float alpha, int32_t sector)
 {
-  MemLayout::wrapper<GPUTPCBaseTrackParamSkeleton, MemLayout::const_reference> t = (*sectorTr).Param();
-  mOrigTrack = sectorTr;
+  MemLayout::wrapper<GPUTPCBaseTrackParamSkeleton, MemLayout::const_reference> t = sectorTr.Param();
+  mOrigTrack = &sectorTr;
   mParam.mX = t.GetX();
   mParam.mY = t.GetY();
   mParam.mZ = t.GetZ();
@@ -43,12 +43,12 @@ GPUd() void GPUTPCGMSectorTrack::Set(const GPUTPCGMMerger* merger, MemLayout::wr
   } else {
     mTZOffset = merger->GetConstantMem()->calibObjects.fastTransformHelper->getCorrMap()->convZOffsetToVertexTime(sector, t.GetZOffset(), merger->Param().continuousMaxTimeBin);
   }
-  mNClusters = (*sectorTr).NHits();
+  mNClusters = sectorTr->NHits();
 }
 
-GPUd() void GPUTPCGMSectorTrack::Set(const GPUTPCGMTrackParam& trk, MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_pointer> sectorTr, float alpha, int32_t sector)
+GPUd() void GPUTPCGMSectorTrack::Set(const GPUTPCGMTrackParam& trk, MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_reference> sectorTr, float alpha, int32_t sector)
 {
-  mOrigTrack = sectorTr;
+  mOrigTrack = &sectorTr;
   mParam.mX = trk.GetX();
   mParam.mY = trk.GetY();
   mParam.mZ = trk.GetZ();
@@ -60,7 +60,7 @@ GPUd() void GPUTPCGMSectorTrack::Set(const GPUTPCGMTrackParam& trk, MemLayout::w
   mAlpha = alpha;
   mSector = sector;
   mTZOffset = trk.GetTZOffset();
-  mNClusters = (*sectorTr).NHits();
+  mNClusters = sectorTr->NHits();
   mParam.mC0 = trk.GetCov(0);
   mParam.mC2 = trk.GetCov(2);
   mParam.mC3 = trk.GetCov(3);
@@ -96,11 +96,11 @@ GPUd() void GPUTPCGMSectorTrack::SetParam2(const GPUTPCGMTrackParam& trk)
 GPUd() bool GPUTPCGMSectorTrack::FilterErrors(const GPUTPCGMMerger* merger, int32_t iSector, float maxSinPhi, float sinPhiMargin)
 {
   float lastX;
-  // float lastX = GPUTPCGeometry::Row2X((*mOrigTrack).Cluster((*mOrigTrack).NClusters() - 1).GetRow()); // TODO: Why is this needed to be set below, Row2X should work, but looses some tracks
+  // float lastX = GPUTPCGeometry::Row2X(mOrigTrack->Cluster(mOrigTrack->NClusters() - 1).GetRow()); // TODO: Why is this needed to be set below, Row2X should work, but looses some tracks
   float y, z;
   int32_t row, index;
   const GPUTPCTracker& trk = merger->GetConstantMem()->tpcTrackers[iSector];
-  const GPUTPCHitId& ic = trk.TrackHits()[(*mOrigTrack).FirstHitID() + (*mOrigTrack).NHits() - 1];
+  const GPUTPCHitId& ic = trk.TrackHits()[mOrigTrack->FirstHitID() + mOrigTrack->NHits() - 1];
   index = trk.Data().ClusterDataIndex(trk.Data().Row(ic.RowIndex()), ic.HitIndex()) + merger->GetConstantMem()->ioPtrs.clustersNative->clusterOffset[iSector][0];
   row = ic.RowIndex();
   const ClusterNative& cl = merger->GetConstantMem()->ioPtrs.clustersNative->clustersLinear[index];
@@ -493,7 +493,7 @@ GPUd() bool GPUTPCGMSectorTrack::TransportToXAlpha(GPUTPCGMMerger* merger, float
 
 GPUd() void GPUTPCGMSectorTrack::CopyBaseTrackCov()
 {
-  const float* GPUrestrict() cov = (*mOrigTrack).Param().mC;
+  const float* GPUrestrict() cov = mOrigTrack->Param().mC;
   mParam.mC0 = cov[0];
   mParam.mC2 = cov[2];
   mParam.mC3 = cov[3];
diff --git a/GPU/GPUTracking/Merger/GPUTPCGMSectorTrack.h b/GPU/GPUTracking/Merger/GPUTPCGMSectorTrack.h
index 412e5fe16d748..fa5c0c850d14c 100644
--- a/GPU/GPUTracking/Merger/GPUTPCGMSectorTrack.h
+++ b/GPU/GPUTracking/Merger/GPUTPCGMSectorTrack.h
@@ -74,9 +74,9 @@ class GPUTPCGMSectorTrack
     mClusterZT[1] = v2;
   }
 
-  GPUd() void Set(const GPUTPCGMTrackParam& trk, MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_pointer> sectorTr, float alpha, int32_t sector);
+  GPUd() void Set(const GPUTPCGMTrackParam& trk, MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_reference> sectorTr, float alpha, int32_t sector);
   GPUd() void SetParam2(const GPUTPCGMTrackParam& trk);
-  GPUd() void Set(const GPUTPCGMMerger* merger, MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_pointer> sectorTr, float alpha, int32_t sector);
+  GPUd() void Set(const GPUTPCGMMerger* merger, MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::const_reference> sectorTr, float alpha, int32_t sector);
   GPUd() void UseParam2() { mParam = mParam2; }
   GPUd() void SetX2(float v) { mParam2.mX = v; }
   GPUd() float X2() const { return mParam2.mX; }
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCBaseTrackParam.h b/GPU/GPUTracking/SectorTracker/GPUTPCBaseTrackParam.h
index e8db8f131a7be..499d25905352e 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCBaseTrackParam.h
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCBaseTrackParam.h
@@ -21,22 +21,29 @@
 namespace o2::gpu
 {
 
-struct Covariance {
-    float m00, m01, m02, m03, m04, m05, m06, m07, m08, m09, m10, m11, m12, m13, m14;
-    constexpr operator float*() { return &m00; }
-    constexpr operator const float*() const { return &m00; }
-    constexpr float& operator[](int32_t i) { return *((&m00) + i); }
-    constexpr const float& operator[](int32_t i) const { return *((&m00) + i); }
+namespace detail
+{
+
+struct GPUTPCCovariance
+{
+  float data[15];
+  constexpr operator float*() { return data; }
+  constexpr operator const float*() const { return data; }
+  constexpr float& operator[](int32_t i) { return data[i]; }
+  constexpr const float& operator[](int32_t i) const { return data[i]; }
 };
 
-struct Parameters {
-    float mY, mZ, mSinPhi, mDzDs, mQPt;
-    constexpr operator float*() { return &mY; }
-    constexpr operator const float*() const { return &mY; }
-    constexpr float& operator[](int32_t i) { return *((&mY) + i); }
-    constexpr const float& operator[](int32_t i) const { return *((&mY) + i); }
+struct GPUTPCParameter
+{
+  float data[5]; // mY, mZ, mSinPhi, mDzDs, mQPt
+  constexpr operator float*() { return data; }
+  constexpr operator const float*() const { return data; }
+  constexpr float& operator[](int32_t i) { return data[i]; }
+  constexpr const float& operator[](int32_t i) const { return data[i]; }
 };
 
+}
+
 /**
  * @class GPUTPCBaseTrackParam
  *
@@ -95,9 +102,9 @@ struct GPUTPCBaseTrackParamSkeleton
   // This is neccessary for performance reasons!!!
   // Changes to Elements of this class therefore must also be applied to TrackletConstructor!!!
   F<float> mX;       // x position
-  F<Covariance> mC;   // the covariance matrix for Y,Z,SinPhi,..
+  F<detail::GPUTPCCovariance> mC;  // the covariance matrix for Y,Z,SinPhi,..
   F<float> mZOffset; // z offset
-  F<Parameters> mP;    // 'active' track parameters: Y, Z, SinPhi, DzDs, q/Pt
+  F<detail::GPUTPCParameter> mP;  // 'active' track parameters: Y, Z, SinPhi, DzDs, q/Pt
 };
 
 // Needed for sorting
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCExtrapolationTracking.cxx b/GPU/GPUTracking/SectorTracker/GPUTPCExtrapolationTracking.cxx
index 81e91e70d0737..e499dfe78c675 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCExtrapolationTracking.cxx
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCExtrapolationTracking.cxx
@@ -110,8 +110,7 @@ GPUd() int32_t GPUTPCExtrapolationTracking::PerformExtrapolationTrackingRun(GPUT
         rowIndex--;
       }
     }
-    //GPUTPCTrack& GPUrestrict() track = tracker.Tracks()[trackId];
-    MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::reference> track = tracker.Tracks()[trackId]; // reference_restrict
+    MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::reference_restrict> track = tracker.Tracks()[trackId];
     track.SetParam(tParam.GetParam());
     track.SetNHits(nHits);
     track.SetFirstHitID(hitId);
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCSectorDebugSortKernels.cxx b/GPU/GPUTracking/SectorTracker/GPUTPCSectorDebugSortKernels.cxx
index d2833442956be..f1bc69ead23bf 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCSectorDebugSortKernels.cxx
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCSectorDebugSortKernels.cxx
@@ -106,6 +106,6 @@ GPUdii() void GPUTPCSectorDebugSortKernels::Thread<GPUTPCSectorDebugSortKernels:
     }
     return tracker.TrackHits()[trk1.FirstHitID()].HitIndex() > tracker.TrackHits()[trk2.FirstHitID()].HitIndex();
   };
-  GPUCommonAlgorithm::sortDeviceDynamic(tracker.Tracks(), &tracker.Tracks()[tracker.CommonMemory()->nLocalTracks], sorter);
-  GPUCommonAlgorithm::sortDeviceDynamic(&tracker.Tracks()[tracker.CommonMemory()->nLocalTracks], &tracker.Tracks()[*tracker.NTracks()], sorter);
+  GPUCommonAlgorithm::sortDeviceDynamic(tracker.Tracks(), tracker.Tracks() + tracker.CommonMemory()->nLocalTracks, sorter);
+  GPUCommonAlgorithm::sortDeviceDynamic(tracker.Tracks() + tracker.CommonMemory()->nLocalTracks, tracker.Tracks() + *tracker.NTracks(), sorter);
 }
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTrack.h b/GPU/GPUTracking/SectorTracker/GPUTPCTrack.h
index 66cc17f2311e5..eca56e2f31ae8 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCTrack.h
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCTrack.h
@@ -21,6 +21,8 @@
 
 namespace o2::gpu
 {
+constexpr MemLayout::Flag GPUTPCTrackLayout = MemLayout::Flag::aos;
+
 /**
  * @class GPUTPCTrack
  *
@@ -36,6 +38,7 @@ class GPUTPCTrackSkeleton
     MEMLAYOUT_APPLY_BINARY(GPUTPCTrackSkeleton, MEMLAYOUT_EXPAND(mFirstHitID), MEMLAYOUT_EXPAND(mNHits), MEMLAYOUT_EXPAND(mLocalTrackId), MEMLAYOUT_EXPAND(mParam))
 
 //#if !defined(GPUCA_GPUCODE)
+//  GPUTPCTrack() : mFirstHitID(0), mNHits(0), mLocalTrackId(-1), mParam() {}
 //  ~GPUTPCTrackSkeleton() = default;
 //#endif //! GPUCA_GPUCODE
 
@@ -79,7 +82,7 @@ struct iterator_traits<MemLayout::wrapper<o2::gpu::GPUTPCTrackSkeleton, MemLayou
     using iterator_category = random_access_iterator_tag;
     using difference_type = MemLayout::ptrdiff_t;
     using value_type = MemLayout::wrapper<o2::gpu::GPUTPCTrackSkeleton, MemLayout::value>;
-    using pointer = void; //MemLayout::wrapper<o2::gpu::GPUTPCTrackSkeleton, MemLayout::pointer>;
+    using pointer = void;
     using reference = MemLayout::wrapper<o2::gpu::GPUTPCTrackSkeleton, MemLayout::reference>;
 };
 
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTrackLinearisation.h b/GPU/GPUTracking/SectorTracker/GPUTPCTrackLinearisation.h
index b13033632bdd7..703c43b4a8b4b 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCTrackLinearisation.h
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCTrackLinearisation.h
@@ -40,7 +40,7 @@ class GPUTPCTrackLinearisation
   GPUTPCTrackLinearisation() : mSinPhi(0), mCosPhi(1), mDzDs(0), mQPt(0) {}
   GPUTPCTrackLinearisation(float SinPhi1, float CosPhi1, float DzDs1, float QPt1) : mSinPhi(SinPhi1), mCosPhi(CosPhi1), mDzDs(DzDs1), mQPt(QPt1) {}
 
-  GPUd() GPUTPCTrackLinearisation(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::const_reference> t); // const_reference_restrict
+  GPUd() GPUTPCTrackLinearisation(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::const_reference_restrict> t);
 
   GPUd() void Set(float SinPhi1, float CosPhi1, float DzDs1, float QPt1);
 
@@ -66,7 +66,7 @@ class GPUTPCTrackLinearisation
   float mQPt;    // QPt
 };
 
-GPUdi() GPUTPCTrackLinearisation::GPUTPCTrackLinearisation(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::const_reference> t) : mSinPhi(t.SinPhi()), mCosPhi(0), mDzDs(t.DzDs()), mQPt(t.QPt()) // const_reference_restrict
+GPUdi() GPUTPCTrackLinearisation::GPUTPCTrackLinearisation(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::const_reference_restrict> t) : mSinPhi(t.SinPhi()), mCosPhi(0), mDzDs(t.DzDs()), mQPt(t.QPt())
 {
   if (mSinPhi > GPUCA_MAX_SIN_PHI) {
     mSinPhi = GPUCA_MAX_SIN_PHI;
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTrackParam.cxx b/GPU/GPUTracking/SectorTracker/GPUTPCTrackParam.cxx
index 42de67eb6b02b..7e99dda354648 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCTrackParam.cxx
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCTrackParam.cxx
@@ -15,7 +15,6 @@
 #include "GPUTPCTrackLinearisation.h"
 #include "GPUTPCTrackParam.h"
 #include "GPUTPCGeometry.h"
-#include "GPUTPCDef.h"
 #include "MemLayout.h"
 
 using namespace o2::gpu;
@@ -31,7 +30,7 @@ using namespace o2::gpu;
 //
 
 template <template <class> class F>
-GPUd() float GPUTPCTrackParamSkeleton<F>::GetDist2(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::const_reference> t) const // const_reference_restrict
+GPUd() float GPUTPCTrackParamSkeleton<F>::GetDist2(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::const_reference_restrict> t) const
 {
   // get squared distance between tracks
 
@@ -42,7 +41,7 @@ GPUd() float GPUTPCTrackParamSkeleton<F>::GetDist2(MemLayout::wrapper<GPUTPCTrac
 }
 
 template <template <class> class F>
-GPUd() float GPUTPCTrackParamSkeleton<F>::GetDistXZ2(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::const_reference> t) const // const_reference_restrict
+GPUd() float GPUTPCTrackParamSkeleton<F>::GetDistXZ2(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::const_reference_restrict> t) const
 {
   // get squared distance between tracks in X&Z
 
@@ -305,12 +304,12 @@ template <template <class> class F>
 GPUd() bool GPUTPCTrackParamSkeleton<F>::TransportToX(float x, float Bz, float maxSinPhi)
 {
   //* Transport the track parameters to X=x
-  GPUTPCTrackLinearisation t0(*static_cast<MemLayout::wrapper<GPUTPCTrackParamSkeleton, F>*>(this));
+  GPUTPCTrackLinearisation t0{{*this}};
   return TransportToX(x, t0, Bz, maxSinPhi);
 }
 
 template <template <class> class F>
-GPUd() bool GPUTPCTrackParamSkeleton<F>::TransportToXWithMaterial(float x, GPUTPCTrackLinearisation& GPUrestrict() t0, GPUTPCTrackFitParam& GPUrestrict() par, float Bz, float maxSinPhi)
+GPUd() bool GPUTPCTrackParamSkeleton<F>::TransportToXWithMaterial(float x, GPUTPCTrackLinearisation& GPUrestrict() t0, detail::GPUTPCTrackFitParam& GPUrestrict() par, float Bz, float maxSinPhi)
 {
   //* Transport the track parameters to X=x  taking into account material budget
 
@@ -329,10 +328,10 @@ GPUd() bool GPUTPCTrackParamSkeleton<F>::TransportToXWithMaterial(float x, GPUTP
 }
 
 template <template <class> class F>
-GPUd() bool GPUTPCTrackParamSkeleton<F>::TransportToXWithMaterial(float x, GPUTPCTrackFitParam& GPUrestrict() par, float Bz, float maxSinPhi)
+GPUd() bool GPUTPCTrackParamSkeleton<F>::TransportToXWithMaterial(float x, detail::GPUTPCTrackFitParam& GPUrestrict() par, float Bz, float maxSinPhi)
 {
   //* Transport the track parameters to X=x  taking into account material budget
-  GPUTPCTrackLinearisation t0(*static_cast<MemLayout::wrapper<GPUTPCTrackParamSkeleton, F>*>(this));
+  GPUTPCTrackLinearisation t0{{*this}};
   return TransportToXWithMaterial(x, t0, par, Bz, maxSinPhi);
 }
 
@@ -341,7 +340,7 @@ GPUd() bool GPUTPCTrackParamSkeleton<F>::TransportToXWithMaterial(float x, float
 {
   //* Transport the track parameters to X=x taking into account material budget
 
-  GPUTPCTrackFitParam par;
+  detail::GPUTPCTrackFitParam par;
   CalculateFitParameters(par);
   return TransportToXWithMaterial(x, par, Bz, maxSinPhi);
 }
@@ -440,7 +439,7 @@ GPUd() float GPUTPCTrackParamSkeleton<F>::ApproximateBetheBloch(float beta2)
 }
 
 template <template <class> class F>
-GPUd() void GPUTPCTrackParamSkeleton<F>::CalculateFitParameters(GPUTPCTrackFitParam& par, float mass)
+GPUd() void GPUTPCTrackParamSkeleton<F>::CalculateFitParameters(detail::GPUTPCTrackFitParam& par, float mass)
 {
   //*!
 
@@ -475,7 +474,7 @@ GPUd() void GPUTPCTrackParamSkeleton<F>::CalculateFitParameters(GPUTPCTrackFitPa
 }
 
 template <template <class> class F>
-GPUd() bool GPUTPCTrackParamSkeleton<F>::CorrectForMeanMaterial(float xOverX0, float xTimesRho, const GPUTPCTrackFitParam& par)
+GPUd() bool GPUTPCTrackParamSkeleton<F>::CorrectForMeanMaterial(float xOverX0, float xTimesRho, const detail::GPUTPCTrackFitParam& par)
 {
   //------------------------------------------------------------------
   // This function corrects the track parameters for the crossed material.
@@ -879,5 +878,5 @@ GPUd() int32_t GPUTPCTrackParamSkeleton<F>::GetPropagatedYZ(float bz, float x, f
 namespace o2::gpu {
   template class GPUTPCTrackParamSkeleton<MemLayout::value>;
   template class GPUTPCTrackParamSkeleton<MemLayout::reference>;
-  //template class GPUTPCTrackParamSkeleton<MemLayout::reference_restrict>;
+  template class GPUTPCTrackParamSkeleton<MemLayout::reference_restrict>;
 } // namespace o2::gpu
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTrackParam.h b/GPU/GPUTracking/SectorTracker/GPUTPCTrackParam.h
index e7b6b02efb722..7070011521b6d 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCTrackParam.h
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCTrackParam.h
@@ -25,6 +25,15 @@ namespace o2::gpu
 
 class GPUTPCTrackLinearisation;
 
+namespace detail
+{
+
+struct GPUTPCTrackFitParam {
+  float bethe, e, theta2, EP2, sigmadE2, k22, k33, k43, k44; // parameters
+};
+
+}
+
 /**
  * @class GPUTPCTrackParam
  *
@@ -39,10 +48,6 @@ class GPUTPCTrackParamSkeleton
   MEMLAYOUT_APPLY_UNARY(mParam, mSignCosPhi, mChi2, mNDF)
   MEMLAYOUT_APPLY_BINARY(GPUTPCTrackParamSkeleton, MEMLAYOUT_EXPAND(mParam), MEMLAYOUT_EXPAND(mSignCosPhi), MEMLAYOUT_EXPAND(mChi2), MEMLAYOUT_EXPAND(mNDF))
 
-  struct GPUTPCTrackFitParam {
-    float bethe, e, theta2, EP2, sigmadE2, k22, k33, k43, k44; // parameters
-  };
-
   GPUd() MemLayout::wrapper<GPUTPCBaseTrackParamSkeleton, MemLayout::const_reference> GetParam() const { return mParam; }
   GPUd() void SetParam(MemLayout::wrapper<GPUTPCBaseTrackParamSkeleton, MemLayout::const_reference> v);
   GPUd() void InitParam();
@@ -98,8 +103,8 @@ class GPUTPCTrackParamSkeleton
   GPUd() void SetChi2(float v) { mChi2 = v; }
   GPUd() void SetNDF(int32_t v) { mNDF = v; }
 
-  GPUd() float GetDist2(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::const_reference> t) const; // const_reference_restrict
-  GPUd() float GetDistXZ2(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::const_reference> t) const; // const_reference_restrict
+  GPUd() float GetDist2(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::const_reference_restrict> t) const;
+  GPUd() float GetDistXZ2(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::const_reference_restrict> t) const;
 
   GPUd() float GetS(float x, float y, float Bz) const;
 
@@ -112,17 +117,17 @@ class GPUTPCTrackParamSkeleton
 
   GPUd() bool TransportToX(float x, float sinPhi0, float cosPhi0, float Bz, float maxSinPhi = GPUCA_MAX_SIN_PHI);
 
-  GPUd() bool TransportToXWithMaterial(float x, GPUTPCTrackLinearisation& t0, GPUTPCTrackFitParam& par, float Bz, float maxSinPhi = GPUCA_MAX_SIN_PHI);
+  GPUd() bool TransportToXWithMaterial(float x, GPUTPCTrackLinearisation& t0, detail::GPUTPCTrackFitParam& par, float Bz, float maxSinPhi = GPUCA_MAX_SIN_PHI);
 
-  GPUd() bool TransportToXWithMaterial(float x, GPUTPCTrackFitParam& par, float Bz, float maxSinPhi = GPUCA_MAX_SIN_PHI);
+  GPUd() bool TransportToXWithMaterial(float x, detail::GPUTPCTrackFitParam& par, float Bz, float maxSinPhi = GPUCA_MAX_SIN_PHI);
 
   GPUd() static float ApproximateBetheBloch(float beta2);
   GPUd() static float BetheBlochGeant(float bg, float kp0 = 2.33f, float kp1 = 0.20f, float kp2 = 3.00f, float kp3 = 173e-9f, float kp4 = 0.49848f);
   GPUd() static float BetheBlochSolid(float bg);
   GPUd() static float BetheBlochGas(float bg);
 
-  GPUd() void CalculateFitParameters(GPUTPCTrackFitParam& par, float mass = 0.13957f);
-  GPUd() bool CorrectForMeanMaterial(float xOverX0, float xTimesRho, const GPUTPCTrackFitParam& par);
+  GPUd() void CalculateFitParameters(detail::GPUTPCTrackFitParam& par, float mass = 0.13957f);
+  GPUd() bool CorrectForMeanMaterial(float xOverX0, float xTimesRho, const detail::GPUTPCTrackFitParam& par);
 
   GPUd() bool Rotate(float alpha, float maxSinPhi = GPUCA_MAX_SIN_PHI);
   GPUd() bool Rotate(float alpha, GPUTPCTrackLinearisation& t0, float maxSinPhi = GPUCA_MAX_SIN_PHI);
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTracker.cxx b/GPU/GPUTracking/SectorTracker/GPUTPCTracker.cxx
index bb3f09aa6f93d..3a61c455431f7 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCTracker.cxx
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCTracker.cxx
@@ -14,12 +14,14 @@
 
 #include "GPUTPCTracker.h"
 #include "GPUTPCRow.h"
-#include "GPUTPCTrack.h"
 #include "GPUCommonMath.h"
+#include "MemLayout.h"
 
 #include "GPUTPCClusterData.h"
 #include "GPUO2DataTypes.h"
 #include "GPUTPCTrackParam.h"
+#include "GPUTPCTracklet.h"
+#include "GPUTPCTrack.h"
 #include "GPUParam.inc"
 #include "GPUTPCConvertImpl.h"
 #include "GPUDefParametersRuntime.h"
@@ -29,6 +31,7 @@
 #include <cmath>
 #include <algorithm>
 #include <stdexcept>
+#include <type_traits>
 
 #include "GPUReconstruction.h"
 #include "GPUMemorySizeScalers.h"
@@ -108,40 +111,41 @@ void GPUTPCTracker::RegisterMemoryAllocation()
   mMemoryResOutput = mRec->RegisterMemoryAllocation(this, &GPUTPCTracker::SetPointersOutput, type, "TPCTrackerTracks");
 }
 
-GPUhd() void GPUTPCTracker::SetPointersTrackletsHelper(void* & mem, GPUTPCTracker::TrackletArrayType<MemLayout::Flag::aos>& tracklets) {
-    computePointerWithAlignment(mem, tracklets, mNMaxTracklets);
-}
+namespace {
+
+template <class Function>
+struct ApplyMemberwise {
+    Function g;
+
+    template <class ...Args>
+    void operator()(Args& ...args) const { (g(args), ...); }
+};
+
+template <class Function>
+struct ApplyRecursive {
+    Function f;
 
-GPUhd() void GPUTPCTracker::SetPointersTrackletsHelper(void* & mem, GPUTPCTracker::TrackletArrayType<MemLayout::Flag::soa>& tracklets) {
-  computePointerWithAlignment(mem, tracklets.mFirstRow, mNMaxTracklets);
-  computePointerWithAlignment(mem, tracklets.mLastRow, mNMaxTracklets);
+    template <class T>
+    void operator()(T * & aos) const { f(aos); }
 
-  computePointerWithAlignment(mem, tracklets.mParam.mX, mNMaxTracklets);
-  computePointerWithAlignment(mem, tracklets.mParam.mC, mNMaxTracklets);
-  computePointerWithAlignment(mem, tracklets.mParam.mZOffset, mNMaxTracklets);
-  computePointerWithAlignment(mem, tracklets.mParam.mP, mNMaxTracklets);
+    template <template <template <class> class> class S>
+    void operator()(S<MemLayout::pointer>& soa) const { soa.apply(ApplyMemberwise<ApplyRecursive>{f}); }
+};
 
-  computePointerWithAlignment(mem, tracklets.mHitWeight, mNMaxTracklets);
-  computePointerWithAlignment(mem, tracklets.mFirstHit, mNMaxTracklets);
 }
 
 GPUhd() void* GPUTPCTracker::SetPointersTracklets(void* mem)
 {
-  SetPointersTrackletsHelper(mem, mTracklets);
+  auto tracklet_helper = [&mem, this](auto& tracklets) -> void { computePointerWithAlignment(mem, tracklets, mNMaxTracklets); };
+  ApplyRecursive{tracklet_helper}(mTracklets);
   computePointerWithAlignment(mem, mTrackletRowHits, mNMaxRowHits);
   return mem;
 }
 
 GPUhd() void* GPUTPCTracker::SetPointersOutput(void* mem)
 {
-  computePointerWithAlignment(mem, mTracks.mFirstHitID, mNMaxTracklets);
-  computePointerWithAlignment(mem, mTracks.mNHits, mNMaxTracklets);
-  computePointerWithAlignment(mem, mTracks.mLocalTrackId, mNMaxTracklets);
-  computePointerWithAlignment(mem, mTracks.mParam.mX, mNMaxTracklets);
-  computePointerWithAlignment(mem, mTracks.mParam.mC, mNMaxTracklets);
-  computePointerWithAlignment(mem, mTracks.mParam.mZOffset, mNMaxTracklets);
-  computePointerWithAlignment(mem, mTracks.mParam.mP, mNMaxTracklets);
-  //computePointerWithAlignment(mem, mTracks, mNMaxTracks);
+  auto track_helper = [&mem, this](auto& tracks) -> void { computePointerWithAlignment(mem, tracks, mNMaxTracks); };
+  ApplyRecursive{track_helper}(mTracks);
   computePointerWithAlignment(mem, mTrackHits, mNMaxTrackHits);
   return mem;
 }
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTracker.h b/GPU/GPUTracking/SectorTracker/GPUTPCTracker.h
index 8fe7b0832f046..768bd06fabd55 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCTracker.h
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCTracker.h
@@ -34,7 +34,7 @@ namespace o2::gpu
 struct GPUTPCClusterData;
 struct GPUParam;
 template <template <class> class F>
-class GPUTPCTrackSkeleton;
+class GPUTPCTrackParamSkeleton;
 class GPUTPCRow;
 
 class GPUTPCTracker : public GPUProcessor
@@ -62,9 +62,6 @@ class GPUTPCTracker : public GPUProcessor
   void DumpTrackletHits(std::ostream& out);         // Same for Track Hits
 #endif
 
-  template <MemLayout::Flag layout>
-  using TrackletArrayType = MemLayout::interface<GPUTPCTrackletSkeleton, MemLayout::pointer, layout>::type;
-
   struct commonMemoryStruct {
     GPUAtomic(uint32_t) nStartHits = 0; // number of start hits
     GPUAtomic(uint32_t) nTracklets = 0; // number of tracklets
@@ -110,9 +107,8 @@ class GPUTPCTracker : public GPUProcessor
   void* SetPointersScratch(void* mem);
   void* SetPointersScratchHost(void* mem);
   void* SetPointersCommon(void* mem);
-  void SetPointersTrackletsHelper(void* & mem, TrackletArrayType<MemLayout::Flag::aos>& tracklets);
-  void SetPointersTrackletsHelper(void* & mem, TrackletArrayType<MemLayout::Flag::soa>& tracklets);
   void* SetPointersTracklets(void* mem);
+  void* SetPointersTracks(void* mem);
   void* SetPointersOutput(void* mem);
   void RegisterMemoryAllocation();
 
@@ -193,14 +189,11 @@ class GPUTPCTracker : public GPUProcessor
   GPUhd() GPUglobalref() GPUTPCHitId* TrackletStartHits() { return mTrackletStartHits; }
   GPUhd() GPUglobalref() GPUTPCHitId* TrackletTmpStartHits() const { return mTrackletTmpStartHits; }
 
-  GPUhd() GPUglobalref() MemLayout::wrapper<GPUTPCTrackletSkeleton, MemLayout::reference> Tracklet(int32_t i) { return mTracklets[i]; } // reference_restrict
-
-  template <MemLayout::Flag layout>
-  GPUhd() GPUglobalref() TrackletArrayType<layout> Tracklets() const { return mTracklets; }
+  GPUhd() GPUglobalref() auto Tracklets() const { return mTracklets; }
   GPUhd() GPUglobalref() calink* TrackletRowHits() const { return mTrackletRowHits; }
 
   GPUhd() GPUglobalref() GPUAtomic(uint32_t) * NTracks() const { return &mCommonMem->nTracks; }
-  GPUhd() GPUglobalref() MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::pointer> Tracks() const { return mTracks; }
+  GPUhd() GPUglobalref() auto Tracks() const { return mTracks; }
   GPUhd() GPUglobalref() GPUAtomic(uint32_t) * NTrackHits() const { return &mCommonMem->nTrackHits; }
   GPUhd() GPUglobalref() GPUTPCHitId* TrackHits() const { return mTrackHits; }
 
@@ -248,12 +241,12 @@ class GPUTPCTracker : public GPUProcessor
   GPUglobalref() char* mGPUTrackletTemp = nullptr;             // Temp Memory for GPU Tracklet Constructor
 
   // event
-  GPUglobalref() commonMemoryStruct* mCommonMem = nullptr;  // common event memory
-  GPUglobalref() GPUTPCHitId* mTrackletStartHits = nullptr; // start hits for the tracklets
-  GPUglobalref() TrackletArrayType<MemLayout::Flag::soa> mTracklets; // tracklets
-  GPUglobalref() calink* mTrackletRowHits = nullptr;        // Hits for each Tracklet in each row
-  GPUglobalref() MemLayout::wrapper<GPUTPCTrackSkeleton, MemLayout::pointer> mTracks; // reconstructed tracks
-  GPUglobalref() GPUTPCHitId* mTrackHits = nullptr;         // array of track hit numbers
+  GPUglobalref() commonMemoryStruct* mCommonMem = nullptr;           // common event memory
+  GPUglobalref() GPUTPCHitId* mTrackletStartHits = nullptr;          // start hits for the tracklets
+  GPUglobalref() MemLayout::interface<GPUTPCTrackletSkeleton, MemLayout::pointer, GPUTPCTrackletLayout>::type mTracklets; // tracklets
+  GPUglobalref() calink* mTrackletRowHits = nullptr;                 // Hits for each Tracklet in each row
+  GPUglobalref() MemLayout::interface<GPUTPCTrackSkeleton, MemLayout::pointer, GPUTPCTrackLayout>::type mTracks;       // reconstructed tracks
+  GPUglobalref() GPUTPCHitId* mTrackHits = nullptr;                  // array of track hit numbers
 
   static int32_t StarthitSortComparison(const void* a, const void* b);
 };
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTrackerDump.cxx b/GPU/GPUTracking/SectorTracker/GPUTPCTrackerDump.cxx
index f291c9cc2fb9a..3c1a7dea70db3 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCTrackerDump.cxx
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCTrackerDump.cxx
@@ -15,6 +15,7 @@
 #include "GPUTPCTracker.h"
 #include "GPUReconstruction.h"
 #include "GPUTPCHitId.h"
+#include "GPUTPCTracklet.h"
 #include "GPUTPCTrack.h"
 #include "GPULogging.h"
 #include "MemLayout.h"
@@ -130,22 +131,22 @@ void GPUTPCTracker::DumpTrackletHits(std::ostream& out)
   std::iota(Ids.begin(), Ids.end(), 0);
   if (mRec->GetProcessingSettings().deterministicGPUReconstruction) {
     std::sort(Ids.begin(), Ids.end(), [this](const int32_t& a, const int32_t& b) {
-      if (this->Tracklet(a).FirstRow() != this->Tracklet(b).FirstRow()) {
-        return this->Tracklet(a).FirstRow() > this->Tracklet(b).FirstRow();
+      if (this->Tracklets()[a].FirstRow() != this->Tracklets()[b].FirstRow()) {
+        return this->Tracklets()[a].FirstRow() > this->Tracklets()[b].FirstRow();
       }
-      if (this->Tracklet(a).LastRow() != this->Tracklet(b).LastRow()) {
-        return this->Tracklet(a).LastRow() > this->Tracklet(b).LastRow();
+      if (this->Tracklets()[a].LastRow() != this->Tracklets()[b].LastRow()) {
+        return this->Tracklets()[a].LastRow() > this->Tracklets()[b].LastRow();
       }
-      if (this->Tracklet(a).Param().Y() != this->Tracklet(b).Param().Y()) {
-        return this->Tracklet(a).Param().Y() > this->Tracklet(b).Param().Y();
+      if (this->Tracklets()[a].Param().Y() != this->Tracklets()[b].Param().Y()) {
+        return this->Tracklets()[a].Param().Y() > this->Tracklets()[b].Param().Y();
       }
-      return this->Tracklet(a).Param().Z() > this->Tracklet(b).Param().Z();
+      return this->Tracklets()[a].Param().Z() > this->Tracklets()[b].Param().Z();
     });
   }
   for (int32_t jj = 0; jj < nTracklets; jj++) {
     const int32_t j = Ids[jj];
-    const auto& tracklet = Tracklet(j);
-    out << "Tracklet " << std::setw(4) << jj << " (Rows: " << Tracklet(j).FirstRow() << " - " << tracklet.LastRow() << ", Weight " << Tracklet(j).HitWeight() << ") ";
+    MemLayout::wrapper<GPUTPCTrackletSkeleton, MemLayout::const_reference> tracklet = Tracklets()[j];
+    out << "Tracklet " << std::setw(4) << jj << " (Rows: " << Tracklets()[j].FirstRow() << " - " << tracklet.LastRow() << ", Weight " << Tracklets()[j].HitWeight() << ") ";
     if (tracklet.LastRow() > tracklet.FirstRow() && (tracklet.FirstRow() >= GPUCA_ROW_COUNT || tracklet.LastRow() >= GPUCA_ROW_COUNT)) {
       GPUError("Error: Tracklet %d First %d Last %d", j, tracklet.FirstRow(), tracklet.LastRow());
       out << " (Error: Tracklet " << j << " First " << tracklet.FirstRow() << " Last " << tracklet.LastRow() << ") ";
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTracklet.h b/GPU/GPUTracking/SectorTracker/GPUTPCTracklet.h
index df069cdb35d60..3d5d5cf734952 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCTracklet.h
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCTracklet.h
@@ -21,6 +21,8 @@
 
 namespace o2::gpu
 {
+constexpr MemLayout::Flag GPUTPCTrackletLayout = MemLayout::Flag::aos;
+
 /**
  * @class GPUTPCTracklet
  *
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.cxx b/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.cxx
index 777b55bbe5594..db7adeb788e38 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.cxx
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.cxx
@@ -30,13 +30,13 @@
 
 using namespace o2::gpu;
 
-GPUdii() void GPUTPCTrackletConstructor::InitTracklet(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference> tParam) // reference_restrict
+GPUdii() void GPUTPCTrackletConstructor::InitTracklet(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference_restrict> tParam)
 {
   // Initialize Tracklet Parameters using default values
   tParam.InitParam();
 }
 
-GPUd() bool GPUTPCTrackletConstructor::CheckCov(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference> tParam) // reference_restrict
+GPUd() bool GPUTPCTrackletConstructor::CheckCov(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference_restrict> tParam)
 {
   bool ok = 1;
   const float* c = tParam.Cov();
@@ -53,7 +53,7 @@ GPUd() bool GPUTPCTrackletConstructor::CheckCov(MemLayout::wrapper<GPUTPCTrackPa
   return (ok);
 }
 
-GPUd() void GPUTPCTrackletConstructor::StoreTracklet(int32_t /*nBlocks*/, int32_t /*nThreads*/, int32_t /*iBlock*/, int32_t /*iThread*/, GPUsharedref() GPUSharedMemory& s, GPUTPCThreadMemory& GPUrestrict() r, GPUconstantref() GPUTPCTracker& GPUrestrict() tracker, MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference> tParam, calink* rowHits) // reference_restrict
+GPUd() void GPUTPCTrackletConstructor::StoreTracklet(int32_t /*nBlocks*/, int32_t /*nThreads*/, int32_t /*iBlock*/, int32_t /*iThread*/, GPUsharedref() GPUSharedMemory& s, GPUTPCThreadMemory& GPUrestrict() r, GPUconstantref() GPUTPCTracker& GPUrestrict() tracker, MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference_restrict> tParam, calink* rowHits)
 {
   // reconstruction of tracklets, tracklet store step
   const uint32_t nHits = r.mLastRow + 1 - r.mFirstRow;
@@ -79,7 +79,7 @@ GPUd() void GPUTPCTrackletConstructor::StoreTracklet(int32_t /*nBlocks*/, int32_
     return;
   }
 
-  GPUglobalref() MemLayout::wrapper<GPUTPCTrackletSkeleton, MemLayout::reference> tracklet = tracker.Tracklet(itrout); // reference_restrict
+  GPUglobalref() MemLayout::wrapper<GPUTPCTrackletSkeleton, MemLayout::reference_restrict> tracklet = tracker.Tracklets()[itrout];
 
   CADEBUG(printf("    Storing tracklet: %d rows\n", nHits));
 
@@ -104,7 +104,7 @@ GPUd() void GPUTPCTrackletConstructor::StoreTracklet(int32_t /*nBlocks*/, int32_
 }
 
 template <class T>
-GPUdic(2, 1) void GPUTPCTrackletConstructor::UpdateTracklet(int32_t /*nBlocks*/, int32_t /*nThreads*/, int32_t /*iBlock*/, int32_t /*iThread*/, GPUsharedref() T& s, GPUTPCThreadMemory& GPUrestrict() r, GPUconstantref() GPUTPCTracker& GPUrestrict() tracker, MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference> tParam, int32_t iRow, calink& rowHit, calink* rowHits) // reference_restrict
+GPUdic(2, 1) void GPUTPCTrackletConstructor::UpdateTracklet(int32_t /*nBlocks*/, int32_t /*nThreads*/, int32_t /*iBlock*/, int32_t /*iThread*/, GPUsharedref() T& s, GPUTPCThreadMemory& GPUrestrict() r, GPUconstantref() GPUTPCTracker& GPUrestrict() tracker, MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference_restrict> tParam, int32_t iRow, calink& rowHit, calink* rowHits)
 {
   // reconstruction of tracklets, tracklets update step
   CA_MAKE_SHARED_REF(GPUTPCRow, row, tracker.Row(iRow), s.mRows[iRow]);
@@ -491,7 +491,7 @@ GPUdii() void GPUTPCTrackletConstructor::Thread(int32_t nBlocks, int32_t nThread
 }
 
 template <> // FIXME: GPUgeneric() needed to make the clang spirv output link correctly
-GPUd() int32_t GPUTPCTrackletConstructor::GPUTPCTrackletConstructorExtrapolationTracking<GPUgeneric() GPUTPCExtrapolationTracking::GPUSharedMemory>(GPUconstantref() GPUTPCTracker& GPUrestrict() tracker, GPUsharedref() GPUTPCExtrapolationTracking::GPUSharedMemory& sMem, MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference> tParam, int32_t row, int32_t increment, int32_t iTracklet, calink* rowHits) // reference_restrict
+GPUd() int32_t GPUTPCTrackletConstructor::GPUTPCTrackletConstructorExtrapolationTracking<GPUgeneric() GPUTPCExtrapolationTracking::GPUSharedMemory>(GPUconstantref() GPUTPCTracker& GPUrestrict() tracker, GPUsharedref() GPUTPCExtrapolationTracking::GPUSharedMemory& sMem, MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference_restrict> tParam, int32_t row, int32_t increment, int32_t iTracklet, calink* rowHits)
 {
   GPUTPCThreadMemory rMem;
   rMem.mISH = iTracklet;
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.h b/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.h
index d6c315a104082..9272cfc9adf0a 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.h
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.h
@@ -72,18 +72,18 @@ class GPUTPCTrackletConstructor : public GPUKernelTemplate
 #endif                // GPUCA_TRACKLET_CONSTRUCTOR_DO_PROFILE
   };
 
-  GPUd() static void InitTracklet(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference> tParam); // reference_restrict
+  GPUd() static void InitTracklet(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference_restrict> tParam);
 
   template <class T>
-  GPUd() static void UpdateTracklet(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() T& s, GPUTPCThreadMemory& r, GPUconstantref() GPUTPCTracker& tracker, MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference> tParam, int32_t iRow, calink& rowHit, calink* rowHits); // reference_restrict
+  GPUd() static void UpdateTracklet(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() T& s, GPUTPCThreadMemory& r, GPUconstantref() GPUTPCTracker& tracker, MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference_restrict> tParam, int32_t iRow, calink& rowHit, calink* rowHits);
 
-  GPUd() static void StoreTracklet(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& s, GPUTPCThreadMemory& r, GPUconstantref() GPUTPCTracker& tracker, MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference> tParam, calink* rowHits); // reference_restrict
+  GPUd() static void StoreTracklet(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& s, GPUTPCThreadMemory& r, GPUconstantref() GPUTPCTracker& tracker, MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference_restrict> tParam, calink* rowHits);
 
-  GPUd() static bool CheckCov(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference> tParam); // reference_restrict
+  GPUd() static bool CheckCov(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference_restrict> tParam);
   GPUd() static void DoTracklet(GPUconstantref() GPUTPCTracker& tracker, GPUsharedref() GPUTPCTrackletConstructor::GPUSharedMemory& sMem, GPUTPCThreadMemory& rMem);
 
   template <class T>
-  GPUd() static int32_t GPUTPCTrackletConstructorExtrapolationTracking(GPUconstantref() GPUTPCTracker& tracker, GPUsharedref() T& sMem, MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference> tParam, int32_t startrow, int32_t increment, int32_t iTracklet, calink* rowHits); // reference_restrict
+  GPUd() static int32_t GPUTPCTrackletConstructorExtrapolationTracking(GPUconstantref() GPUTPCTracker& tracker, GPUsharedref() T& sMem, MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference_restrict> tParam, int32_t startrow, int32_t increment, int32_t iTracklet, calink* rowHits);
 
   typedef GPUconstantref() GPUTPCTracker processorType;
   GPUhdi() constexpr static GPUDataTypes::RecoStep GetRecoStep() { return GPUDataTypes::RecoStep::TPCSectorTracking; }
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTrackletSelector.cxx b/GPU/GPUTracking/SectorTracker/GPUTPCTrackletSelector.cxx
index 592998af560e2..8b6b020e41da7 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCTrackletSelector.cxx
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCTrackletSelector.cxx
@@ -40,7 +40,7 @@ GPUdii() void GPUTPCTrackletSelector::Thread<0>(int32_t nBlocks, int32_t nThread
   for (int32_t itr = s.mItr0 + iThread; itr < s.mNTracklets; itr += s.mNThreadsTotal) {
     GPUbarrierWarp();
 
-    GPUglobalref() MemLayout::wrapper<GPUTPCTrackletSkeleton, MemLayout::reference> tracklet = tracker.Tracklet(itr); // reference_restrict
+    GPUglobalref() MemLayout::wrapper<GPUTPCTrackletSkeleton, MemLayout::reference_restrict> tracklet = tracker.Tracklets()[itr];
 
     int32_t firstRow = tracklet.FirstRow();
     int32_t lastRow = tracklet.LastRow();
diff --git a/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx b/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx
index 5a46e07c8ee2c..c87ff65de8da5 100644
--- a/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx
+++ b/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx
@@ -287,7 +287,7 @@ GPUDisplay::vboList GPUDisplay::DrawTracks(const GPUTPCTracker& tracker, int32_t
   }
   size_t startCount = mVertexBufferStart[iSector].size();
   for (uint32_t i = (global ? tracker.CommonMemory()->nLocalTracks : 0); i < (global ? *tracker.NTracks() : tracker.CommonMemory()->nLocalTracks); i++) {
-    GPUTPCTrack& track = tracker.Tracks()[i];
+    GPUTPCTrackSkeleton<MemLayout::reference> track = tracker.Tracks()[i];
     size_t startCountInner = mVertexBuffer[iSector].size();
     for (int32_t j = 0; j < track.NHits(); j++) {
       const GPUTPCHitId& hit = tracker.TrackHits()[track.FirstHitID() + j];