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 db57e7ec06d4b..a1bbdd930b4cb 100644
--- a/GPU/Common/GPUCommonAlgorithm.h
+++ b/GPU/Common/GPUCommonAlgorithm.h
@@ -28,19 +28,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 +225,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 +235,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 +249,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 +259,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..be74e2143ae37 100644
--- a/GPU/Common/GPUCommonAlgorithmThrust.h
+++ b/GPU/Common/GPUCommonAlgorithmThrust.h
@@ -25,6 +25,7 @@
 
 #include "GPUCommonDef.h"
 #include "GPUCommonHelpers.h"
+#include "GPUTPCTrack.h"
 
 #ifndef __HIPCC__ // CUDA
 #include <cub/cub.cuh>
@@ -81,19 +82,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..250ad36435474 100644
--- a/GPU/Common/MemLayout.h
+++ b/GPU/Common/MemLayout.h
@@ -1,257 +1,308 @@
 #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 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 {
-    size_t i;
+    MemLayout::size_t i;
+    template <class... Args>
+    constexpr SF operator()(Args& ...args) const { return {args[i]...}; }
     template <class... Args>
-    constexpr SF operator()(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...}; }
 };
 
 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(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<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); }
 };
 
 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(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>>{}); }
 };
 
 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>>{});
+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(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;
     }
-    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);
+    constexpr wrapper& operator=(const wrapper& other) {
+        Base::apply(other, CopyAssignment{});
+        return *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);
+    constexpr wrapper& operator=(const wrapper<S, const_reference>& 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, reference_restrict>& 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, const_reference_restrict>& other) {
+        Base::apply(other, CopyAssignment{});
+        return *this;
     }
-};
 
-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>>{});
-    }
+    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 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>>{});
+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;
     }
-    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);
+    constexpr wrapper& operator=(const wrapper& other) {
+        Base::apply(other, CopyAssignment{});
+        return *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 wrapper& operator=(const wrapper<S, reference>& other) {
+        Base::apply(other, CopyAssignment{});
+        return *this;
     }
-    constexpr S<pointer_restrtict> operator& () {
-        return static_cast<Derived*>(this)->apply(GetPointer<S<pointer_restrtict>>{});
+    constexpr wrapper& operator=(const wrapper<S, const_reference>& other) {
+        Base::apply(other, CopyAssignment{});
+        return *this;
     }
-    constexpr S<const_pointer_restrict> operator& () const {
-        return static_cast<const Derived*>(this)->apply(GetPointer<S<const_pointer_restrict>>{});
+    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 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_reference> : public S<const_reference> {
+    using Base = S<const_reference>;
+
+    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_restrict>& 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,
-    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; }
+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>
+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<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<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; }
+};
+
+template <template <template <class> class> class S>
+struct wrapper<S, const_pointer> : public S<const_pointer> {
+    using Base = S<const_pointer>;
+
+    constexpr wrapper() = default;
+    constexpr wrapper(Base b) : Base{static_cast<Base&&>(b)} {}
+    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<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 reference operator*() { return handle[index]; }
+    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 };
+
+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(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>\
+    constexpr STRUCT_NAME apply(const STRUCT_NAME<F_other>& other, Function&& f) const { return {__VA_ARGS__}; }\
 
 #endif // MEMLAYOUT_H
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 6713b0a906bf1..32f87719ab2d8 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::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 9d2578731a30c..ea96587a5ff41 100644
--- a/GPU/GPUTracking/Global/GPUChainTracking.cxx
+++ b/GPU/GPUTracking/Global/GPUChainTracking.cxx
@@ -549,12 +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]);
+    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 2a9ce862d86f7..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,7 +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;
-    std::unique_ptr<GPUTPCTrackSkeleton<MemLayout::value>[]> 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/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..83dc127e1f82e 100644
--- a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx
+++ b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx
@@ -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_reference> inTrack, float alpha, int32_t sector)
 {
   GPUTPCGMPropagator prop;
   prop.SetMaterialTPC();
@@ -497,14 +497,14 @@ 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
+  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) {
@@ -513,14 +513,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.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->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];
@@ -550,11 +550,11 @@ 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_reference> 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.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) {
@@ -574,12 +574,11 @@ 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];
   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;
     }
@@ -611,10 +610,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_reference> sectorTr = trk.Tracks()[itr];
     GPUTPCGMSectorTrack track;
     SetTrackClusterZT(track, iSector, sectorTr);
     if (Param().rec.tpc.mergerCovSource == 0) {
@@ -642,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;
   }
 }
@@ -710,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) {
diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMerger.h b/GPU/GPUTracking/Merger/GPUTPCGMMerger.h
index 54a541ebe0fd6..0bad9628d1ab6 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,6 @@ class GPUTPCGMMerger : public GPUProcessor
     GPUAtomic(uint32_t) nMergedTrackClusters;
     GPUAtomic(uint32_t) nO2Tracks;
     GPUAtomic(uint32_t) nO2ClusRefs;
-    const GPUTPCTrack* firstExtrapolatedTracks[NSECTORS];
     GPUAtomic(uint32_t) tmpCounter[2 * NSECTORS];
     GPUAtomic(uint32_t) nLooperMatchCandidates;
   };
