[Thrust] Perform asynchronous allocations by default for the `par_nosync` policy #4204

brycelelbach · 2025-03-19T18:35:03Z

This will make algorithms (like scans) that don't have a computation-dependent result but do temporary allocation properly asynchronous under par_nosync.

This is implemented by adding overloads of (get|return)_temporary_buffer for par_nosync and it's corresponding stream-attached form. The overloads just use cuda(Malloc|Free)Async directly for now. It falls back to cuda(Malloc|Free) in the case that async allocation fails, because I'm told that async allocation doesn't work on some platforms (notably Windows TCC driver).

When libcu++ memory resources are non-experimental, instead of specializing and using cuda(Malloc|FreeAsync directly, we should check whether the memory resource attached to the policy supports async, and for par_nosync we should default to using the memory resource's async APIs. We'd also have to give par_nosync a default policy.

However, for the time being, this works and allows much broader use of par_nosync.

Implements #4117.

copy-pr-bot · 2025-03-19T18:35:07Z

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

bernhardmgruber · 2025-03-20T01:01:44Z

/ok to test

bernhardmgruber · 2025-03-20T01:02:00Z

pre-commit.ci autofix

…ync` policy. This will make algorithms (like scans) that don't have a computation-dependent result but do temporary allocation properly asynchronous under `par_nosync`.

thrust/thrust/system/cuda/detail/temporary_buffer.h

github-actions · 2025-03-20T02:50:14Z

🟨 CI finished in 1h 31m: Pass: 54%/97 | Total: 1d 21h | Avg: 28m 18s | Max: 1h 15m | Hits: 75%/57652

🟨 thrust: Pass: 4%/45 | Total: 6h 52m | Avg: 9m 10s | Max: 58m 21s | Hits: 77%/3552

🚨 cudacxx_family: nvcc 🚨
  🟩 ClangCUDA          Pass: 100%/2   | Total: 48m 48s | Avg: 24m 24s | Max: 24m 43s | Hits:  77%/3552  
  🔥 nvcc               Pass:   0%/43  | Total:  6h 03m | Avg:  8m 27s | Max: 58m 21s
🟨 cudacxx
  🟩 ClangCUDA18        Pass: 100%/2   | Total: 48m 48s | Avg: 24m 24s | Max: 24m 43s | Hits:  77%/3552  
  🟥 nvcc12.0           Pass:   0%/5   | Total:  1h 08m | Avg: 13m 39s | Max: 49m 45s
  🟥 nvcc12.6           Pass:   0%/2   | Total: 14m 34s | Avg:  7m 17s | Max:  7m 18s
  🟥 nvcc12.8           Pass:   0%/36  | Total:  4h 40m | Avg:  7m 48s | Max: 58m 21s
🟥 cmake_options
  🟥 -DTHRUST_DISPATCH_TYPE=Force32bit Pass:   0%/2   | Total:  3m 43s | Avg:  1m 51s | Max:  3m 43s
🟨 cpu
  🟨 amd64              Pass:   4%/43  | Total:  6h 43m | Avg:  9m 22s | Max: 58m 21s | Hits:  77%/3552  
  🟥 arm64              Pass:   0%/2   | Total:  9m 17s | Avg:  4m 38s | Max:  4m 47s
🟨 ctk
  🟥 12.0               Pass:   0%/5   | Total:  1h 08m | Avg: 13m 39s | Max: 49m 45s
  🟥 12.6               Pass:   0%/2   | Total: 14m 34s | Avg:  7m 17s | Max:  7m 18s
  🟨 12.8               Pass:   5%/38  | Total:  5h 29m | Avg:  8m 40s | Max: 58m 21s | Hits:  77%/3552  
🟨 cxx
  🟥 Clang14            Pass:   0%/4   | Total: 18m 15s | Avg:  4m 33s | Max:  4m 50s
  🟥 Clang15            Pass:   0%/2   | Total:  8m 47s | Avg:  4m 23s | Max:  4m 30s
  🟥 Clang16            Pass:   0%/2   | Total:  9m 16s | Avg:  4m 38s | Max:  4m 46s
  🟥 Clang17            Pass:   0%/2   | Total:  9m 14s | Avg:  4m 37s | Max:  4m 38s
  🟨 Clang18            Pass:  28%/7   | Total:  1h 02m | Avg:  8m 56s | Max: 24m 43s | Hits:  77%/3552  
  🟥 GCC7               Pass:   0%/2   | Total:  9m 20s | Avg:  4m 40s | Max:  4m 47s
  🟥 GCC8               Pass:   0%/1   | Total:  4m 15s | Avg:  4m 15s | Max:  4m 15s
  🟥 GCC9               Pass:   0%/2   | Total:  9m 30s | Avg:  4m 45s | Max:  4m 46s
  🟥 GCC10              Pass:   0%/2   | Total:  8m 41s | Avg:  4m 20s | Max:  4m 21s
  🟥 GCC11              Pass:   0%/2   | Total:  9m 08s | Avg:  4m 34s | Max:  4m 43s
  🟥 GCC12              Pass:   0%/2   | Total:  9m 03s | Avg:  4m 31s | Max:  4m 36s
  🟥 GCC13              Pass:   0%/10  | Total: 25m 13s | Avg:  2m 31s | Max:  4m 47s
  🟥 MSVC14.29          Pass:   0%/2   | Total:  1h 40m | Avg: 50m 28s | Max: 51m 12s
  🟥 MSVC14.42          Pass:   0%/3   | Total:  1h 53m | Avg: 37m 54s | Max: 58m 21s
  🟥 NVHPC25.1          Pass:   0%/2   | Total: 14m 34s | Avg:  7m 17s | Max:  7m 18s
🟨 cxx_family
  🟨 Clang              Pass:  11%/17  | Total:  1h 48m | Avg:  6m 21s | Max: 24m 43s | Hits:  77%/3552  
  🟥 GCC                Pass:   0%/21  | Total:  1h 15m | Avg:  3m 34s | Max:  4m 47s
  🟥 MSVC               Pass:   0%/5   | Total:  3h 34m | Avg: 42m 56s | Max: 58m 21s
  🟥 NVHPC              Pass:   0%/2   | Total: 14m 34s | Avg:  7m 17s | Max:  7m 18s
🟨 gpu
  🟥 h100               Pass:   0%/2   | Total:  3m 15s | Avg:  1m 37s | Max:  3m 15s
  🟨 rtx2080            Pass:   6%/33  | Total:  5h 37m | Avg: 10m 14s | Max: 55m 23s | Hits:  77%/3552  
  🟥 rtx4090            Pass:   0%/10  | Total:  1h 11m | Avg:  7m 08s | Max: 58m 21s
🟨 jobs
  🟨 Build              Pass:   5%/38  | Total:  6h 52m | Avg: 10m 51s | Max: 58m 21s | Hits:  77%/3552  
  🟥 TestCPU            Pass:   0%/3  
  🟥 TestGPU            Pass:   0%/4  
🟥 sm
  🟥 90                 Pass:   0%/2   | Total:  3m 15s | Avg:  1m 37s | Max:  3m 15s
  🟥 90;90a;100         Pass:   0%/1   | Total:  4m 27s | Avg:  4m 27s | Max:  4m 27s
🟨 std
  🟨 17                 Pass:   5%/20  | Total:  4h 17m | Avg: 12m 52s | Max: 55m 23s | Hits:  77%/1776  
  🟨 20                 Pass:   4%/23  | Total:  2h 31m | Avg:  6m 34s | Max: 58m 21s | Hits:  77%/1776

🟥 python: Pass: 0%/1 | Total: 12m 48s | Avg: 12m 48s | Max: 12m 48s

🟥 cpu
  🟥 amd64              Pass:   0%/1   | Total: 12m 48s | Avg: 12m 48s | Max: 12m 48s
🟥 ctk
  🟥 12.8               Pass:   0%/1   | Total: 12m 48s | Avg: 12m 48s | Max: 12m 48s
🟥 cudacxx
  🟥 nvcc12.8           Pass:   0%/1   | Total: 12m 48s | Avg: 12m 48s | Max: 12m 48s
🟥 cudacxx_family
  🟥 nvcc               Pass:   0%/1   | Total: 12m 48s | Avg: 12m 48s | Max: 12m 48s
🟥 cxx
  🟥 GCC13              Pass:   0%/1   | Total: 12m 48s | Avg: 12m 48s | Max: 12m 48s
🟥 cxx_family
  🟥 GCC                Pass:   0%/1   | Total: 12m 48s | Avg: 12m 48s | Max: 12m 48s
🟥 gpu
  🟥 rtx2080            Pass:   0%/1   | Total: 12m 48s | Avg: 12m 48s | Max: 12m 48s
🟥 jobs
  🟥 Test               Pass:   0%/1   | Total: 12m 48s | Avg: 12m 48s | Max: 12m 48s

🟩 cub: Pass: 100%/45 | Total: 1d 14h | Avg: 50m 40s | Max: 1h 15m | Hits: 75%/53780

🟩 cpu
  🟩 amd64              Pass: 100%/43  | Total:  1d 12h | Avg: 50m 25s | Max:  1h 15m | Hits:  75%/51336 
  🟩 arm64              Pass: 100%/2   | Total:  1h 51m | Avg: 55m 54s | Max:  1h 03m | Hits:  69%/2444  
🟩 ctk
  🟩 12.0               Pass: 100%/5   | Total:  4h 48m | Avg: 57m 46s | Max:  1h 04m | Hits:  70%/5940  
  🟩 12.6               Pass: 100%/2   | Total:  2h 25m | Avg:  1h 12m | Max:  1h 15m | Hits:  69%/2260  
  🟩 12.8               Pass: 100%/38  | Total:  1d 06h | Avg: 48m 35s | Max:  1h 05m | Hits:  76%/45580 
🟩 cudacxx
  🟩 ClangCUDA18        Pass: 100%/2   | Total:  2h 01m | Avg:  1h 00m | Max:  1h 01m | Hits:  75%/2108  
  🟩 nvcc12.0           Pass: 100%/5   | Total:  4h 48m | Avg: 57m 46s | Max:  1h 04m | Hits:  70%/5940  
  🟩 nvcc12.6           Pass: 100%/2   | Total:  2h 25m | Avg:  1h 12m | Max:  1h 15m | Hits:  69%/2260  
  🟩 nvcc12.8           Pass: 100%/36  | Total:  1d 04h | Avg: 47m 54s | Max:  1h 05m | Hits:  76%/43472 
🟩 cudacxx_family
  🟩 ClangCUDA          Pass: 100%/2   | Total:  2h 01m | Avg:  1h 00m | Max:  1h 01m | Hits:  75%/2108  
  🟩 nvcc               Pass: 100%/43  | Total:  1d 11h | Avg: 50m 11s | Max:  1h 15m | Hits:  75%/51672 
🟩 cxx
  🟩 Clang14            Pass: 100%/4   | Total:  4h 11m | Avg:  1h 02m | Max:  1h 05m | Hits:  69%/4896  
  🟩 Clang15            Pass: 100%/2   | Total:  2h 02m | Avg:  1h 01m | Max:  1h 02m | Hits:  69%/2444  
  🟩 Clang16            Pass: 100%/2   | Total:  2h 02m | Avg:  1h 01m | Max:  1h 04m | Hits:  69%/2444  
  🟩 Clang17            Pass: 100%/2   | Total:  2h 03m | Avg:  1h 01m | Max:  1h 02m | Hits:  69%/2444  
  🟩 Clang18            Pass: 100%/7   | Total:  5h 58m | Avg: 51m 13s | Max:  1h 03m | Hits:  80%/8218  
  🟩 GCC7               Pass: 100%/2   | Total:  1h 45m | Avg: 52m 58s | Max: 53m 39s | Hits:  69%/2448  
  🟩 GCC8               Pass: 100%/1   | Total: 46m 56s | Avg: 46m 56s | Max: 46m 56s | Hits:  69%/1224  
  🟩 GCC9               Pass: 100%/2   | Total:  1h 42m | Avg: 51m 19s | Max: 51m 30s | Hits:  69%/2448  
  🟩 GCC10              Pass: 100%/2   | Total:  1h 35m | Avg: 47m 30s | Max: 48m 37s | Hits:  69%/2448  
  🟩 GCC11              Pass: 100%/2   | Total:  1h 40m | Avg: 50m 08s | Max: 53m 01s | Hits:  69%/2444  
  🟩 GCC12              Pass: 100%/2   | Total:  1h 47m | Avg: 53m 33s | Max: 55m 23s | Hits:  69%/2444  
  🟩 GCC13              Pass: 100%/11  | Total:  6h 06m | Avg: 33m 16s | Max: 50m 14s | Hits:  85%/13442 
  🟩 MSVC14.29          Pass: 100%/2   | Total:  1h 55m | Avg: 57m 54s | Max: 58m 20s | Hits:  74%/2088  
  🟩 MSVC14.42          Pass: 100%/2   | Total:  1h 56m | Avg: 58m 03s | Max: 58m 41s | Hits:  74%/2088  
  🟩 NVHPC25.1          Pass: 100%/2   | Total:  2h 25m | Avg:  1h 12m | Max:  1h 15m | Hits:  69%/2260  
🟩 cxx_family
  🟩 Clang              Pass: 100%/17  | Total: 16h 19m | Avg: 57m 36s | Max:  1h 05m | Hits:  73%/20446 
  🟩 GCC                Pass: 100%/22  | Total: 15h 23m | Avg: 41m 59s | Max: 55m 23s | Hits:  77%/26898 
  🟩 MSVC               Pass: 100%/4   | Total:  3h 51m | Avg: 57m 58s | Max: 58m 41s | Hits:  74%/4176  
  🟩 NVHPC              Pass: 100%/2   | Total:  2h 25m | Avg:  1h 12m | Max:  1h 15m | Hits:  69%/2260  
🟩 gpu
  🟩 h100               Pass: 100%/3   | Total:  1h 16m | Avg: 25m 32s | Max: 28m 26s | Hits:  89%/3666  
  🟩 rtx2080            Pass: 100%/34  | Total:  1d 08h | Avg: 57m 14s | Max:  1h 15m | Hits:  70%/40338 
  🟩 rtxa6000           Pass: 100%/8   | Total:  4h 17m | Avg: 32m 09s | Max:  1h 03m | Hits:  92%/9776  
🟩 jobs
  🟩 Build              Pass: 100%/37  | Total:  1d 10h | Avg: 56m 22s | Max:  1h 15m | Hits:  70%/44004 
  🟩 DeviceLaunch       Pass: 100%/1   | Total: 27m 02s | Avg: 27m 02s | Max: 27m 02s | Hits:  99%/1222  
  🟩 GraphCapture       Pass: 100%/1   | Total: 20m 37s | Avg: 20m 37s | Max: 20m 37s | Hits:  99%/1222  
  🟩 HostLaunch         Pass: 100%/3   | Total:  1h 20m | Avg: 26m 51s | Max: 28m 26s | Hits:  99%/3666  
  🟩 TestGPU            Pass: 100%/3   | Total:  1h 05m | Avg: 21m 58s | Max: 22m 09s | Hits:  99%/3666  
🟩 sm
  🟩 90                 Pass: 100%/3   | Total:  1h 16m | Avg: 25m 32s | Max: 28m 26s | Hits:  89%/3666  
  🟩 90;90a;100         Pass: 100%/1   | Total: 46m 42s | Avg: 46m 42s | Max: 46m 42s | Hits:  69%/1222  
🟩 std
  🟩 17                 Pass: 100%/20  | Total: 19h 12m | Avg: 57m 36s | Max:  1h 15m | Hits:  70%/23662 
  🟩 20                 Pass: 100%/25  | Total: 18h 48m | Avg: 45m 07s | Max:  1h 10m | Hits:  79%/30118

🟩 stdpar: Pass: 100%/4 | Total: 16m 52s | Avg: 4m 13s | Max: 5m 08s

🟩 cpu
  🟩 amd64              Pass: 100%/2   | Total: 10m 08s | Avg:  5m 04s | Max:  5m 08s
  🟩 arm64              Pass: 100%/2   | Total:  6m 44s | Avg:  3m 22s | Max:  3m 22s
🟩 ctk
  🟩 12.6               Pass: 100%/4   | Total: 16m 52s | Avg:  4m 13s | Max:  5m 08s
🟩 cudacxx
  🟩 nvcc12.6           Pass: 100%/4   | Total: 16m 52s | Avg:  4m 13s | Max:  5m 08s
🟩 cudacxx_family
  🟩 nvcc               Pass: 100%/4   | Total: 16m 52s | Avg:  4m 13s | Max:  5m 08s
🟩 cxx
  🟩 NVHPC25.1          Pass: 100%/4   | Total: 16m 52s | Avg:  4m 13s | Max:  5m 08s
🟩 cxx_family
  🟩 NVHPC              Pass: 100%/4   | Total: 16m 52s | Avg:  4m 13s | Max:  5m 08s
🟩 gpu
  🟩 rtx2080            Pass: 100%/4   | Total: 16m 52s | Avg:  4m 13s | Max:  5m 08s
🟩 jobs
  🟩 Build              Pass: 100%/4   | Total: 16m 52s | Avg:  4m 13s | Max:  5m 08s
🟩 std
  🟩 17                 Pass: 100%/2   | Total:  8m 22s | Avg:  4m 11s | Max:  5m 00s
  🟩 20                 Pass: 100%/2   | Total:  8m 30s | Avg:  4m 15s | Max:  5m 08s

🟩 cccl_c_parallel: Pass: 100%/2 | Total: 24m 16s | Avg: 12m 08s | Max: 21m 58s | Hits: 98%/320

🟩 cpu
  🟩 amd64              Pass: 100%/2   | Total: 24m 16s | Avg: 12m 08s | Max: 21m 58s | Hits:  98%/320   
🟩 ctk
  🟩 12.8               Pass: 100%/2   | Total: 24m 16s | Avg: 12m 08s | Max: 21m 58s | Hits:  98%/320   
🟩 cudacxx
  🟩 nvcc12.8           Pass: 100%/2   | Total: 24m 16s | Avg: 12m 08s | Max: 21m 58s | Hits:  98%/320   
🟩 cudacxx_family
  🟩 nvcc               Pass: 100%/2   | Total: 24m 16s | Avg: 12m 08s | Max: 21m 58s | Hits:  98%/320   
🟩 cxx
  🟩 GCC13              Pass: 100%/2   | Total: 24m 16s | Avg: 12m 08s | Max: 21m 58s | Hits:  98%/320   
🟩 cxx_family
  🟩 GCC                Pass: 100%/2   | Total: 24m 16s | Avg: 12m 08s | Max: 21m 58s | Hits:  98%/320   
🟩 gpu
  🟩 rtx2080            Pass: 100%/2   | Total: 24m 16s | Avg: 12m 08s | Max: 21m 58s | Hits:  98%/320   
🟩 jobs
  🟩 Build              Pass: 100%/1   | Total:  2m 18s | Avg:  2m 18s | Max:  2m 18s | Hits:  98%/160   
  🟩 Test               Pass: 100%/1   | Total: 21m 58s | Avg: 21m 58s | Max: 21m 58s | Hits:  98%/160

👃 Inspect Changes

Modifications in project?

	Project
	CCCL Infrastructure
	libcu++
	CUB
+/-	Thrust
	CUDA Experimental
	stdpar
	python
	CCCL C Parallel Library
	Catch2Helper

Modifications in project or dependencies?

	Project
	CCCL Infrastructure
	libcu++
+/-	CUB
+/-	Thrust
	CUDA Experimental
+/-	stdpar
+/-	python
+/-	CCCL C Parallel Library
+/-	Catch2Helper

🏃‍ Runner counts (total jobs: 97)

#	Runner
68	`linux-amd64-cpu16`
9	`windows-amd64-cpu16`
6	`linux-arm64-cpu16`
6	`linux-amd64-gpu-rtxa6000-latest-1`
3	`linux-amd64-gpu-h100-latest-1`
3	`linux-amd64-gpu-rtx4090-latest-1`
2	`linux-amd64-gpu-rtx2080-latest-1`

miscco

I believe there is no reason to not use asynchronous allocations.

par_nosync relates to synchronization at the end of the algorithm and we should thrive the make the internals as efficient as possible

miscco · 2025-03-20T07:01:18Z

thrust/thrust/system/cuda/detail/temporary_buffer.h

+  void* ptr;
+  cudaError_t status = cudaMallocAsync(&ptr, sizeof(T) * n, nullptr);
+
+  if (status != cudaSuccess)
+  {
+    cudaGetLastError(); // Clear the CUDA global error state.
+
+    // That didn't work. We could be somewhere where async allocation isn't
+    // supported like Windows, so try again with cudaMalloc.
+    status = cudaMalloc(&ptr, sizeof(T) * n);
+
+    if (status != cudaSuccess)
+    {
+      throw system::detail::bad_alloc(cuda_category().message(status).c_str());
+    }
+  }


That should be a compile time switch and not runtime

What do you mean? Status is a runtime variable.

miscco · 2025-03-20T07:02:40Z

thrust/thrust/system/cuda/detail/temporary_buffer.h

+template <typename Pointer>
+_CCCL_HOST void return_temporary_buffer(par_nosync_t&, Pointer ptr, ::cuda::std::ptrdiff_t n)
+{
+  void* void_ptr = reinterpret_pointer_cast<void*>(ptr);


In libcu++ we have __voidify which does the proper cast to void thingy

I think the point here is to cast a thrust::device_pointer<T> into it's underlying T*. Since T* converts to void* automatically, we should be able to just thrust::raw_pointer_cast(ptr) here.

Bernhard is right, the intent is to go from device_pointer to T*. reinterpert_pointer_cast is what we use elsewhere in *_temporary_buffer machinery which is why we used it. I also think he's right that we can just use raw_pointer_cast here.

thrust/thrust/system/cuda/detail/temporary_buffer.h

Co-authored-by: Bernhard Manfred Gruber <[email protected]>

brycelelbach · 2025-03-22T13:42:07Z

Ideally, we need a way to test the fallback cases of synchronous malloc/free, which I'm told isn't available everywhere. I'll check with the driver team and confirm the platforms where this is the case.

github-actions · 2025-03-22T15:10:27Z

🟨 CI finished in 1h 28m: Pass: 55%/97 | Total: 1d 22h | Avg: 28m 59s | Max: 1h 13m | Hits: 75%/57652

🟨 thrust: Pass: 4%/45 | Total: 6h 53m | Avg: 9m 10s | Max: 57m 33s | Hits: 77%/3552

🚨 cudacxx_family: nvcc 🚨
  🟩 ClangCUDA          Pass: 100%/2   | Total: 50m 31s | Avg: 25m 15s | Max: 25m 44s | Hits:  77%/3552  
  🔥 nvcc               Pass:   0%/43  | Total:  6h 02m | Avg:  8m 25s | Max: 57m 33s
🟨 cudacxx
  🟩 ClangCUDA18        Pass: 100%/2   | Total: 50m 31s | Avg: 25m 15s | Max: 25m 44s | Hits:  77%/3552  
  🟥 nvcc12.0           Pass:   0%/5   | Total:  1h 10m | Avg: 14m 09s | Max: 51m 57s
  🟥 nvcc12.6           Pass:   0%/2   | Total: 14m 26s | Avg:  7m 13s | Max:  7m 15s
  🟥 nvcc12.8           Pass:   0%/36  | Total:  4h 37m | Avg:  7m 42s | Max: 57m 33s
🟥 cmake_options
  🟥 -DTHRUST_DISPATCH_TYPE=Force32bit Pass:   0%/2   | Total:  3m 58s | Avg:  1m 59s | Max:  3m 58s
🟨 cpu
  🟨 amd64              Pass:   4%/43  | Total:  6h 43m | Avg:  9m 23s | Max: 57m 33s | Hits:  77%/3552  
  🟥 arm64              Pass:   0%/2   | Total:  9m 21s | Avg:  4m 40s | Max:  4m 43s
🟨 ctk
  🟥 12.0               Pass:   0%/5   | Total:  1h 10m | Avg: 14m 09s | Max: 51m 57s
  🟥 12.6               Pass:   0%/2   | Total: 14m 26s | Avg:  7m 13s | Max:  7m 15s
  🟨 12.8               Pass:   5%/38  | Total:  5h 27m | Avg:  8m 37s | Max: 57m 33s | Hits:  77%/3552  
🟨 cxx
  🟥 Clang14            Pass:   0%/4   | Total: 18m 17s | Avg:  4m 34s | Max:  4m 56s
  🟥 Clang15            Pass:   0%/2   | Total:  8m 41s | Avg:  4m 20s | Max:  4m 27s
  🟥 Clang16            Pass:   0%/2   | Total:  9m 00s | Avg:  4m 30s | Max:  4m 33s
  🟥 Clang17            Pass:   0%/2   | Total:  9m 23s | Avg:  4m 41s | Max:  4m 50s
  🟨 Clang18            Pass:  28%/7   | Total:  1h 04m | Avg:  9m 10s | Max: 25m 44s | Hits:  77%/3552  
  🟥 GCC7               Pass:   0%/2   | Total:  9m 16s | Avg:  4m 38s | Max:  4m 46s
  🟥 GCC8               Pass:   0%/1   | Total:  4m 30s | Avg:  4m 30s | Max:  4m 30s
  🟥 GCC9               Pass:   0%/2   | Total:  9m 02s | Avg:  4m 31s | Max:  4m 36s
  🟥 GCC10              Pass:   0%/2   | Total:  8m 28s | Avg:  4m 14s | Max:  4m 17s
  🟥 GCC11              Pass:   0%/2   | Total:  8m 23s | Avg:  4m 11s | Max:  4m 12s
  🟥 GCC12              Pass:   0%/2   | Total:  8m 42s | Avg:  4m 21s | Max:  4m 22s
  🟥 GCC13              Pass:   0%/10  | Total: 25m 03s | Avg:  2m 30s | Max:  4m 38s
  🟥 MSVC14.29          Pass:   0%/2   | Total:  1h 49m | Avg: 54m 45s | Max: 57m 33s
  🟥 MSVC14.42          Pass:   0%/3   | Total:  1h 46m | Avg: 35m 23s | Max: 53m 32s
  🟥 NVHPC25.1          Pass:   0%/2   | Total: 14m 26s | Avg:  7m 13s | Max:  7m 15s
🟨 cxx_family
  🟨 Clang              Pass:  11%/17  | Total:  1h 49m | Avg:  6m 26s | Max: 25m 44s | Hits:  77%/3552  
  🟥 GCC                Pass:   0%/21  | Total:  1h 13m | Avg:  3m 29s | Max:  4m 46s
  🟥 MSVC               Pass:   0%/5   | Total:  3h 35m | Avg: 43m 07s | Max: 57m 33s
  🟥 NVHPC              Pass:   0%/2   | Total: 14m 26s | Avg:  7m 13s | Max:  7m 15s
🟨 gpu
  🟥 h100               Pass:   0%/2   | Total:  3m 07s | Avg:  1m 33s | Max:  3m 07s
  🟨 rtx2080            Pass:   6%/33  | Total:  5h 44m | Avg: 10m 26s | Max: 57m 33s | Hits:  77%/3552  
  🟥 rtx4090            Pass:   0%/10  | Total:  1h 05m | Avg:  6m 33s | Max: 52m 37s
🟨 jobs
  🟨 Build              Pass:   5%/38  | Total:  6h 53m | Avg: 10m 52s | Max: 57m 33s | Hits:  77%/3552  
  🟥 TestCPU            Pass:   0%/3  
  🟥 TestGPU            Pass:   0%/4  
🟥 sm
  🟥 90                 Pass:   0%/2   | Total:  3m 07s | Avg:  1m 33s | Max:  3m 07s
  🟥 90;90a;100         Pass:   0%/1   | Total:  4m 25s | Avg:  4m 25s | Max:  4m 25s
🟨 std
  🟨 17                 Pass:   5%/20  | Total:  4h 23m | Avg: 13m 10s | Max: 57m 33s | Hits:  77%/1776  
  🟨 20                 Pass:   4%/23  | Total:  2h 25m | Avg:  6m 19s | Max: 52m 37s | Hits:  77%/1776

🟩 cub: Pass: 100%/45 | Total: 1d 14h | Avg: 51m 00s | Max: 1h 13m | Hits: 75%/53780

🟩 cpu
  🟩 amd64              Pass: 100%/43  | Total:  1d 12h | Avg: 50m 49s | Max:  1h 13m | Hits:  75%/51336 
  🟩 arm64              Pass: 100%/2   | Total:  1h 49m | Avg: 54m 49s | Max:  1h 01m | Hits:  69%/2444  
🟩 ctk
  🟩 12.0               Pass: 100%/5   | Total:  4h 54m | Avg: 58m 55s | Max:  1h 06m | Hits:  70%/5940  
  🟩 12.6               Pass: 100%/2   | Total:  2h 25m | Avg:  1h 12m | Max:  1h 13m | Hits:  69%/2260  
  🟩 12.8               Pass: 100%/38  | Total:  1d 06h | Avg: 48m 48s | Max:  1h 09m | Hits:  76%/45580 
🟩 cudacxx
  🟩 ClangCUDA18        Pass: 100%/2   | Total:  2h 00m | Avg:  1h 00m | Max:  1h 00m | Hits:  75%/2108  
  🟩 nvcc12.0           Pass: 100%/5   | Total:  4h 54m | Avg: 58m 55s | Max:  1h 06m | Hits:  70%/5940  
  🟩 nvcc12.6           Pass: 100%/2   | Total:  2h 25m | Avg:  1h 12m | Max:  1h 13m | Hits:  69%/2260  
  🟩 nvcc12.8           Pass: 100%/36  | Total:  1d 04h | Avg: 48m 10s | Max:  1h 09m | Hits:  76%/43472 
🟩 cudacxx_family
  🟩 ClangCUDA          Pass: 100%/2   | Total:  2h 00m | Avg:  1h 00m | Max:  1h 00m | Hits:  75%/2108  
  🟩 nvcc               Pass: 100%/43  | Total:  1d 12h | Avg: 50m 35s | Max:  1h 13m | Hits:  75%/51672 
🟩 cxx
  🟩 Clang14            Pass: 100%/4   | Total:  4h 13m | Avg:  1h 03m | Max:  1h 06m | Hits:  69%/4896  
  🟩 Clang15            Pass: 100%/2   | Total:  2h 00m | Avg:  1h 00m | Max:  1h 00m | Hits:  69%/2444  
  🟩 Clang16            Pass: 100%/2   | Total:  2h 09m | Avg:  1h 04m | Max:  1h 07m | Hits:  69%/2444  
  🟩 Clang17            Pass: 100%/2   | Total:  2h 10m | Avg:  1h 05m | Max:  1h 09m | Hits:  69%/2444  
  🟩 Clang18            Pass: 100%/7   | Total:  5h 54m | Avg: 50m 36s | Max:  1h 03m | Hits:  80%/8218  
  🟩 GCC7               Pass: 100%/2   | Total:  1h 48m | Avg: 54m 05s | Max: 58m 40s | Hits:  69%/2448  
  🟩 GCC8               Pass: 100%/1   | Total: 47m 05s | Avg: 47m 05s | Max: 47m 05s | Hits:  69%/1224  
  🟩 GCC9               Pass: 100%/2   | Total:  1h 45m | Avg: 52m 44s | Max: 52m 56s | Hits:  69%/2448  
  🟩 GCC10              Pass: 100%/2   | Total:  1h 36m | Avg: 48m 15s | Max: 48m 31s | Hits:  69%/2448  
  🟩 GCC11              Pass: 100%/2   | Total:  1h 40m | Avg: 50m 10s | Max: 52m 02s | Hits:  69%/2444  
  🟩 GCC12              Pass: 100%/2   | Total:  1h 40m | Avg: 50m 23s | Max: 53m 12s | Hits:  69%/2444  
  🟩 GCC13              Pass: 100%/11  | Total:  6h 15m | Avg: 34m 10s | Max: 52m 04s | Hits:  85%/13442 
  🟩 MSVC14.29          Pass: 100%/2   | Total:  1h 53m | Avg: 56m 39s | Max: 58m 26s | Hits:  75%/2088  
  🟩 MSVC14.42          Pass: 100%/2   | Total:  1h 54m | Avg: 57m 17s | Max: 58m 36s | Hits:  75%/2088  
  🟩 NVHPC25.1          Pass: 100%/2   | Total:  2h 25m | Avg:  1h 12m | Max:  1h 13m | Hits:  69%/2260  
🟩 cxx_family
  🟩 Clang              Pass: 100%/17  | Total: 16h 27m | Avg: 58m 04s | Max:  1h 09m | Hits:  73%/20446 
  🟩 GCC                Pass: 100%/22  | Total: 15h 34m | Avg: 42m 27s | Max: 58m 40s | Hits:  77%/26898 
  🟩 MSVC               Pass: 100%/4   | Total:  3h 47m | Avg: 56m 58s | Max: 58m 36s | Hits:  75%/4176  
  🟩 NVHPC              Pass: 100%/2   | Total:  2h 25m | Avg:  1h 12m | Max:  1h 13m | Hits:  69%/2260  
🟩 gpu
  🟩 h100               Pass: 100%/3   | Total:  1h 16m | Avg: 25m 35s | Max: 29m 10s | Hits:  89%/3666  
  🟩 rtx2080            Pass: 100%/34  | Total:  1d 08h | Avg: 57m 33s | Max:  1h 13m | Hits:  70%/40338 
  🟩 rtxa6000           Pass: 100%/8   | Total:  4h 21m | Avg: 32m 41s | Max:  1h 00m | Hits:  92%/9776  
🟩 jobs
  🟩 Build              Pass: 100%/37  | Total:  1d 10h | Avg: 56m 27s | Max:  1h 13m | Hits:  70%/44004 
  🟩 DeviceLaunch       Pass: 100%/1   | Total: 27m 44s | Avg: 27m 44s | Max: 27m 44s | Hits:  99%/1222  
  🟩 GraphCapture       Pass: 100%/1   | Total: 24m 22s | Avg: 24m 22s | Max: 24m 22s | Hits:  99%/1222  
  🟩 HostLaunch         Pass: 100%/3   | Total:  1h 23m | Avg: 27m 56s | Max: 29m 10s | Hits:  99%/3666  
  🟩 TestGPU            Pass: 100%/3   | Total:  1h 10m | Avg: 23m 31s | Max: 25m 06s | Hits:  99%/3666  
🟩 sm
  🟩 90                 Pass: 100%/3   | Total:  1h 16m | Avg: 25m 35s | Max: 29m 10s | Hits:  89%/3666  
  🟩 90;90a;100         Pass: 100%/1   | Total: 46m 01s | Avg: 46m 01s | Max: 46m 01s | Hits:  69%/1222  
🟩 std
  🟩 17                 Pass: 100%/20  | Total: 19h 15m | Avg: 57m 45s | Max:  1h 13m | Hits:  70%/23662 
  🟩 20                 Pass: 100%/25  | Total: 19h 00m | Avg: 45m 36s | Max:  1h 12m | Hits:  79%/30118

🟩 stdpar: Pass: 100%/4 | Total: 17m 06s | Avg: 4m 16s | Max: 5m 14s

🟩 cpu
  🟩 amd64              Pass: 100%/2   | Total: 10m 01s | Avg:  5m 00s | Max:  5m 14s
  🟩 arm64              Pass: 100%/2   | Total:  7m 05s | Avg:  3m 32s | Max:  3m 33s
🟩 ctk
  🟩 12.6               Pass: 100%/4   | Total: 17m 06s | Avg:  4m 16s | Max:  5m 14s
🟩 cudacxx
  🟩 nvcc12.6           Pass: 100%/4   | Total: 17m 06s | Avg:  4m 16s | Max:  5m 14s
🟩 cudacxx_family
  🟩 nvcc               Pass: 100%/4   | Total: 17m 06s | Avg:  4m 16s | Max:  5m 14s
🟩 cxx
  🟩 NVHPC25.1          Pass: 100%/4   | Total: 17m 06s | Avg:  4m 16s | Max:  5m 14s
🟩 cxx_family
  🟩 NVHPC              Pass: 100%/4   | Total: 17m 06s | Avg:  4m 16s | Max:  5m 14s
🟩 gpu
  🟩 rtx2080            Pass: 100%/4   | Total: 17m 06s | Avg:  4m 16s | Max:  5m 14s
🟩 jobs
  🟩 Build              Pass: 100%/4   | Total: 17m 06s | Avg:  4m 16s | Max:  5m 14s
🟩 std
  🟩 17                 Pass: 100%/2   | Total:  8m 20s | Avg:  4m 10s | Max:  4m 47s
  🟩 20                 Pass: 100%/2   | Total:  8m 46s | Avg:  4m 23s | Max:  5m 14s

🟩 cccl_c_parallel: Pass: 100%/2 | Total: 18m 12s | Avg: 9m 06s | Max: 15m 51s | Hits: 98%/320

🟩 cpu
  🟩 amd64              Pass: 100%/2   | Total: 18m 12s | Avg:  9m 06s | Max: 15m 51s | Hits:  98%/320   
🟩 ctk
  🟩 12.8               Pass: 100%/2   | Total: 18m 12s | Avg:  9m 06s | Max: 15m 51s | Hits:  98%/320   
🟩 cudacxx
  🟩 nvcc12.8           Pass: 100%/2   | Total: 18m 12s | Avg:  9m 06s | Max: 15m 51s | Hits:  98%/320   
🟩 cudacxx_family
  🟩 nvcc               Pass: 100%/2   | Total: 18m 12s | Avg:  9m 06s | Max: 15m 51s | Hits:  98%/320   
🟩 cxx
  🟩 GCC13              Pass: 100%/2   | Total: 18m 12s | Avg:  9m 06s | Max: 15m 51s | Hits:  98%/320   
🟩 cxx_family
  🟩 GCC                Pass: 100%/2   | Total: 18m 12s | Avg:  9m 06s | Max: 15m 51s | Hits:  98%/320   
🟩 gpu
  🟩 rtx2080            Pass: 100%/2   | Total: 18m 12s | Avg:  9m 06s | Max: 15m 51s | Hits:  98%/320   
🟩 jobs
  🟩 Build              Pass: 100%/1   | Total:  2m 21s | Avg:  2m 21s | Max:  2m 21s | Hits:  98%/160   
  🟩 Test               Pass: 100%/1   | Total: 15m 51s | Avg: 15m 51s | Max: 15m 51s | Hits:  98%/160

🟩 python: Pass: 100%/1 | Total: 1h 08m | Avg: 1h 08m | Max: 1h 08m

🟩 cpu
  🟩 amd64              Pass: 100%/1   | Total:  1h 08m | Avg:  1h 08m | Max:  1h 08m
🟩 ctk
  🟩 12.8               Pass: 100%/1   | Total:  1h 08m | Avg:  1h 08m | Max:  1h 08m
🟩 cudacxx
  🟩 nvcc12.8           Pass: 100%/1   | Total:  1h 08m | Avg:  1h 08m | Max:  1h 08m
🟩 cudacxx_family
  🟩 nvcc               Pass: 100%/1   | Total:  1h 08m | Avg:  1h 08m | Max:  1h 08m
🟩 cxx
  🟩 GCC13              Pass: 100%/1   | Total:  1h 08m | Avg:  1h 08m | Max:  1h 08m
🟩 cxx_family
  🟩 GCC                Pass: 100%/1   | Total:  1h 08m | Avg:  1h 08m | Max:  1h 08m
🟩 gpu
  🟩 rtx2080            Pass: 100%/1   | Total:  1h 08m | Avg:  1h 08m | Max:  1h 08m
🟩 jobs
  🟩 Test               Pass: 100%/1   | Total:  1h 08m | Avg:  1h 08m | Max:  1h 08m

👃 Inspect Changes

Modifications in project?

	Project
	CCCL Infrastructure
	libcu++
	CUB
+/-	Thrust
	CUDA Experimental
	stdpar
	python
	CCCL C Parallel Library
	Catch2Helper

Modifications in project or dependencies?

	Project
	CCCL Infrastructure
	libcu++
+/-	CUB
+/-	Thrust
	CUDA Experimental
+/-	stdpar
+/-	python
+/-	CCCL C Parallel Library
+/-	Catch2Helper

🏃‍ Runner counts (total jobs: 97)

#	Runner
68	`linux-amd64-cpu16`
9	`windows-amd64-cpu16`
6	`linux-arm64-cpu16`
6	`linux-amd64-gpu-rtxa6000-latest-1`
3	`linux-amd64-gpu-h100-latest-1`
3	`linux-amd64-gpu-rtx4090-latest-1`
2	`linux-amd64-gpu-rtx2080-latest-1`

brycelelbach requested a review from a team as a code owner March 19, 2025 18:35

brycelelbach requested a review from elstehle March 19, 2025 18:35

brycelelbach requested review from bernhardmgruber and gevtushenko March 19, 2025 18:36

[Thrust] Perform asynchronous allocations by default for the `par_nos…

dfb3936

…ync` policy. This will make algorithms (like scans) that don't have a computation-dependent result but do temporary allocation properly asynchronous under `par_nosync`.

bernhardmgruber force-pushed the pr/thrust/par_nosync_async_allocation branch from d6560a0 to ee5c351 Compare March 20, 2025 01:13

bernhardmgruber reviewed Mar 20, 2025

View reviewed changes

thrust/thrust/system/cuda/detail/temporary_buffer.h Outdated Show resolved Hide resolved

Cleanup

7a3edf1

bernhardmgruber force-pushed the pr/thrust/par_nosync_async_allocation branch from ee5c351 to 7a3edf1 Compare March 20, 2025 01:15

miscco reviewed Mar 20, 2025

View reviewed changes

Apply suggestions from code review to par_nosync async allocation

71e8556

Co-authored-by: Bernhard Manfred Gruber <[email protected]>

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[Thrust] Perform asynchronous allocations by default for the `par_nosync` policy #4204

[Thrust] Perform asynchronous allocations by default for the `par_nosync` policy #4204

brycelelbach commented Mar 19, 2025 •

edited

Loading

copy-pr-bot bot commented Mar 19, 2025

bernhardmgruber commented Mar 20, 2025

bernhardmgruber commented Mar 20, 2025

github-actions bot commented Mar 20, 2025

🟨 thrust: Pass: 4%/45 | Total: 6h 52m | Avg: 9m 10s | Max: 58m 21s | Hits: 77%/3552

🟥 python: Pass: 0%/1 | Total: 12m 48s | Avg: 12m 48s | Max: 12m 48s

🟩 cub: Pass: 100%/45 | Total: 1d 14h | Avg: 50m 40s | Max: 1h 15m | Hits: 75%/53780

🟩 stdpar: Pass: 100%/4 | Total: 16m 52s | Avg: 4m 13s | Max: 5m 08s

🟩 cccl_c_parallel: Pass: 100%/2 | Total: 24m 16s | Avg: 12m 08s | Max: 21m 58s | Hits: 98%/320

👃 Inspect Changes

Modifications in project?

Modifications in project or dependencies?

🏃‍ Runner counts (total jobs: 97)

miscco left a comment

miscco Mar 20, 2025

brycelelbach Mar 22, 2025

miscco Mar 20, 2025

bernhardmgruber Mar 21, 2025

brycelelbach Mar 22, 2025

brycelelbach commented Mar 22, 2025

github-actions bot commented Mar 22, 2025

🟨 thrust: Pass: 4%/45 | Total: 6h 53m | Avg: 9m 10s | Max: 57m 33s | Hits: 77%/3552

🟩 cub: Pass: 100%/45 | Total: 1d 14h | Avg: 51m 00s | Max: 1h 13m | Hits: 75%/53780

🟩 stdpar: Pass: 100%/4 | Total: 17m 06s | Avg: 4m 16s | Max: 5m 14s

🟩 cccl_c_parallel: Pass: 100%/2 | Total: 18m 12s | Avg: 9m 06s | Max: 15m 51s | Hits: 98%/320

🟩 python: Pass: 100%/1 | Total: 1h 08m | Avg: 1h 08m | Max: 1h 08m

👃 Inspect Changes

Modifications in project?

Modifications in project or dependencies?

🏃‍ Runner counts (total jobs: 97)

[Thrust] Perform asynchronous allocations by default for the par_nosync policy #4204

Are you sure you want to change the base?

[Thrust] Perform asynchronous allocations by default for the par_nosync policy #4204

Conversation

brycelelbach commented Mar 19, 2025 • edited Loading

copy-pr-bot bot commented Mar 19, 2025

bernhardmgruber commented Mar 20, 2025

bernhardmgruber commented Mar 20, 2025

github-actions bot commented Mar 20, 2025

🟨 thrust: Pass: 4%/45 | Total: 6h 52m | Avg: 9m 10s | Max: 58m 21s | Hits: 77%/3552

🟥 python: Pass: 0%/1 | Total: 12m 48s | Avg: 12m 48s | Max: 12m 48s

🟩 cub: Pass: 100%/45 | Total: 1d 14h | Avg: 50m 40s | Max: 1h 15m | Hits: 75%/53780

🟩 stdpar: Pass: 100%/4 | Total: 16m 52s | Avg: 4m 13s | Max: 5m 08s

🟩 cccl_c_parallel: Pass: 100%/2 | Total: 24m 16s | Avg: 12m 08s | Max: 21m 58s | Hits: 98%/320

👃 Inspect Changes

Modifications in project?

Modifications in project or dependencies?

🏃‍ Runner counts (total jobs: 97)

miscco left a comment

Choose a reason for hiding this comment

miscco Mar 20, 2025

Choose a reason for hiding this comment

brycelelbach Mar 22, 2025

Choose a reason for hiding this comment

miscco Mar 20, 2025

Choose a reason for hiding this comment

bernhardmgruber Mar 21, 2025

Choose a reason for hiding this comment

brycelelbach Mar 22, 2025

Choose a reason for hiding this comment

brycelelbach commented Mar 22, 2025

github-actions bot commented Mar 22, 2025

🟨 thrust: Pass: 4%/45 | Total: 6h 53m | Avg: 9m 10s | Max: 57m 33s | Hits: 77%/3552

🟩 cub: Pass: 100%/45 | Total: 1d 14h | Avg: 51m 00s | Max: 1h 13m | Hits: 75%/53780

🟩 stdpar: Pass: 100%/4 | Total: 17m 06s | Avg: 4m 16s | Max: 5m 14s

🟩 cccl_c_parallel: Pass: 100%/2 | Total: 18m 12s | Avg: 9m 06s | Max: 15m 51s | Hits: 98%/320

🟩 python: Pass: 100%/1 | Total: 1h 08m | Avg: 1h 08m | Max: 1h 08m

👃 Inspect Changes

Modifications in project?

Modifications in project or dependencies?

🏃‍ Runner counts (total jobs: 97)

[Thrust] Perform asynchronous allocations by default for the `par_nosync` policy #4204

[Thrust] Perform asynchronous allocations by default for the `par_nosync` policy #4204

brycelelbach commented Mar 19, 2025 •

edited

Loading