[libcu++] Uncomment some tests and fix launch include after launch was moved to libcu++ (#6966)

* Uncomment some tests and fix launch include after launch was moved to libcu++

* Try to make GCC 7 happy

Author: pciolkosz

libcudacxx/test/libcudacxx/cuda/ccclrt/hierarchy/hierarchy_smoke.cu

@@ -8,6 +8,8 @@
 //
 //===----------------------------------------------------------------------===//
 
+#include <cuda/launch>
+
 #include <iostream>
 
 #include <cooperative_groups.h>
@@ -566,16 +568,12 @@ C2H_TEST("Trivially constructable", "[hierarchy]")
 
 C2H_TEST("cuda::distribute", "[hierarchy]")
 {
-  /*
   int numElements               = 50000;
   constexpr int threadsPerBlock = 256;
-  auto config                   =
-  cuda::distribute<threadsPerBlock>(numElements);
+  auto config                   = cuda::distribute<threadsPerBlock>(numElements);
 
   CCCLRT_REQUIRE(config.dims.count(cuda::thread, cuda::block) == 256);
-  CCCLRT_REQUIRE(config.dims.count(cuda::block, cuda::grid) == (numElements +
-  threadsPerBlock - 1) / threadsPerBlock);
-  */
+  CCCLRT_REQUIRE(config.dims.count(cuda::block, cuda::grid) == (numElements + threadsPerBlock - 1) / threadsPerBlock);
 }
 
 C2H_TEST("hierarchy merge", "[hierarchy]")

libcudacxx/test/libcudacxx/cuda/containers/buffer/transform.cu

@@ -15,13 +15,12 @@
 
 #include <cuda/algorithm>
 #include <cuda/buffer>
+#include <cuda/launch>
 #include <cuda/memory_resource>
 #include <cuda/std/tuple>
 #include <cuda/std/type_traits>
 #include <cuda/std/utility>
 
-#include <cuda/experimental/launch.cuh>
-
 #include <algorithm>
 
 #include "helper.h"

libcudacxx/test/libcudacxx/cuda/memory_resource/resources/device_memory_resource.cu

@@ -8,6 +8,7 @@
 //
 //===----------------------------------------------------------------------===//
 
+#include <cuda/launch>
 #include <cuda/memory_pool>
 #include <cuda/memory_resource>
 #include <cuda/std/cstdint>
@@ -374,22 +375,21 @@ C2H_CCCLRT_TEST("device_memory_resource comparison", "[memory_resource]")
 
 C2H_CCCLRT_TEST("Async memory resource access", "")
 {
-  /* disable until we move the launch API to libcudacxx
   if (cuda::devices.size() > 1)
   {
     auto peers = cuda::devices[0].peers();
     if (peers.size() > 0)
     {
       cuda::device_memory_pool pool{cuda::devices[0]};
-      cuda::stream_ref stream{peers.front()};
+      cuda::stream stream{peers.front()};
       CCCLRT_CHECK(pool.is_accessible_from(cuda::devices[0]));
 
       auto allocate_and_check_access = [&](auto& resource) {
         auto* ptr1  = resource.allocate(stream, sizeof(int));
         auto* ptr2  = resource.allocate_sync(sizeof(int));
-        auto config = cudax::distribute<1>(1);
-        cudax::launch(stream, config, test::assign_42{}, (int*) ptr1);
-        cudax::launch(stream, config, test::assign_42{}, (int*) ptr2);
+        auto config = cuda::distribute<1>(1);
+        cuda::launch(stream, config, test::assign_42{}, (int*) ptr1);
+        cuda::launch(stream, config, test::assign_42{}, (int*) ptr2);
         stream.sync();
         resource.deallocate(stream, ptr1, sizeof(int));
         resource.deallocate_sync(ptr2, sizeof(int));
@@ -400,7 +400,7 @@ C2H_CCCLRT_TEST("Async memory resource access", "")
       CCCLRT_CHECK(pool.is_accessible_from(peers.front()));
       allocate_and_check_access(pool);
 
-      cudax::device_memory_pool_ref resource{pool};
+      cuda::device_memory_pool_ref resource{pool};
       CCCLRT_CHECK(resource.is_accessible_from(peers.front()));
       allocate_and_check_access(resource);
 
@@ -421,7 +421,9 @@ C2H_CCCLRT_TEST("Async memory resource access", "")
 
       // Check if enable can include the device on which the pool resides
       {
-        std::vector peers_ext(peers.begin(), peers.end());
+        // Separate insert call because GCC 7 doesn't like the constructor from iterators
+        std::vector<cuda::device_ref> peers_ext;
+        peers_ext.insert(peers_ext.end(), peers.begin(), peers.end());
         peers_ext.push_back(cuda::devices[0]);
         pool.enable_access_from(peers_ext);
 
@@ -439,5 +441,4 @@ C2H_CCCLRT_TEST("Async memory resource access", "")
       }
     }
   }
-  */
 }