@@ -153,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, 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_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 d86a63ec61c53..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, const GPUTPCTrack* 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)
 {
-  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();
@@ -46,9 +46,9 @@ GPUd() void GPUTPCGMSectorTrack::Set(const GPUTPCGMMerger* merger, const GPUTPCT
   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_reference> sectorTr, float alpha, int32_t sector)
 {
-  mOrigTrack = sectorTr;
+  mOrigTrack = &sectorTr;
   mParam.mX = trk.GetX();
   mParam.mY = trk.GetY();
   mParam.mZ = trk.GetZ();
@@ -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..fa5c0c850d14c 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_reference> 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_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; }
@@ -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..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
  *
@@ -45,11 +52,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,54 +98,25 @@ 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!!!
   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
 };
 
-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..e499dfe78c675 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,7 @@ GPUd() int32_t GPUTPCExtrapolationTracking::PerformExtrapolationTrackingRun(GPUT
         rowIndex--;
       }
     }
-    GPUTPCTrack& GPUrestrict() track = tracker.Tracks()[trackId];
+    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 7981ef5af26d8..f1bc69ead23bf 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();
     }
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTrack.h b/GPU/GPUTracking/SectorTracker/GPUTPCTrack.h
index de4b12b6f0d4a..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
  *
@@ -29,55 +31,61 @@ 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)
+//  GPUTPCTrack() : mFirstHitID(0), mNHits(0), mLocalTrackId(-1), mParam() {}
+//  ~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;
+    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..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(GPUTPCTrackParamSkeleton<MemLayout::const_reference_restrict> t);
+  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(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_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 d141ce7e4a7f5..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(GPUTPCTrackParamSkeleton<MemLayout::const_reference_restrict> t) const
+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(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_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(*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(*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.
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTrackParam.h b/GPU/GPUTracking/SectorTracker/GPUTPCTrackParam.h
index 5536114c88752..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
  *
@@ -33,19 +42,14 @@ 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)
-
-  struct GPUTPCTrackFitParam {
-    float bethe, e, theta2, EP2, sigmadE2, k22, k33, k43, k44; // parameters
-  };
+  MEMLAYOUT_APPLY_UNARY(mParam, mSignCosPhi, mChi2, mNDF)
+  MEMLAYOUT_APPLY_BINARY(GPUTPCTrackParamSkeleton, MEMLAYOUT_EXPAND(mParam), MEMLAYOUT_EXPAND(mSignCosPhi), MEMLAYOUT_EXPAND(mChi2), MEMLAYOUT_EXPAND(mNDF))
 
-  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 +103,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_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;
 
@@ -113,17 +117,17 @@ class GPUTPCTrackParamSkeleton : public MemLayout::CRTP<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);
@@ -146,10 +150,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 +164,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..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,33 +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, 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 5af0939188391..768bd06fabd55 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"
 
@@ -33,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
@@ -61,9 +62,6 @@ class GPUTPCTracker : public GPUProcessor
   void DumpTrackletHits(std::ostream& out);         // Same for Track Hits
 #endif
 
-  template <MemLayout::Flag layout>
-  using TrackletArrayType = MemLayout::wrapper<GPUTPCTrackletSkeleton, MemLayout::pointer, layout>::type;
-
   struct commonMemoryStruct {
     GPUAtomic(uint32_t) nStartHits = 0; // number of start hits
     GPUAtomic(uint32_t) nTracklets = 0; // number of tracklets
@@ -84,13 +82,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);
@@ -109,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();
 
@@ -192,14 +189,11 @@ 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]; }
-
-  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() GPUTPCTrackSkeleton<MemLayout::value>* 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; }
 
@@ -247,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() GPUTPCTrackSkeleton<MemLayout::value>* mTracks = nullptr; // 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 3c3f113aa8c21..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"
@@ -102,7 +103,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) << " ";
@@ -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 95bc83e8ee65f..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
  *
@@ -28,49 +30,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..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(GPUTPCTrackParamSkeleton<MemLayout::reference_restrict> tParam)
+GPUdii() void GPUTPCTrackletConstructor::InitTracklet(MemLayout::wrapper<GPUTPCTrackParamSkeleton, MemLayout::reference_restrict> tParam)
 {
   // 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_restrict> tParam)
 {
   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_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() GPUTPCTrackletSkeleton<MemLayout::reference_restrict> tracklet = tracker.Tracklet(itrout);
+  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, 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_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]);
@@ -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_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 c4792b9e00ec5..9272cfc9adf0a 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_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, 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_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, 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_restrict> tParam, calink* rowHits);
 
+  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, 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_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 01823a641d92d..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() GPUTPCTrackletSkeleton<MemLayout::reference_restrict> tracklet = tracker.Tracklet(itr);
+    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 c25efb21e2279..c87ff65de8da5 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++) {
@@ -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];