diff --git a/src/atlas/array/Array.h b/src/atlas/array/Array.h
index 1e3958ee0..aa9379ca5 100644
--- a/src/atlas/array/Array.h
+++ b/src/atlas/array/Array.h
@@ -13,7 +13,10 @@
 #include <memory>
 #include <vector>
 
+#include "eckit/config/Parametrisation.h"
+
 #include "atlas/library/config.h"
+#include "atlas/util/Config.h"
 
 #include "atlas/util/Object.h"
 
@@ -21,6 +24,10 @@
 #include "atlas/array/DataType.h"
 #include "atlas/array_fwd.h"
 
+namespace eckit {
+class Parametrisation;
+}
+
 namespace atlas {
 namespace array {
 
@@ -40,38 +47,48 @@ class Array : public util::Object {
     Array() = default;
     virtual ~Array();
 
-    static Array* create(array::DataType, const ArrayShape&);
+    static Array* create(array::DataType, const ArrayShape&, const eckit::Parametrisation& param 
+            = util::Config());
 
-    static Array* create(array::DataType, const ArrayShape&, const ArrayLayout&);
+    static Array* create(array::DataType, const ArrayShape&, const ArrayLayout&,
+            const eckit::Parametrisation& param = util::Config());
 
-    static Array* create(array::DataType, ArraySpec&&);
+    static Array* create(array::DataType, ArraySpec&&,
+            const eckit::Parametrisation& param = util::Config());
 
-    static Array* create(ArraySpec&&);
+    static Array* create(ArraySpec&&,
+            const eckit::Parametrisation& param = util::Config());
 
     virtual size_t footprint() const = 0;
 
     template <typename Value>
-    static Array* create(idx_t size0);
+    static Array* create(idx_t size0, const eckit::Parametrisation& param = util::Config());
     template <typename Value>
-    static Array* create(idx_t size0, idx_t size1);
+    static Array* create(idx_t size0, idx_t size1, const eckit::Parametrisation& param = util::Config());
     template <typename Value>
-    static Array* create(idx_t size0, idx_t size1, idx_t size2);
+    static Array* create(idx_t size0, idx_t size1, idx_t size2,
+            const eckit::Parametrisation& param = util::Config());
     template <typename Value>
-    static Array* create(idx_t size0, idx_t size1, idx_t size2, idx_t size3);
+    static Array* create(idx_t size0, idx_t size1, idx_t size2, idx_t size3,
+            const eckit::Parametrisation& param = util::Config());
     template <typename Value>
-    static Array* create(idx_t size0, idx_t size1, idx_t size2, idx_t size3, idx_t size4);
+    static Array* create(idx_t size0, idx_t size1, idx_t size2, idx_t size3, idx_t size4,
+            const eckit::Parametrisation& param = util::Config());
 
     template <typename Value>
-    static Array* create(const ArrayShape& shape);
+    static Array* create(const ArrayShape& shape, const eckit::Parametrisation& param = util::Config());
 
     template <typename Value>
-    static Array* create(const ArrayShape& shape, const ArrayLayout& layout);
+    static Array* create(const ArrayShape& shape, const ArrayLayout& layout,
+            const eckit::Parametrisation& param = util::Config());
 
     template <typename Value>
-    static Array* wrap(Value* data, const ArrayShape& shape);
+    static Array* wrap(Value* data, const ArrayShape& shape, const eckit::Parametrisation& param 
+            = util::Config());
 
     template <typename Value>
-    static Array* wrap(Value* data, const ArraySpec& spec);
+    static Array* wrap(Value* data, const ArraySpec& spec, const eckit::Parametrisation& param 
+            = util::Config());
 
     idx_t bytes() const { return datatype().size() * spec().allocatedSize(); }
 
@@ -202,19 +219,21 @@ class Array : public util::Object {
 template <typename Value>
 class ArrayT : public Array {
 public:
-    ArrayT(idx_t size0);
-    ArrayT(idx_t size0, idx_t size1);
-    ArrayT(idx_t size0, idx_t size1, idx_t size2);
-    ArrayT(idx_t size0, idx_t size1, idx_t size2, idx_t size3);
-    ArrayT(idx_t size0, idx_t size1, idx_t size2, idx_t size3, idx_t size4);
+    ArrayT(idx_t size0, const eckit::Parametrisation& param = util::Config());
+    ArrayT(idx_t size0, idx_t size1, const eckit::Parametrisation& param = util::Config());
+    ArrayT(idx_t size0, idx_t size1, idx_t size2, const eckit::Parametrisation& param = util::Config());
+    ArrayT(idx_t size0, idx_t size1, idx_t size2, idx_t size3,
+            const eckit::Parametrisation& param = util::Config());
+    ArrayT(idx_t size0, idx_t size1, idx_t size2, idx_t size3, idx_t size4,
+            const eckit::Parametrisation& param = util::Config());
 
-    ArrayT(ArraySpec&&);
+    ArrayT(ArraySpec&&, const eckit::Parametrisation& param = util::Config());
 
-    ArrayT(const ArrayShape&);
+    ArrayT(const ArrayShape&, const eckit::Parametrisation& param = util::Config());
 
-    ArrayT(const ArrayShape&, const ArrayAlignment&);
+    ArrayT(const ArrayShape&, const ArrayAlignment&, const eckit::Parametrisation& param = util::Config());
 
-    ArrayT(const ArrayShape&, const ArrayLayout&);
+    ArrayT(const ArrayShape&, const ArrayLayout&, const eckit::Parametrisation& param = util::Config());
 
     virtual void insert(idx_t idx1, idx_t size1);
 
@@ -234,7 +253,7 @@ class ArrayT : public Array {
 
     // This constructor is used through the Array::create() or the Array::wrap()
     // methods
-    ArrayT(ArrayDataStore*, const ArraySpec&);
+    ArrayT(ArrayDataStore*, const ArraySpec&, const eckit::Parametrisation& param = util::Config());
 
     virtual size_t footprint() const;
 
diff --git a/src/atlas/array/ArrayDataStore.h b/src/atlas/array/ArrayDataStore.h
index f5963ecf0..9377c9246 100644
--- a/src/atlas/array/ArrayDataStore.h
+++ b/src/atlas/array/ArrayDataStore.h
@@ -73,6 +73,9 @@ class ArrayDataStore {
     Value* deviceData() {
         return static_cast<Value*>(voidDeviceData());
     }
+private:
+    bool device_memory_pinned_ = false;
+    bool device_memory_mapped_ = false;
 };
 
 #ifndef DOXYGEN_SHOULD_SKIP_THIS
diff --git a/src/atlas/array/native/NativeArray.cc b/src/atlas/array/native/NativeArray.cc
index 2f2442ba6..3106166d1 100644
--- a/src/atlas/array/native/NativeArray.cc
+++ b/src/atlas/array/native/NativeArray.cc
@@ -10,12 +10,15 @@
 
 #include <iostream>
 
+#include "eckit/config/Parametrisation.h"
+
 #include "atlas/array.h"
 #include "atlas/array/ArrayDataStore.h"
 #include "atlas/array/MakeView.h"
 #include "atlas/array/helpers/ArrayInitializer.h"
 #include "atlas/array/helpers/ArrayWriter.h"
 #include "atlas/array/native/NativeDataStore.h"
+#include "atlas/util/Config.h"
 #include "atlas/runtime/Exception.h"
 
 using namespace atlas::array::helpers;
@@ -24,61 +27,63 @@ namespace atlas {
 namespace array {
 
 template <typename Value>
-Array* Array::create(idx_t dim0) {
-    return new ArrayT<Value>(dim0);
+Array* Array::create(idx_t dim0, const eckit::Parametrisation& param) {
+    return new ArrayT<Value>(dim0, param);
 }
 template <typename Value>
-Array* Array::create(idx_t dim0, idx_t dim1) {
-    return new ArrayT<Value>(dim0, dim1);
+Array* Array::create(idx_t dim0, idx_t dim1, const eckit::Parametrisation& param) {
+    return new ArrayT<Value>(dim0, dim1, param);
 }
 template <typename Value>
-Array* Array::create(idx_t dim0, idx_t dim1, idx_t dim2) {
-    return new ArrayT<Value>(dim0, dim1, dim2);
+Array* Array::create(idx_t dim0, idx_t dim1, idx_t dim2, const eckit::Parametrisation& param) {
+    return new ArrayT<Value>(dim0, dim1, dim2, param);
 }
 template <typename Value>
-Array* Array::create(idx_t dim0, idx_t dim1, idx_t dim2, idx_t dim3) {
-    return new ArrayT<Value>(dim0, dim1, dim2, dim3);
+Array* Array::create(idx_t dim0, idx_t dim1, idx_t dim2, idx_t dim3, const eckit::Parametrisation& param) {
+    return new ArrayT<Value>(dim0, dim1, dim2, dim3, param);
 }
 template <typename Value>
-Array* Array::create(idx_t dim0, idx_t dim1, idx_t dim2, idx_t dim3, idx_t dim4) {
-    return new ArrayT<Value>(dim0, dim1, dim2, dim3, dim4);
+Array* Array::create(idx_t dim0, idx_t dim1, idx_t dim2, idx_t dim3, idx_t dim4,
+        const eckit::Parametrisation& param) {
+    return new ArrayT<Value>(dim0, dim1, dim2, dim3, dim4, param);
 }
 template <typename Value>
-Array* Array::create(const ArrayShape& shape) {
-    return new ArrayT<Value>(shape);
+Array* Array::create(const ArrayShape& shape, const eckit::Parametrisation& param) {
+    return new ArrayT<Value>(shape, param);
 }
 template <typename Value>
-Array* Array::create(const ArrayShape& shape, const ArrayLayout& layout) {
-    return new ArrayT<Value>(shape, layout);
+Array* Array::create(const ArrayShape& shape, const ArrayLayout& layout,
+        const eckit::Parametrisation& param) {
+    return new ArrayT<Value>(shape, layout, param);
 }
 template <typename Value>
-Array* Array::wrap(Value* data, const ArrayShape& shape) {
+Array* Array::wrap(Value* data, const ArrayShape& shape, const eckit::Parametrisation& param) {
     size_t size = 1;
     for (int i = 0; i < shape.size(); ++i) {
         size *= shape[i];
     }
-    return new ArrayT<Value>(new native::WrappedDataStore<Value>(data, size), shape);
+    return new ArrayT<Value>(new native::WrappedDataStore<Value>(data, size, param), shape, param);
 }
 template <typename Value>
-Array* Array::wrap(Value* data, const ArraySpec& spec) {
+Array* Array::wrap(Value* data, const ArraySpec& spec, const eckit::Parametrisation& param) {
     size_t size = spec.size();
-    return new ArrayT<Value>(new native::WrappedDataStore<Value>(data, size), spec);
+    return new ArrayT<Value>(new native::WrappedDataStore<Value>(data, size, param), spec, param);
 }
 
 Array::~Array() = default;
 
-Array* Array::create(DataType datatype, const ArrayShape& shape) {
+Array* Array::create(DataType datatype, const ArrayShape& shape, const eckit::Parametrisation& param) {
     switch (datatype.kind()) {
         case DataType::KIND_REAL64:
-            return new ArrayT<double>(shape);
+            return new ArrayT<double>(shape, param);
         case DataType::KIND_REAL32:
-            return new ArrayT<float>(shape);
+            return new ArrayT<float>(shape, param);
         case DataType::KIND_INT32:
-            return new ArrayT<int>(shape);
+            return new ArrayT<int>(shape, param);
         case DataType::KIND_INT64:
-            return new ArrayT<long>(shape);
+            return new ArrayT<long>(shape, param);
         case DataType::KIND_UINT64:
-            return new ArrayT<unsigned long>(shape);
+            return new ArrayT<unsigned long>(shape, param);
         default: {
             std::stringstream err;
             err << "data kind " << datatype.kind() << " not recognised.";
@@ -87,18 +92,18 @@ Array* Array::create(DataType datatype, const ArrayShape& shape) {
     }
 }
 
-Array* Array::create(DataType datatype, ArraySpec&& spec) {
+Array* Array::create(DataType datatype, ArraySpec&& spec, const eckit::Parametrisation& param) {
     switch (datatype.kind()) {
         case DataType::KIND_REAL64:
-            return new ArrayT<double>(std::move(spec));
+            return new ArrayT<double>(std::move(spec), param);
         case DataType::KIND_REAL32:
-            return new ArrayT<float>(std::move(spec));
+            return new ArrayT<float>(std::move(spec), param);
         case DataType::KIND_INT32:
-            return new ArrayT<int>(std::move(spec));
+            return new ArrayT<int>(std::move(spec), param);
         case DataType::KIND_INT64:
-            return new ArrayT<long>(std::move(spec));
+            return new ArrayT<long>(std::move(spec), param);
         case DataType::KIND_UINT64:
-            return new ArrayT<unsigned long>(std::move(spec));
+            return new ArrayT<unsigned long>(std::move(spec), param);
         default: {
             std::stringstream err;
             err << "data kind " << datatype.kind() << " not recognised.";
@@ -107,72 +112,75 @@ Array* Array::create(DataType datatype, ArraySpec&& spec) {
     }
 }
 
-Array* Array::create(ArraySpec&& spec) {
-    return create(spec.datatype(), std::move(spec));
+Array* Array::create(ArraySpec&& spec, const eckit::Parametrisation& param) {
+    return create(spec.datatype(), std::move(spec), param);
 }
 
 
 template <typename Value>
-ArrayT<Value>::ArrayT(ArrayDataStore* ds, const ArraySpec& spec) {
+ArrayT<Value>::ArrayT(ArrayDataStore* ds, const ArraySpec& spec, const eckit::Parametrisation& param) {
     data_store_ = std::unique_ptr<ArrayDataStore>(ds);
     spec_       = spec;
 }
 
 template <typename Value>
-ArrayT<Value>::ArrayT(idx_t dim0) {
+ArrayT<Value>::ArrayT(idx_t dim0, const eckit::Parametrisation& param) {
     spec_       = ArraySpec(make_shape(dim0));
-    data_store_ = std::make_unique<native::DataStore<Value>>(spec_.size());
+    data_store_ = std::make_unique<native::DataStore<Value>>(spec_.size(), param);
 }
 template <typename Value>
-ArrayT<Value>::ArrayT(idx_t dim0, idx_t dim1) {
+ArrayT<Value>::ArrayT(idx_t dim0, idx_t dim1, const eckit::Parametrisation& param) {
     spec_       = ArraySpec(make_shape(dim0, dim1));
-    data_store_ = std::make_unique<native::DataStore<Value>>(spec_.size());
+    data_store_ = std::make_unique<native::DataStore<Value>>(spec_.size(), param);
 }
 template <typename Value>
-ArrayT<Value>::ArrayT(idx_t dim0, idx_t dim1, idx_t dim2) {
+ArrayT<Value>::ArrayT(idx_t dim0, idx_t dim1, idx_t dim2, const eckit::Parametrisation& param) {
     spec_       = ArraySpec(make_shape(dim0, dim1, dim2));
-    data_store_ = std::make_unique<native::DataStore<Value>>(spec_.size());
+    data_store_ = std::make_unique<native::DataStore<Value>>(spec_.size(), param);
 }
 template <typename Value>
-ArrayT<Value>::ArrayT(idx_t dim0, idx_t dim1, idx_t dim2, idx_t dim3) {
+ArrayT<Value>::ArrayT(idx_t dim0, idx_t dim1, idx_t dim2, idx_t dim3, const eckit::Parametrisation& param) {
     spec_       = ArraySpec(make_shape(dim0, dim1, dim2, dim3));
-    data_store_ = std::make_unique<native::DataStore<Value>>(spec_.size());
+    data_store_ = std::make_unique<native::DataStore<Value>>(spec_.size(), param);
 }
 template <typename Value>
-ArrayT<Value>::ArrayT(idx_t dim0, idx_t dim1, idx_t dim2, idx_t dim3, idx_t dim4) {
+ArrayT<Value>::ArrayT(idx_t dim0, idx_t dim1, idx_t dim2, idx_t dim3, idx_t dim4,
+        const eckit::Parametrisation& param) {
     spec_       = ArraySpec(make_shape(dim0, dim1, dim2, dim3, dim4));
-    data_store_ = std::make_unique<native::DataStore<Value>>(spec_.size());
+    data_store_ = std::make_unique<native::DataStore<Value>>(spec_.size(), param);
 }
 
 template <typename Value>
-ArrayT<Value>::ArrayT(const ArrayShape& shape) {
+ArrayT<Value>::ArrayT(const ArrayShape& shape, const eckit::Parametrisation& param) {
     ATLAS_ASSERT(shape.size() > 0);
     size_t size = 1;
     for (size_t j = 0; j < shape.size(); ++j) {
         size *= size_t(shape[j]);
     }
-    data_store_ = std::make_unique<native::DataStore<Value>>(size);
+    data_store_ = std::make_unique<native::DataStore<Value>>(size, param);
     spec_       = ArraySpec(shape);
 }
 
 template <typename Value>
-ArrayT<Value>::ArrayT(const ArrayShape& shape, const ArrayAlignment& alignment) {
+ArrayT<Value>::ArrayT(const ArrayShape& shape, const ArrayAlignment& alignment,
+        const eckit::Parametrisation& param) {
     spec_       = ArraySpec(shape, alignment);
-    data_store_ = std::make_unique<native::DataStore<Value>>(spec_.allocatedSize());
+    data_store_ = std::make_unique<native::DataStore<Value>>(spec_.allocatedSize(), param);
 }
 
 template <typename Value>
-ArrayT<Value>::ArrayT(const ArrayShape& shape, const ArrayLayout& layout) {
+ArrayT<Value>::ArrayT(const ArrayShape& shape, const ArrayLayout& layout, 
+        const eckit::Parametrisation& param) {
     spec_       = ArraySpec(shape);
-    data_store_ = std::make_unique<native::DataStore<Value>>(spec_.size());
+    data_store_ = std::make_unique<native::DataStore<Value>>(spec_.size(), param);
     for (size_t j = 0; j < layout.size(); ++j) {
         ATLAS_ASSERT(spec_.layout()[j] == layout[j]);
     }
 }
 
 template <typename Value>
-ArrayT<Value>::ArrayT(ArraySpec&& spec): Array(std::move(spec)) {
-    data_store_ = std::make_unique<native::DataStore<Value>>(spec_.allocatedSize());
+ArrayT<Value>::ArrayT(ArraySpec&& spec, const eckit::Parametrisation& param): Array(std::move(spec)) {
+    data_store_ = std::make_unique<native::DataStore<Value>>(spec_.allocatedSize(), param);
 }
 
 template <typename Value>
@@ -301,59 +309,59 @@ bool ArrayT<Value>::accMapped() const {
 
 //------------------------------------------------------------------------------
 
-template Array* Array::create<int>(idx_t);
-template Array* Array::create<long>(idx_t);
-template Array* Array::create<float>(idx_t);
-template Array* Array::create<double>(idx_t);
-template Array* Array::create<long unsigned>(idx_t);
-
-template Array* Array::create<int>(idx_t, idx_t);
-template Array* Array::create<long>(idx_t, idx_t);
-template Array* Array::create<float>(idx_t, idx_t);
-template Array* Array::create<double>(idx_t, idx_t);
-template Array* Array::create<long unsigned>(idx_t, idx_t);
-
-template Array* Array::create<int>(idx_t, idx_t, idx_t);
-template Array* Array::create<long>(idx_t, idx_t, idx_t);
-template Array* Array::create<float>(idx_t, idx_t, idx_t);
-template Array* Array::create<double>(idx_t, idx_t, idx_t);
-template Array* Array::create<long unsigned>(idx_t, idx_t, idx_t);
-
-template Array* Array::create<int>(idx_t, idx_t, idx_t, idx_t);
-template Array* Array::create<long>(idx_t, idx_t, idx_t, idx_t);
-template Array* Array::create<float>(idx_t, idx_t, idx_t, idx_t);
-template Array* Array::create<double>(idx_t, idx_t, idx_t, idx_t);
-template Array* Array::create<long unsigned>(idx_t, idx_t, idx_t, idx_t);
-
-template Array* Array::create<int>(idx_t, idx_t, idx_t, idx_t, idx_t);
-template Array* Array::create<long>(idx_t, idx_t, idx_t, idx_t, idx_t);
-template Array* Array::create<float>(idx_t, idx_t, idx_t, idx_t, idx_t);
-template Array* Array::create<double>(idx_t, idx_t, idx_t, idx_t, idx_t);
-template Array* Array::create<long unsigned>(idx_t, idx_t, idx_t, idx_t, idx_t);
-
-template Array* Array::create<int>(const ArrayShape&);
-template Array* Array::create<long>(const ArrayShape&);
-template Array* Array::create<float>(const ArrayShape&);
-template Array* Array::create<double>(const ArrayShape&);
-template Array* Array::create<long unsigned>(const ArrayShape&);
-
-template Array* Array::create<int>(const ArrayShape&, const ArrayLayout&);
-template Array* Array::create<long>(const ArrayShape&, const ArrayLayout&);
-template Array* Array::create<float>(const ArrayShape&, const ArrayLayout&);
-template Array* Array::create<double>(const ArrayShape&, const ArrayLayout&);
-template Array* Array::create<long unsigned>(const ArrayShape&, const ArrayLayout&);
-
-template Array* Array::wrap<int>(int*, const ArrayShape&);
-template Array* Array::wrap<long>(long*, const ArrayShape&);
-template Array* Array::wrap<float>(float*, const ArrayShape&);
-template Array* Array::wrap<double>(double*, const ArrayShape&);
-template Array* Array::wrap<long unsigned>(long unsigned*, const ArrayShape&);
-
-template Array* Array::wrap<int>(int*, const ArraySpec&);
-template Array* Array::wrap<long>(long*, const ArraySpec&);
-template Array* Array::wrap<float>(float*, const ArraySpec&);
-template Array* Array::wrap<double>(double*, const ArraySpec&);
-template Array* Array::wrap<long unsigned>(long unsigned*, const ArraySpec&);
+template Array* Array::create<int>(idx_t, const eckit::Parametrisation&);
+template Array* Array::create<long>(idx_t, const eckit::Parametrisation&);
+template Array* Array::create<float>(idx_t, const eckit::Parametrisation&);
+template Array* Array::create<double>(idx_t, const eckit::Parametrisation&);
+template Array* Array::create<long unsigned>(idx_t, const eckit::Parametrisation&);
+
+template Array* Array::create<int>(idx_t, idx_t, const eckit::Parametrisation&);
+template Array* Array::create<long>(idx_t, idx_t, const eckit::Parametrisation&);
+template Array* Array::create<float>(idx_t, idx_t, const eckit::Parametrisation&);
+template Array* Array::create<double>(idx_t, idx_t, const eckit::Parametrisation&);
+template Array* Array::create<long unsigned>(idx_t, idx_t, const eckit::Parametrisation&);
+
+template Array* Array::create<int>(idx_t, idx_t, idx_t, const eckit::Parametrisation&);
+template Array* Array::create<long>(idx_t, idx_t, idx_t, const eckit::Parametrisation&);
+template Array* Array::create<float>(idx_t, idx_t, idx_t, const eckit::Parametrisation&);
+template Array* Array::create<double>(idx_t, idx_t, idx_t, const eckit::Parametrisation&);
+template Array* Array::create<long unsigned>(idx_t, idx_t, idx_t, const eckit::Parametrisation&);
+
+template Array* Array::create<int>(idx_t, idx_t, idx_t, idx_t, const eckit::Parametrisation&);
+template Array* Array::create<long>(idx_t, idx_t, idx_t, idx_t, const eckit::Parametrisation&);
+template Array* Array::create<float>(idx_t, idx_t, idx_t, idx_t, const eckit::Parametrisation&);
+template Array* Array::create<double>(idx_t, idx_t, idx_t, idx_t, const eckit::Parametrisation&);
+template Array* Array::create<long unsigned>(idx_t, idx_t, idx_t, idx_t, const eckit::Parametrisation&);
+
+template Array* Array::create<int>(idx_t, idx_t, idx_t, idx_t, idx_t, const eckit::Parametrisation&);
+template Array* Array::create<long>(idx_t, idx_t, idx_t, idx_t, idx_t, const eckit::Parametrisation&);
+template Array* Array::create<float>(idx_t, idx_t, idx_t, idx_t, idx_t, const eckit::Parametrisation&);
+template Array* Array::create<double>(idx_t, idx_t, idx_t, idx_t, idx_t, const eckit::Parametrisation&);
+template Array* Array::create<long unsigned>(idx_t, idx_t, idx_t, idx_t, idx_t, const eckit::Parametrisation&);
+
+template Array* Array::create<int>(const ArrayShape&, const eckit::Parametrisation&);
+template Array* Array::create<long>(const ArrayShape&, const eckit::Parametrisation&);
+template Array* Array::create<float>(const ArrayShape&, const eckit::Parametrisation&);
+template Array* Array::create<double>(const ArrayShape&, const eckit::Parametrisation&);
+template Array* Array::create<long unsigned>(const ArrayShape&, const eckit::Parametrisation&);
+
+template Array* Array::create<int>(const ArrayShape&, const ArrayLayout&, const eckit::Parametrisation&);
+template Array* Array::create<long>(const ArrayShape&, const ArrayLayout&, const eckit::Parametrisation&);
+template Array* Array::create<float>(const ArrayShape&, const ArrayLayout&, const eckit::Parametrisation&);
+template Array* Array::create<double>(const ArrayShape&, const ArrayLayout&, const eckit::Parametrisation&);
+template Array* Array::create<long unsigned>(const ArrayShape&, const ArrayLayout&, const eckit::Parametrisation&);
+
+template Array* Array::wrap<int>(int*, const ArrayShape&, const eckit::Parametrisation&);
+template Array* Array::wrap<long>(long*, const ArrayShape&, const eckit::Parametrisation&);
+template Array* Array::wrap<float>(float*, const ArrayShape&, const eckit::Parametrisation&);
+template Array* Array::wrap<double>(double*, const ArrayShape&, const eckit::Parametrisation&);
+template Array* Array::wrap<long unsigned>(long unsigned*, const ArrayShape&, const eckit::Parametrisation&);
+
+template Array* Array::wrap<int>(int*, const ArraySpec&, const eckit::Parametrisation&);
+template Array* Array::wrap<long>(long*, const ArraySpec&, const eckit::Parametrisation&);
+template Array* Array::wrap<float>(float*, const ArraySpec&, const eckit::Parametrisation&);
+template Array* Array::wrap<double>(double*, const ArraySpec&, const eckit::Parametrisation&);
+template Array* Array::wrap<long unsigned>(long unsigned*, const ArraySpec&, const eckit::Parametrisation&);
 
 template class ArrayT<int>;
 template class ArrayT<long>;
diff --git a/src/atlas/array/native/NativeDataStore.h b/src/atlas/array/native/NativeDataStore.h
index 2499219b8..043e02d16 100644
--- a/src/atlas/array/native/NativeDataStore.h
+++ b/src/atlas/array/native/NativeDataStore.h
@@ -101,11 +101,21 @@ void initialise(Value[], size_t) {}
 template <typename Value>
 class DataStore : public ArrayDataStore {
 public:
-    DataStore(size_t size): size_(size) {
+    DataStore(size_t size, const eckit::Parametrisation& param): size_(size) {
         allocateHost();
         initialise(host_data_, size_);
 #if ATLAS_HAVE_CUDA
         device_updated_ = false;
+        std::string param_in;
+        if (param.get("host_memory_pinned", param_in)) {
+            host_memory_pinned_ = (param_in == "true");
+        }
+        if (param.get("host_mapped_mapped", param_in)) {
+            host_memory_mapped_ = (param_in == "true");
+        }
+        if (! host_memory_pinned_ && host_memory_mapped_) {
+            throw_AssertionFailed("Host memory can not be mapped when it is not pinned.", Here());
+        }
 #else
         device_data_ = host_data_;
 #endif
@@ -124,13 +134,15 @@ class DataStore : public ArrayDataStore {
             }
             allocateDevice();
         }
-        if (atlas::Library::instance().traceMemory()) {
-            Log::trace() << "updateDevice(" << name() << ") : cudaMemcpyHostToDevice( device_ptr:" << device_data_ << " , host_ptr:"<< host_data_ << " , " << eckit::Bytes(size_*sizeof(Value)) << " ) " << std::endl;
-        }
+        if (! host_memory_mapped_) {
+            if (atlas::Library::instance().traceMemory()) {
+                Log::trace() << "updateDevice(" << name() << ") : cudaMemcpyHostToDevice( device_ptr:" << device_data_ << " , host_ptr:"<< host_data_ << " , " << eckit::Bytes(size_*sizeof(Value)) << " ) " << std::endl;
+            }
 
-        cudaError_t err = cudaMemcpy(device_data_, host_data_, size_*sizeof(Value), cudaMemcpyHostToDevice);
-        if (err != cudaSuccess) {
-            throw_AssertionFailed("Failed to updateDevice("+std::string(name())+") : "+std::string(cudaGetErrorString(err)), Here());
+            cudaError_t err = cudaMemcpy(device_data_, host_data_, size_*sizeof(Value), cudaMemcpyHostToDevice);
+            if (err != cudaSuccess) {
+                throw_AssertionFailed("Failed to updateDevice("+std::string(name())+") : "+std::string(cudaGetErrorString(err)), Here());
+            }
         }
         device_updated_ = true;
 #endif
@@ -142,9 +154,17 @@ class DataStore : public ArrayDataStore {
             if (atlas::Library::instance().traceMemory()) {
                 Log::trace() << "updateHost(" << name() << ") : cudaMemcpyDeviceToHost( host_ptr:" << host_data_ << " , device_ptr:"<< device_data_ << " , " << eckit::Bytes(size_*sizeof(Value)) << " ) " << std::endl;
             }
-            cudaError_t err = cudaMemcpy(host_data_, device_data_, size_*sizeof(Value), cudaMemcpyDeviceToHost);
-            if (err != cudaSuccess) {
-                throw_AssertionFailed("Failed to updateHost("+std::string(name())+") : "+std::string(cudaGetErrorString(err)), Here());
+            if (!host_memory_mapped_) {
+                cudaError_t err = cudaMemcpy(host_data_, device_data_, size_*sizeof(Value), cudaMemcpyDeviceToHost);
+                if (err != cudaSuccess) {
+                    throw_AssertionFailed("Failed to updateHost("+std::string(name())+") : "+std::string(cudaGetErrorString(err)), Here());
+                }
+            }
+            else {
+                cudaError_t err = cudaDeviceSynchronize();
+                if (err != cudaSuccess) {
+                    throw_AssertionFailed("Failed to sync device: "+std::string(cudaGetErrorString(err)), Here());
+                }
             }
             host_updated_ = true;
         }
@@ -182,9 +202,23 @@ class DataStore : public ArrayDataStore {
         }
         if (size_) {
             size_t bytes = sizeof(Value)*size_;
-            cudaError_t err = cudaMalloc((void**)&device_data_, bytes);
-            if (err != cudaSuccess) {
-                throw_AssertionFailed("allocateDevice("+std::string(name())+") Failed to allocate GPU memory: " + std::string(cudaGetErrorString(err)), Here());
+            if (host_memory_pinned_) {
+                cudaError_t err = cudaHostRegister(host_data_, bytes, cudaHostRegisterMapped);
+                if (err != cudaSuccess) {
+                    throw_AssertionFailed("Failed to get device pointer: "+std::string(cudaGetErrorString(err)), Here());
+                }
+            }
+            if (host_memory_mapped_) {
+                cudaError_t err = cudaHostGetDevicePointer((void**)&device_data_, host_data_, 0);
+                if (err != cudaSuccess) {
+                    throw_AssertionFailed("Failed to get device pointer: "+std::string(cudaGetErrorString(err)), Here());
+                }
+            }
+            else {
+                cudaError_t err = cudaMalloc((void**)&device_data_, bytes);
+                if (err != cudaSuccess) {
+                    throw_AssertionFailed("allocateDevice("+std::string(name())+") Failed to allocate GPU memory: " + std::string(cudaGetErrorString(err)), Here());
+                }
             }
             if (atlas::Library::instance().traceMemory()) {
                 Log::trace() << "allocateDevice(" << name() << ") : cudaMalloc( device_ptr:" << device_data_ << " , " << eckit::Bytes(bytes) << " )" << std::endl;
@@ -202,11 +236,21 @@ class DataStore : public ArrayDataStore {
             if (atlas::Library::instance().traceMemory()) {
                 Log::trace() << "deallocateDevice(" << name() << ") : cudaFree( device_ptr:" << device_data_ << " , " << eckit::Bytes(bytes) << " )" << std::endl;
             }
-            cudaError_t err = cudaFree(device_data_);
-            if (err != cudaSuccess) {
-                throw_AssertionFailed("Failed to deallocateDevice("+std::string(name())+") : " + std::string(cudaGetErrorString(err)), Here());
-            }
+
             accUnmap();
+
+            if (! host_memory_mapped_) {
+                cudaError_t err = cudaFree(device_data_);
+                if (err != cudaSuccess) {
+                    throw_AssertionFailed("Failed to deallocateDevice("+std::string(name())+") : " + std::string(cudaGetErrorString(err)), Here());
+                }
+            }
+            if (host_memory_pinned_) {
+                cudaError_t err = cudaHostUnregister(host_data_);
+                if (err != cudaSuccess) {
+                    throw_AssertionFailed("Failed to make memory pageable again: " + std::string(cudaGetErrorString(err)), Here());
+                }
+            }
             device_data_ = nullptr;
             device_allocated_ = false;
         }
@@ -310,6 +354,8 @@ class DataStore : public ArrayDataStore {
     size_t size_;
     Value* host_data_;
     mutable Value* device_data_{nullptr};
+    bool host_memory_pinned_ = false;
+    bool host_memory_mapped_ = false;
 
     mutable bool host_updated_{true};
     mutable bool device_updated_{true};
@@ -323,9 +369,17 @@ class DataStore : public ArrayDataStore {
 template <typename Value>
 class WrappedDataStore : public ArrayDataStore {
 public:
-    WrappedDataStore(Value* host_data, size_t size): host_data_(host_data), size_(size) {
+    WrappedDataStore(Value* host_data, size_t size, const eckit::Parametrisation& param): 
+        host_data_(host_data), size_(size) {
 #if ATLAS_HAVE_CUDA
         device_updated_ = false;
+        std::string param_in;
+        if (! param.get("host_memory_pinned", param_in)) {
+            host_memory_pinned_ = (param_in == "true");
+        }
+        if (! param.get("host_memory_pinned", param_in)) {
+            host_memory_mapped_ = (param_in == "true");
+        }
 #else
         device_data_ = host_data_;
 #endif
@@ -347,9 +401,11 @@ class WrappedDataStore : public ArrayDataStore {
     void updateHost() const override {
 #if ATLAS_HAVE_CUDA
         if (device_allocated_) {
-            cudaError_t err = cudaMemcpy(host_data_, device_data_, size_*sizeof(Value), cudaMemcpyDeviceToHost);
-            if (err != cudaSuccess) {
-                throw_AssertionFailed("Failed to updateHost: "+std::string(cudaGetErrorString(err)), Here());
+            if (not host_memory_pinned_) {
+                cudaError_t err = cudaMemcpy(host_data_, device_data_, size_*sizeof(Value), cudaMemcpyDeviceToHost);
+                if (err != cudaSuccess) {
+                    throw_AssertionFailed("Failed to updateHost: "+std::string(cudaGetErrorString(err)), Here());
+                }
             }
             host_updated_ = true;
         }
@@ -454,6 +510,8 @@ class WrappedDataStore : public ArrayDataStore {
     size_t size_;
     Value* host_data_;
     mutable Value* device_data_;
+    bool host_memory_pinned_ = false;
+    bool host_memory_mapped_ = false;
 
     mutable bool host_updated_{true};
     mutable bool device_updated_{true};
diff --git a/src/atlas/field/Field.cc b/src/atlas/field/Field.cc
index 905db5634..95d754c43 100644
--- a/src/atlas/field/Field.cc
+++ b/src/atlas/field/Field.cc
@@ -32,21 +32,26 @@ std::ostream& operator<<(std::ostream& os, const Field& f) {
 
 Field::Field(const eckit::Parametrisation& config): Handle(Implementation::create(config)) {}
 
-Field::Field(const std::string& name, array::DataType datatype, const array::ArrayShape& shape):
-    Handle(Implementation::create(name, datatype, shape)) {}
+Field::Field(const std::string& name, array::DataType datatype, const array::ArrayShape& shape,
+        const eckit::Parametrisation& param):
+    Handle(Implementation::create(name, datatype, shape, param)) {}
 
-Field::Field(const std::string& name, array::DataType datatype, array::ArraySpec&& spec):
-    Handle(Implementation::create(name, datatype, std::move(spec))) {}
+Field::Field(const std::string& name, array::DataType datatype, array::ArraySpec&& spec,
+        const eckit::Parametrisation& param):
+    Handle(Implementation::create(name, datatype, std::move(spec), param)) {}
 
-Field::Field(const std::string& name, array::Array* array): Handle(Implementation::create(name, array)) {}
+Field::Field(const std::string& name, array::Array* array, const eckit::Parametrisation& param):
+    Handle(Implementation::create(name, array, param)) {}
 
 template <typename DATATYPE>
-Field::Field(const std::string& name, DATATYPE* data, const array::ArraySpec& spec):
-    Handle(Implementation::wrap(name, data, spec)) {}
+Field::Field(const std::string& name, DATATYPE* data, const array::ArraySpec& spec,
+        const eckit::Parametrisation& param):
+    Handle(Implementation::wrap(name, data, spec, param)) {}
 
 template <typename DATATYPE>
-Field::Field(const std::string& name, DATATYPE* data, const array::ArrayShape& shape):
-    Handle(Implementation::wrap(name, data, shape)) {}
+Field::Field(const std::string& name, DATATYPE* data, const array::ArrayShape& shape,
+        const eckit::Parametrisation& param):
+    Handle(Implementation::wrap(name, data, shape, param)) {}
 
 /// @brief Implicit conversion to Array
 Field::operator const array::Array&() const {
@@ -268,14 +273,22 @@ void Field::reactivateHostWriteViews() const {
 
 // ------------------------------------------------------------------
 
-template Field::Field(const std::string&, float*, const array::ArraySpec&);
-template Field::Field(const std::string&, float*, const array::ArrayShape&);
-template Field::Field(const std::string&, double*, const array::ArraySpec&);
-template Field::Field(const std::string&, double*, const array::ArrayShape&);
-template Field::Field(const std::string&, long*, const array::ArraySpec&);
-template Field::Field(const std::string&, long*, const array::ArrayShape&);
-template Field::Field(const std::string&, int*, const array::ArraySpec&);
-template Field::Field(const std::string&, int*, const array::ArrayShape&);
+template Field::Field(const std::string&, float*, const array::ArraySpec&,
+        const eckit::Parametrisation& param);
+template Field::Field(const std::string&, float*, const array::ArrayShape&,
+        const eckit::Parametrisation& param);
+template Field::Field(const std::string&, double*, const array::ArraySpec&,
+        const eckit::Parametrisation& param);
+template Field::Field(const std::string&, double*, const array::ArrayShape&,
+        const eckit::Parametrisation& param);
+template Field::Field(const std::string&, long*, const array::ArraySpec&,
+        const eckit::Parametrisation& param);
+template Field::Field(const std::string&, long*, const array::ArrayShape&,
+        const eckit::Parametrisation& param);
+template Field::Field(const std::string&, int*, const array::ArraySpec&,
+        const eckit::Parametrisation& param);
+template Field::Field(const std::string&, int*, const array::ArrayShape&,
+        const eckit::Parametrisation& param);
 
 // ------------------------------------------------------------------
 
diff --git a/src/atlas/field/Field.h b/src/atlas/field/Field.h
index 8bac234f0..9717b1694 100644
--- a/src/atlas/field/Field.h
+++ b/src/atlas/field/Field.h
@@ -69,23 +69,28 @@ class Field : DOXYGEN_HIDE(public util::ObjectHandle<field::FieldImpl>) {
     Field(const eckit::Parametrisation&);
 
     /// @brief Create field with given name, Datatype and ArrayShape
-    Field(const std::string& name, array::DataType, const array::ArrayShape& = array::ArrayShape());
+    Field(const std::string& name, array::DataType, const array::ArrayShape& = array::ArrayShape(),
+            const eckit::Parametrisation& = util::Config());
 
     /// @brief Create field with given name, Datatype and ArraySpec
-    Field(const std::string& name, array::DataType, array::ArraySpec&&);
+    Field(const std::string& name, array::DataType, array::ArraySpec&&,
+            const eckit::Parametrisation& = util::Config());
 
     /// @brief Create field with given name, and take ownership of given Array
-    Field(const std::string& name, array::Array*);
+    Field(const std::string& name, array::Array*,
+            const eckit::Parametrisation& = util::Config());
 
     /// @brief Create field by wrapping existing data, Datatype of template and
     /// ArraySpec
     template <typename DATATYPE>
-    Field(const std::string& name, DATATYPE* data, const array::ArraySpec&);
+    Field(const std::string& name, DATATYPE* data, const array::ArraySpec&,
+            const eckit::Parametrisation& =util::Config());
 
     /// @brief Create field by wrapping existing data, Datatype of template and
     /// ArrayShape
     template <typename DATATYPE>
-    Field(const std::string& name, DATATYPE* data, const array::ArrayShape&);
+    Field(const std::string& name, DATATYPE* data, const array::ArrayShape&,
+            const eckit::Parametrisation& = util::Config());
 
     /// @brief Deep copy
     Field clone(const eckit::Parametrisation& = util::Config()) const;
@@ -199,14 +204,22 @@ class Field : DOXYGEN_HIDE(public util::ObjectHandle<field::FieldImpl>) {
     void reactivateHostWriteViews() const;
 };
 
-extern template Field::Field(const std::string&, float*, const array::ArraySpec&);
-extern template Field::Field(const std::string&, float*, const array::ArrayShape&);
-extern template Field::Field(const std::string&, double*, const array::ArraySpec&);
-extern template Field::Field(const std::string&, double*, const array::ArrayShape&);
-extern template Field::Field(const std::string&, long*, const array::ArraySpec&);
-extern template Field::Field(const std::string&, long*, const array::ArrayShape&);
-extern template Field::Field(const std::string&, int*, const array::ArraySpec&);
-extern template Field::Field(const std::string&, int*, const array::ArrayShape&);
+extern template Field::Field(const std::string&, float*, const array::ArraySpec&,
+        const eckit::Parametrisation&);
+extern template Field::Field(const std::string&, float*, const array::ArrayShape&,
+        const eckit::Parametrisation&);
+extern template Field::Field(const std::string&, double*, const array::ArraySpec&,
+        const eckit::Parametrisation&);
+extern template Field::Field(const std::string&, double*, const array::ArrayShape&,
+        const eckit::Parametrisation&);
+extern template Field::Field(const std::string&, long*, const array::ArraySpec&,
+        const eckit::Parametrisation&);
+extern template Field::Field(const std::string&, long*, const array::ArrayShape&,
+        const eckit::Parametrisation&);
+extern template Field::Field(const std::string&, int*, const array::ArraySpec&,
+        const eckit::Parametrisation&);
+extern template Field::Field(const std::string&, int*, const array::ArrayShape&,
+        const eckit::Parametrisation&);
 
 //------------------------------------------------------------------------------------------------------
 
diff --git a/src/atlas/field/FieldCreatorArraySpec.cc b/src/atlas/field/FieldCreatorArraySpec.cc
index 3967bca2e..46c76d860 100644
--- a/src/atlas/field/FieldCreatorArraySpec.cc
+++ b/src/atlas/field/FieldCreatorArraySpec.cc
@@ -65,8 +65,20 @@ FieldImpl* FieldCreatorArraySpec::createField(const eckit::Parametrisation& para
     for (size_t i = 0; i < s.size(); ++i) {
         Log::trace() << s[i] << (i < s.size() - 1 ? "," : "");
     }
-    Log::trace() << "]" << std::endl;
-    auto field = FieldImpl::create(name, datatype, array::ArraySpec(std::move(s), array::ArrayAlignment(alignment)));
+    Log::trace() << "]";
+
+#if ATLAS_HAVE_ACC
+    bool pinned_mem = 0;
+    bool mapped_mem = 0;
+    params.get("cuda_pinned", pinned_mem);
+    params.get("cuda_pinned_mapped", mapped_mem);
+    if (pinned_mem) {
+        Log::trace() << " cuda_(mapped, pinned) ->(" << pinned_mem << ", " << mapped_mem << ")";
+    }
+#endif
+    Log::trace() << std::endl;
+
+    auto field = FieldImpl::create(name, datatype, array::ArraySpec(std::move(s), array::ArrayAlignment(alignment)), params);
     field->callbackOnDestruction([field]() { Log::trace() << "Destroy field " << field->name() << std::endl; });
     return field;
 }
diff --git a/src/atlas/field/detail/FieldImpl.cc b/src/atlas/field/detail/FieldImpl.cc
index f4d5ec889..09b9e1534 100644
--- a/src/atlas/field/detail/FieldImpl.cc
+++ b/src/atlas/field/detail/FieldImpl.cc
@@ -44,38 +44,43 @@ FieldImpl* FieldImpl::create(const eckit::Parametrisation& params) {
     }
 }
 
-FieldImpl* FieldImpl::create(const std::string& name, array::DataType datatype, const array::ArrayShape& shape) {
-    return new FieldImpl(name, datatype, shape);
+FieldImpl* FieldImpl::create(const std::string& name, array::DataType datatype, const array::ArrayShape& shape,
+        const eckit::Parametrisation& param) {
+    return new FieldImpl(name, datatype, shape, param);
 }
 
-FieldImpl* FieldImpl::create(const std::string& name, array::DataType datatype, array::ArraySpec&& spec) {
-    return new FieldImpl(name, datatype, std::move(spec));
+FieldImpl* FieldImpl::create(const std::string& name, array::DataType datatype, array::ArraySpec&& spec,
+        const eckit::Parametrisation& param) {
+    return new FieldImpl(name, datatype, std::move(spec), param);
 }
 
-FieldImpl* FieldImpl::create(const std::string& name, array::Array* array) {
-    return new FieldImpl(name, array);
+FieldImpl* FieldImpl::create(const std::string& name, array::Array* array, 
+        const eckit::Parametrisation& param) {
+    return new FieldImpl(name, array, param);
 }
 
 // -------------------------------------------------------------------------
 
-FieldImpl::FieldImpl(const std::string& name, array::DataType datatype, const array::ArrayShape& shape)
+FieldImpl::FieldImpl(const std::string& name, array::DataType datatype, const array::ArrayShape& shape,
+        const eckit::Parametrisation& param)
 #if ATLAS_HAVE_FUNCTIONSPACE
     :functionspace_(new FunctionSpace())
 #endif
 {
-    array_ = array::Array::create(datatype, shape);
+    array_ = array::Array::create(datatype, shape, param);
     array_->attach();
     rename(name);
     set_levels(0);
     set_variables(0);
 }
 
-FieldImpl::FieldImpl(const std::string& name, array::DataType datatype, array::ArraySpec&& spec)
+FieldImpl::FieldImpl(const std::string& name, array::DataType datatype, array::ArraySpec&& spec,
+        const eckit::Parametrisation& param)
 #if ATLAS_HAVE_FUNCTIONSPACE
     :functionspace_(new FunctionSpace())
 #endif
 {
-    array_ = array::Array::create(datatype, std::move(spec));
+    array_ = array::Array::create(datatype, std::move(spec), param);
     array_->attach();
     rename(name);
     set_levels(0);
@@ -83,11 +88,12 @@ FieldImpl::FieldImpl(const std::string& name, array::DataType datatype, array::A
 }
 
 
-FieldImpl::FieldImpl(const std::string& name, array::Array* array)
+FieldImpl::FieldImpl(const std::string& name, array::Array* array, const eckit::Parametrisation& param)
 #if ATLAS_HAVE_FUNCTIONSPACE
     :functionspace_(new FunctionSpace())
 #endif
 {
+    // TODO: pass param to Field
     array_ = array;
     array_->attach();
     rename(name);
diff --git a/src/atlas/field/detail/FieldImpl.h b/src/atlas/field/detail/FieldImpl.h
index 066f2cba7..7e9fdea5f 100644
--- a/src/atlas/field/detail/FieldImpl.h
+++ b/src/atlas/field/detail/FieldImpl.h
@@ -47,37 +47,45 @@ class FieldImpl : public util::Object {
     static FieldImpl* create(const eckit::Parametrisation&);
 
     /// @brief Create field with given name, Datatype and ArrayShape
-    static FieldImpl* create(const std::string& name, array::DataType, const array::ArrayShape& = array::ArrayShape());
+    static FieldImpl* create(const std::string& name, array::DataType, const array::ArrayShape& = 
+            array::ArrayShape(), const eckit::Parametrisation& param = util::Config());
 
     /// @brief Create field with given name, Datatype and ArrayShape
-    static FieldImpl* create(const std::string& name, array::DataType, array::ArraySpec&&);
+    static FieldImpl* create(const std::string& name, array::DataType, array::ArraySpec&&,
+            const eckit::Parametrisation& param = util::Config());
 
     /// @brief Create field with given name, Datatype of template and ArrayShape
     template <typename DATATYPE>
-    static FieldImpl* create(const std::string& name, const array::ArrayShape& = array::ArrayShape());
+    static FieldImpl* create(const std::string& name, const array::ArrayShape& = array::ArrayShape(),
+            const eckit::Parametrisation& param = util::Config());
 
     /// @brief Create field with given name, and take ownership of given Array
-    static FieldImpl* create(const std::string& name, array::Array*);
+    static FieldImpl* create(const std::string& name, array::Array*,
+            const eckit::Parametrisation& param = util::Config());
 
     /// @brief Create field by wrapping existing data, Datatype of template and
     /// ArraySpec
     template <typename DATATYPE>
-    static FieldImpl* wrap(const std::string& name, DATATYPE* data, const array::ArraySpec&);
+    static FieldImpl* wrap(const std::string& name, DATATYPE* data, const array::ArraySpec&, 
+            const eckit::Parametrisation& param = util::Config());
 
     /// @brief Create field by wrapping existing data, Datatype of template and
     /// ArrayShape
     template <typename DATATYPE>
-    static FieldImpl* wrap(const std::string& name, DATATYPE* data, const array::ArrayShape&);
+    static FieldImpl* wrap(const std::string& name, DATATYPE* data, const array::ArrayShape&,
+            const eckit::Parametrisation& param = util::Config());
 
 private:  // Private constructors to force use of static create functions
     /// Allocate new Array internally
-    FieldImpl(const std::string& name, array::DataType, const array::ArrayShape&);
+    FieldImpl(const std::string& name, array::DataType, const array::ArrayShape&,
+            const eckit::Parametrisation& param = util::Config());
 
     /// Allocate new Array internally
-    FieldImpl(const std::string& name, array::DataType, array::ArraySpec&&);
+    FieldImpl(const std::string& name, array::DataType, array::ArraySpec&&,
+            const eckit::Parametrisation& param = util::Config());
 
     /// Transfer ownership of Array
-    FieldImpl(const std::string& name, array::Array*);
+    FieldImpl(const std::string& name, array::Array*, const eckit::Parametrisation& param = util::Config());
 
 public:  // Destructor
     virtual ~FieldImpl();
@@ -262,20 +270,23 @@ class FieldObserver {
 //----------------------------------------------------------------------------------------------------------------------
 
 template <typename DATATYPE>
-FieldImpl* FieldImpl::create(const std::string& name, const array::ArrayShape& shape) {
-    return create(name, array::DataType::create<DATATYPE>(), shape);
+FieldImpl* FieldImpl::create(const std::string& name, const array::ArrayShape& shape,
+        const eckit::Parametrisation& param) {
+    return create(name, array::DataType::create<DATATYPE>(), shape, param);
 }
 
 template <typename DATATYPE>
-FieldImpl* FieldImpl::wrap(const std::string& name, DATATYPE* data, const array::ArraySpec& spec) {
-    FieldImpl* wrapped = create(name, array::Array::wrap(data, spec));
+FieldImpl* FieldImpl::wrap(const std::string& name, DATATYPE* data, const array::ArraySpec& spec,
+        const eckit::Parametrisation& param) {
+    FieldImpl* wrapped = create(name, array::Array::wrap(data, spec), param);
     wrapped->set_dirty(false);
     return wrapped;
 }
 
 template <typename DATATYPE>
-FieldImpl* FieldImpl::wrap(const std::string& name, DATATYPE* data, const array::ArrayShape& shape) {
-    FieldImpl* wrapped = create(name, array::Array::wrap(data, shape));
+FieldImpl* FieldImpl::wrap(const std::string& name, DATATYPE* data, const array::ArrayShape& shape,
+        const eckit::Parametrisation& param) {
+    FieldImpl* wrapped = create(name, array::Array::wrap(data, shape, param));
     wrapped->set_dirty(false);
     return wrapped;
 }
diff --git a/src/tests/field/CMakeLists.txt b/src/tests/field/CMakeLists.txt
index 3f3fa3303..c9758e768 100644
--- a/src/tests/field/CMakeLists.txt
+++ b/src/tests/field/CMakeLists.txt
@@ -34,6 +34,17 @@ ecbuild_add_test( TARGET atlas_test_field_acc
 if( TEST atlas_test_field_acc )
     set_tests_properties ( atlas_test_field_acc PROPERTIES LABELS "gpu;acc")
 endif()
+ecbuild_add_test( TARGET atlas_test_field_pinning
+  SOURCES  test_field_pinning.cc
+  LIBS     atlas OpenACC::OpenACC_CXX
+  ENVIRONMENT ${ATLAS_TEST_ENVIRONMENT}
+  CONDITION atlas_HAVE_ACC
+)
+if( TEST atlas_test_field_pinning )
+    target_compile_options( atlas_test_field_pinning PRIVATE "${ACC_C_FLAGS}")
+    target_link_options( atlas_test_field_pinning PRIVATE "${ACC_C_FLAGS}")
+    set_tests_properties ( atlas_test_field_pinning PROPERTIES LABELS "gpu;acc")
+endif()
 
 if( HAVE_FCTEST )