Update tests to use local_accessor

aidan.belton · aidan.belton · commit a12a084acf43 · 2022-06-21T15:43:39.000+01:00
diff --git a/SYCL/AtomicRef/add.h b/SYCL/AtomicRef/add.h
@@ -30,8 +30,7 @@ void add_fetch_local_test(queue q, size_t N) {
        auto sum = sum_buf.template get_access<access::mode::read_write>(cgh);
        auto out =
            output_buf.template get_access<access::mode::discard_write>(cgh);
-       accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
-                                                                           cgh);
+       local_accessor<T, 1> loc(1, cgh);
 
        cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
          int gid = it.get_global_id(0);
diff --git a/SYCL/AtomicRef/and.h b/SYCL/AtomicRef/and.h
@@ -30,8 +30,7 @@ void and_local_test(queue q) {
        auto cum = cum_buf.template get_access<access::mode::read_write>(cgh);
        auto out =
            output_buf.template get_access<access::mode::discard_write>(cgh);
-       accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
-                                                                           cgh);
+       local_accessor<T, 1> loc(1, cgh);
 
        cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
          int gid = it.get_global_id(0);
diff --git a/SYCL/AtomicRef/atomic_memory_order_acq_rel.cpp b/SYCL/AtomicRef/atomic_memory_order_acq_rel.cpp
@@ -76,8 +76,7 @@ template <memory_order order> void test_acquire_local() {
     q.submit([&](handler &cgh) {
        auto error =
            error_buf.template get_access<access::mode::read_write>(cgh);
-       accessor<int, 1, access::mode::read_write, access::target::local> val(
-           2, cgh);
+       local_accessor<int, 1> val(2, cgh);
        cgh.parallel_for(
            nd_range<1>(global_size, local_size), [=](nd_item<1> it) {
              size_t lid = it.get_local_id(0);
@@ -168,8 +167,7 @@ template <memory_order order> void test_release_local() {
     q.submit([&](handler &cgh) {
        auto error =
            error_buf.template get_access<access::mode::read_write>(cgh);
-       accessor<int, 1, access::mode::read_write, access::target::local> val(
-           2, cgh);
+       local_accessor<int, 1> val(2, cgh);
        cgh.parallel_for(
            nd_range<1>(global_size, local_size), [=](nd_item<1> it) {
              size_t lid = it.get_local_id(0);
diff --git a/SYCL/AtomicRef/atomic_memory_order_seq_cst.cpp b/SYCL/AtomicRef/atomic_memory_order_seq_cst.cpp
@@ -120,8 +120,7 @@ template <memory_order order> void test_local() {
 
   q.submit([&](handler &cgh) {
      auto res = res_buf.template get_access<access::mode::discard_write>(cgh);
-     accessor<int, 1, access::mode::read_write, access::target::local> val(2,
-                                                                           cgh);
+     local_accessor<int, 1> val(2, cgh);
      cgh.parallel_for(nd_range<1>(N_items, N_items), [=](nd_item<1> it) {
        val[0] = 0;
        it.barrier(access::fence_space::local_space);
diff --git a/SYCL/AtomicRef/compare_exchange.h b/SYCL/AtomicRef/compare_exchange.h
@@ -32,8 +32,7 @@ void compare_exchange_local_test(queue q, size_t N) {
                cgh);
        auto out =
            output_buf.template get_access<access::mode::discard_write>(cgh);
-       accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
-                                                                           cgh);
+       local_accessor<T, 1> loc(1, cgh);
 
        cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
          int gid = it.get_global_id(0);
diff --git a/SYCL/AtomicRef/exchange.h b/SYCL/AtomicRef/exchange.h
@@ -30,8 +30,7 @@ void exchange_local_test(queue q, size_t N) {
        auto cum = cum_buf.template get_access<access::mode::read_write>(cgh);
        auto out =
            output_buf.template get_access<access::mode::discard_write>(cgh);
-       accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
-                                                                           cgh);
+       local_accessor<T, 1> loc(1, cgh);
 
        cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
          int gid = it.get_global_id(0);
diff --git a/SYCL/AtomicRef/load.h b/SYCL/AtomicRef/load.h
@@ -31,8 +31,7 @@ void load_local_test(queue q, size_t N) {
        auto ld = load_buf.template get_access<access::mode::read_write>(cgh);
        auto out =
            output_buf.template get_access<access::mode::discard_write>(cgh);
-       accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
-                                                                           cgh);
+       local_accessor<T, 1> loc(1, cgh);
        cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
          int gid = it.get_global_id(0);
          if (gid == 0)
diff --git a/SYCL/AtomicRef/max.h b/SYCL/AtomicRef/max.h
@@ -30,8 +30,7 @@ void max_local_test(queue q, size_t N) {
        auto cum = cum_buf.template get_access<access::mode::read_write>(cgh);
        auto out =
            output_buf.template get_access<access::mode::discard_write>(cgh);
-       accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
-                                                                           cgh);
+       local_accessor<T, 1> loc(1, cgh);
 
        cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
          int gid = it.get_global_id(0);
diff --git a/SYCL/AtomicRef/min.h b/SYCL/AtomicRef/min.h
@@ -30,8 +30,7 @@ void min_local_test(queue q, size_t N) {
        auto cum = cum_buf.template get_access<access::mode::read_write>(cgh);
        auto out =
            output_buf.template get_access<access::mode::discard_write>(cgh);
-       accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
-                                                                           cgh);
+       local_accessor<T, 1> loc(1, cgh);
 
        cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
          int gid = it.get_global_id(0);
diff --git a/SYCL/AtomicRef/or.h b/SYCL/AtomicRef/or.h
@@ -30,8 +30,7 @@ void or_local_test(queue q) {
        auto cum = cum_buf.template get_access<access::mode::read_write>(cgh);
        auto out =
            output_buf.template get_access<access::mode::discard_write>(cgh);
-       accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
-                                                                           cgh);
+       local_accessor<T, 1> loc(1, cgh);
 
        cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
          int gid = it.get_global_id(0);
diff --git a/SYCL/AtomicRef/store.h b/SYCL/AtomicRef/store.h
@@ -51,8 +51,7 @@ void store_local_test(queue q, size_t N) {
     buffer<T> store_buf(&store, 1);
     q.submit([&](handler &cgh) {
       auto st = store_buf.template get_access<access::mode::read_write>(cgh);
-      accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
-                                                                          cgh);
+      local_accessor<T, 1> loc(1, cgh);
       cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
         size_t gid = it.get_global_id(0);
         auto atm = AtomicRef<T, memory_order::relaxed, scope, space>(loc[0]);
diff --git a/SYCL/AtomicRef/sub.h b/SYCL/AtomicRef/sub.h
@@ -30,8 +30,7 @@ void sub_fetch_local_test(queue q, size_t N) {
        auto sum = sum_buf.template get_access<access::mode::read_write>(cgh);
        auto out =
            output_buf.template get_access<access::mode::discard_write>(cgh);
-       accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
-                                                                           cgh);
+       local_accessor<T, 1> loc(1, cgh);
 
        cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
          int gid = it.get_global_id(0);
diff --git a/SYCL/AtomicRef/xor.h b/SYCL/AtomicRef/xor.h
@@ -30,8 +30,7 @@ void xor_local_test(queue q) {
        auto cum = cum_buf.template get_access<access::mode::read_write>(cgh);
        auto out =
            output_buf.template get_access<access::mode::discard_write>(cgh);
-       accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
-                                                                           cgh);
+       local_accessor<T, 1> loc(1, cgh);
 
        cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
          int gid = it.get_global_id(0);
diff --git a/SYCL/Basic/device_event.cpp b/SYCL/Basic/device_event.cpp
@@ -76,9 +76,7 @@ int test_strideN(size_t stride) {
 
     myQueue.submit([&](handler &cgh) {
       auto out_ptr = out_buf.get_access<access::mode::write>(cgh);
-      accessor<cl::sycl::cl_int, 1, access::mode::read_write,
-               access::target::local>
-          local_acc(range<1>(16), cgh);
+      local_accessor<cl::sycl::cl_int, 1> local_acc(range<1>(16), cgh);
 
       // Create work-groups with 16 work items in each group.
       auto myRange = nd_range<1>(range<1>(nElems), range<1>(workGroupSize));
diff --git a/SYCL/Basic/group_async_copy.cpp b/SYCL/Basic/group_async_copy.cpp
@@ -112,8 +112,7 @@ template <typename T> int test(size_t Stride) {
   Q.submit([&](handler &CGH) {
      auto In = InBuf.template get_access<access::mode::read>(CGH);
      auto Out = OutBuf.template get_access<access::mode::write>(CGH);
-     accessor<T, 1, access::mode::read_write, access::target::local> Local(
-         range<1>{WorkGroupSize}, CGH);
+     local_accessor<T, 1> Local(range<1>{WorkGroupSize}, CGH);
 
      nd_range<1> NDR{range<1>(NElems), range<1>(WorkGroupSize)};
      CGH.parallel_for<KernelName<T>>(NDR, [=](nd_item<1> NDId) {
diff --git a/SYCL/Basic/multi_ptr.cpp b/SYCL/Basic/multi_ptr.cpp
@@ -71,8 +71,7 @@ template <typename T> void testMultPtr() {
       accessor<T, 1, access::mode::read_write, access::target::device,
                access::placeholder::false_t>
           accessorData_2(bufferData_2, cgh);
-      accessor<T, 1, access::mode::read_write, access::target::local>
-          localAccessor(numOfItems, cgh);
+      local_accessor<T, 1> localAccessor(numOfItems, cgh);
 
       cgh.parallel_for<class testMultPtrKernel<T>>(range<1>{10}, [=](id<1> wiID) {
         auto ptr_1 = make_ptr<T, access::address_space::global_space>(
@@ -136,9 +135,7 @@ template <typename T> void testMultPtrArrowOperator() {
       accessor<point<T>, 1, access::mode::read, access::target::constant_buffer,
                access::placeholder::false_t>
           accessorData_2(bufferData_2, cgh);
-      accessor<point<T>, 1, access::mode::read_write, access::target::local,
-               access::placeholder::false_t>
-          accessorData_3(1, cgh);
+      local_accessor<point<T>, 1> accessorData_3(1, cgh);
       accessor<point<T>, 1, access::mode::read, access::target::device,
                access::placeholder::false_t>
           accessorData_4(bufferData_4, cgh);
diff --git a/SYCL/DeviceLib/ITTAnnotations/barrier.cpp b/SYCL/DeviceLib/ITTAnnotations/barrier.cpp
@@ -24,8 +24,7 @@ int main() {
     // ITT start/finish annotations and ITT wg_barrier/wi_resume annotations.
     q.submit([&](handler &cgh) {
       auto acc = buf.get_access<access::mode::read_write>(cgh);
-      accessor<int, 1, access::mode::read_write, access::target::local>
-          local_acc(local_range, cgh);
+      local_accessor<int, 1> local_acc(local_range, cgh);
       cgh.parallel_for<class simple_barrier_kernel>(
           nd_range<1>(num_items, local_range), [=](nd_item<1> item) {
             size_t idx = item.get_global_linear_id();
diff --git a/SYCL/DeviceLib/string_test.cpp b/SYCL/DeviceLib/string_test.cpp
@@ -390,10 +390,7 @@ bool kernel_test_memcpy_addr_space(sycl::queue &deviceQueue) {
                      sycl::access::placeholder::false_t>
           src_acc(buffer1, cgh);
 
-      sycl::accessor<char, 1, sycl::access::mode::read_write,
-                     sycl::access::target::local,
-                     sycl::access::placeholder::false_t>
-          local_acc(sycl::range<1>(16), cgh);
+      sycl::local_accessor<char, 1> local_acc(sycl::range<1>(16), cgh);
 
       sycl::accessor<char, 1, sycl::access::mode::write,
                      sycl::access::target::device,
diff --git a/SYCL/DiscardEvents/discard_events_accessors.cpp b/SYCL/DiscardEvents/discard_events_accessors.cpp
@@ -64,10 +64,7 @@ int main(int Argc, const char *Argv[]) {
   RunKernelHelper(Q, [&](int *Harray) {
     Q.submit([&](sycl::handler &CGH) {
       const size_t LocalMemSize = BUFFER_SIZE;
-      using LocalAccessor =
-          sycl::accessor<int, 1, sycl::access::mode::read_write,
-                         sycl::access::target::local>;
-      LocalAccessor LocalAcc(LocalMemSize, CGH);
+      sycl::local_accessor<int, 1> LocalAcc(LocalMemSize, CGH);
 
       CGH.parallel_for<class kernel_using_local_memory>(
           Range, [=](sycl::item<1> itemID) {
diff --git a/SYCL/GroupAlgorithm/SYCL2020/sort.cpp b/SYCL/GroupAlgorithm/SYCL2020/sort.cpp
@@ -112,9 +112,7 @@ int test_sort_over_group(sycl::queue &q, std::size_t local,
               << std::endl;
   q.submit([&](sycl::handler &h) {
      auto aI1 = sycl::accessor(bufI1, h);
-     sycl::accessor<std::byte, 1, sycl::access_mode::read_write,
-                    sycl::access::target::local>
-         scratch({local_memory_size}, h);
+     sycl::local_accessor<std::byte, 1> scratch({local_memory_size}, h);
 
      h.parallel_for<sort_over_group_kernel_name<int_wrapper<dim>, T, Compare>>(
          sycl::nd_range<dim>(local_range, local_range),
@@ -167,9 +165,7 @@ int test_joint_sort(sycl::queue &q, std::size_t n_items, std::size_t local,
               << std::endl;
   q.submit([&](sycl::handler &h) {
      auto aI1 = sycl::accessor(bufI1, h);
-     sycl::accessor<std::byte, 1, sycl::access_mode::read_write,
-                    sycl::access::target::local>
-         scratch({local_memory_size}, h);
+     sycl::local_accessor<std::byte, 1> scratch({local_memory_size}, h);
 
      h.parallel_for<joint_sort_kernel_name<T, Compare>>(
          sycl::nd_range<1>{{n_groups * local}, {local}},
diff --git a/SYCL/GroupAlgorithm/barrier.cpp b/SYCL/GroupAlgorithm/barrier.cpp
@@ -22,10 +22,8 @@ void basic() {
 
     q.submit([&](handler &cgh) {
       auto acc = buf.get_access<access::mode::read_write>(cgh);
-      accessor<int, 1, access::mode::read_write, access::target::local> loc(
-          N, cgh);
-      accessor<barrier, 1, access::mode::read_write, access::target::local>
-          loc_barrier(2, cgh);
+      local_accessor<int, 1> loc(N, cgh);
+      local_accessor<barrier, 1> loc_barrier(2, cgh);
       cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> item) {
         size_t idx = item.get_local_linear_id();
         loc[idx] = acc[idx];
@@ -69,10 +67,8 @@ void interface() {
       auto data_acc = data_buf.get_access<access::mode::read_write>(cgh);
       auto test1_acc = test1_buf.get_access<access::mode::read_write>(cgh);
       auto test2_acc = test2_buf.get_access<access::mode::read_write>(cgh);
-      accessor<int, 1, access::mode::read_write, access::target::local> loc(
-          N, cgh);
-      accessor<barrier, 1, access::mode::read_write, access::target::local>
-          loc_barrier(2, cgh);
+      local_accessor<int, 1> loc(N, cgh);
+      local_accessor<barrier, 1> loc_barrier(2, cgh);
       cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> item) {
         size_t idx = item.get_local_linear_id();
         if (idx == 0) {
diff --git a/SYCL/Regression/group.cpp b/SYCL/Regression/group.cpp
@@ -187,9 +187,7 @@ bool group__async_work_group_copy() {
 
       Q.submit([&](handler &cgh) {
         auto AccGlobal = Buf.get_access<access::mode::read_write>(cgh);
-        accessor<DataType, DIMS, access::mode::read_write,
-                 access::target::local>
-            AccLocal(LocalRange, cgh);
+        local_accessor<DataType, DIMS> AccLocal(LocalRange, cgh);
 
         cgh.parallel_for<class group__async_work_group_copy>(
             nd_range<2>{GlobalRange, LocalRange}, [=](nd_item<DIMS> I) {
diff --git a/SYCL/Regression/local-arg-align.cpp b/SYCL/Regression/local-arg-align.cpp
@@ -26,10 +26,8 @@ int main(int argc, char *argv[]) {
 
   q.submit([&](sycl::handler &h) {
      // Use two local buffers, one with an int and one with a double4
-     accessor<cl_int, 1, access::mode::read_write, access::target::local> a(1,
-                                                                            h);
-     accessor<double4, 1, access::mode::read_write, access::target::local> b(1,
-                                                                             h);
+     local_accessor<cl_int, 1> a(1, h);
+     local_accessor<double4, 1> b(1, h);
 
      auto ares = res.get_access<access::mode::read_write>(h);
 
diff --git a/SYCL/Regression/zero_size_local_accessor.cpp b/SYCL/Regression/zero_size_local_accessor.cpp
@@ -17,9 +17,7 @@
 int main() {
   sycl::queue Q;
   Q.submit([&](sycl::handler &CGH) {
-     sycl::accessor<uint8_t, 1, sycl::access_mode::read_write,
-                    sycl::access::target::local>
-         ZeroSizeLocalAcc(sycl::range<1>(0), CGH);
+     sycl::local_accessor<uint8_t, 1> ZeroSizeLocalAcc(sycl::range<1>(0), CGH);
      CGH.single_task([=]() {
        if (ZeroSizeLocalAcc.get_range()[0])
          ZeroSizeLocalAcc[0] = 1;
diff --git a/SYCL/SubGroup/load_store.cpp b/SYCL/SubGroup/load_store.cpp
@@ -44,8 +44,7 @@ template <typename T, int N> void check(queue &Queue) {
     Queue.submit([&](handler &cgh) {
       auto acc = syclbuf.template get_access<access::mode::read_write>(cgh);
       auto sgsizeacc = sgsizebuf.get_access<access::mode::read_write>(cgh);
-      accessor<T, 1, access::mode::read_write, access::target::local> LocalMem(
-          {L + max_sg_size * N}, cgh);
+      local_accessor<T, 1> LocalMem({L + max_sg_size * N}, cgh);
       cgh.parallel_for<sycl_subgr<T, N>>(NdRange, [=](nd_item<1> NdItem) {
         ext::oneapi::sub_group SG = NdItem.get_sub_group();
         auto SGid = SG.get_group_id().get(0);
@@ -132,8 +131,7 @@ template <typename T> void check(queue &Queue) {
     Queue.submit([&](handler &cgh) {
       auto acc = syclbuf.template get_access<access::mode::read_write>(cgh);
       auto sgsizeacc = sgsizebuf.get_access<access::mode::read_write>(cgh);
-      accessor<T, 1, access::mode::read_write, access::target::local> LocalMem(
-          {L}, cgh);
+      local_accessor<T, 1> LocalMem({L}, cgh);
       cgh.parallel_for<sycl_subgr<T, 0>>(NdRange, [=](nd_item<1> NdItem) {
         ext::oneapi::sub_group SG = NdItem.get_sub_group();
         if (NdItem.get_global_id(0) == 0)
diff --git a/SYCL/SubGroup/sub_group_as.cpp b/SYCL/SubGroup/sub_group_as.cpp
@@ -36,9 +36,7 @@ int main(int argc, char *argv[]) {
     queue.submit([&](cl::sycl::handler &cgh) {
       auto global = buf.get_access<cl::sycl::access::mode::read_write,
                                    cl::sycl::access::target::device>(cgh);
-      sycl::accessor<int, 1, sycl::access::mode::read_write,
-                     sycl::access::target::local>
-          local(N, cgh);
+      sycl::local_accessor<int, 1> local(N, cgh);
 
       cgh.parallel_for<class test>(
           cl::sycl::nd_range<1>(N, 32), [=](cl::sycl::nd_item<1> it) {
diff --git a/SYCL/SubGroup/sub_group_as_vec.cpp b/SYCL/SubGroup/sub_group_as_vec.cpp
@@ -38,9 +38,7 @@ int main(int argc, char *argv[]) {
     queue.submit([&](cl::sycl::handler &cgh) {
       auto global = buf.get_access<cl::sycl::access::mode::read_write,
                                    cl::sycl::access::target::device>(cgh);
-      sycl::accessor<sycl::vec<int, 2>, 1, sycl::access::mode::read_write,
-                     sycl::access::target::local>
-          local(N, cgh);
+      sycl::local_accessor<sycl::vec<int, 2>, 1> local(N, cgh);
       cgh.parallel_for<class test>(
           cl::sycl::nd_range<1>(N, 32), [=](cl::sycl::nd_item<1> it) {
             cl::sycl::ext::oneapi::sub_group sg = it.get_sub_group();
diff --git a/SYCL/XPTI/buffer/accessors.cpp b/SYCL/XPTI/buffer/accessors.cpp
@@ -28,7 +28,7 @@ int main() {
     // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID2:.*]]|2014|1025|{{.*}}accessors.cpp:[[# @LINE + 1]]:15
     auto A2 = Buf.get_access<mode::write>(cgh);
     // CHECK: {{[0-9]+}}|Construct accessor|0x0|[[ACCID3:.*]]|2016|1026|{{.*}}accessors.cpp:[[# @LINE + 1]]:61
-    sycl::accessor<int, 1, mode::read_write, target::local> A3(Range, cgh);
+    sycl::local_accessor<int, 1> A3(Range, cgh);
     // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID4:.*]]|2014|1027|{{.*}}accessors.cpp:[[# @LINE + 1]]:15
     auto A4 = Buf.get_access<mode::discard_write>(cgh);
     // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID5:.*]]|2014|1028|{{.*}}accessors.cpp:[[# @LINE + 1]]:15
diff --git a/SYCL/XPTI/kernel/basic.cpp b/SYCL/XPTI/kernel/basic.cpp
@@ -57,7 +57,7 @@ int main() {
         // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID1:.+]]|2014|1026|{{.*}}.cpp:[[# @LINE + 1]]:19
         auto A1 = Buf.get_access<mode::read_write>(cgh);
         // CHECK: {{[0-9]+}}|Construct accessor|0x0|[[ACCID2:.*]]|2016|1026|{{.*}}.cpp:[[# @LINE + 1]]:65
-        sycl::accessor<int, 1, mode::read_write, target::local> A2(Range, cgh);
+        sycl::local_accessor<int, 1> A2(Range, cgh);
         // CHECK-OPT:Node create|{{.*}}FillBuffer{{.*}}|{{.*}}.cpp:[[# @LINE - 6 ]]:3|{5, 1, 1}, {0, 0, 0}, {0, 0, 0}, 6
         // CHECK-NOOPT:Node create|{{.*}}FillBuffer{{.*}}|{{.*}}.cpp:[[# @LINE - 7 ]]:3|{5, 1, 1}, {0, 0, 0}, {0, 0, 0}, 12
         cgh.parallel_for<class FillBuffer>(
