Use variadic interfaces in agent launcher #1804

bernhardmgruber · 2024-06-04T10:34:14Z

I just ran over the agent_launcher.h file in thrust which contained a lot of pre-C++11 workarounds for not having variadic template parameters. This PR replaces those by variadic versions that were already present, but disabled, in the file.

SASS of thrust.test.merge is identical, except for mangled names.

miscco · 2024-06-04T11:02:49Z

thrust/thrust/system/cuda/detail/core/agent_launcher.h

 THRUST_DETAIL_KERNEL_ATTRIBUTES void __launch_bounds__(Agent::ptx_plan::BLOCK_THREADS)
-  _kernel_agent_vshmem(char* vshmem, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8)
+  _kernel_agent_vshmem(char* vshmem, Args... args)
 {
  extern __shared__ char shmem[];
  vshmem = vshmem == NULL ? shmem : vshmem + blockIdx.x * temp_storage_size<typename Agent::ptx_plan>::value;


While we are at it, can we drop the NULL too?

There you go: #1805

github-actions · 2024-06-04T12:51:22Z

🟩 CI finished in 2h 14m: Pass: 100%/249 | Total: 4d 15h | Avg: 26m 51s | Max: 50m 44s | Hits: 71%/248310

🟩 cub: Pass: 100%/131 | Total: 2d 16h | Avg: 29m 27s | Max: 46m 55s | Hits: 71%/109044

🟩 cpu
  🟩 amd64              Pass: 100%/123 | Total:  2d 12h | Avg: 29m 16s | Max: 46m 55s | Hits:  72%/102236
  🟩 arm64              Pass: 100%/8   | Total:  4h 20m | Avg: 32m 30s | Max: 34m 40s | Hits:  61%/6808  
🟩 ctk
  🟩 11.1               Pass: 100%/15  | Total:  6h 41m | Avg: 26m 44s | Max: 46m 55s | Hits:  62%/11554 
  🟩 11.8               Pass: 100%/3   | Total:  2h 06m | Avg: 42m 09s | Max: 42m 32s | Hits:  61%/2553  
  🟩 12.4               Pass: 100%/113 | Total:  2d 07h | Avg: 29m 29s | Max: 42m 58s | Hits:  73%/94937 
🟩 cudacxx_full
  🟩 clang-cuda17       Pass: 100%/2   | Total: 40m 18s | Avg: 20m 09s | Max: 20m 51s | Hits:  67%/1408  
  🟩 nvcc11.1           Pass: 100%/15  | Total:  6h 41m | Avg: 26m 44s | Max: 46m 55s | Hits:  62%/11554 
  🟩 nvcc11.8           Pass: 100%/3   | Total:  2h 06m | Avg: 42m 09s | Max: 42m 32s | Hits:  61%/2553  
  🟩 nvcc12.4           Pass: 100%/111 | Total:  2d 06h | Avg: 29m 39s | Max: 42m 58s | Hits:  73%/93529 
🟩 cudacxx_name
  🟩 clang-cuda         Pass: 100%/2   | Total: 40m 18s | Avg: 20m 09s | Max: 20m 51s | Hits:  67%/1408  
  🟩 nvcc               Pass: 100%/129 | Total:  2d 15h | Avg: 29m 36s | Max: 46m 55s | Hits:  71%/107636
🟩 cxx_full
  🟩 clang9             Pass: 100%/6   | Total:  2h 50m | Avg: 28m 27s | Max: 32m 46s | Hits:  62%/4884  
  🟩 clang10            Pass: 100%/3   | Total:  1h 42m | Avg: 34m 09s | Max: 35m 23s | Hits:  62%/2559  
  🟩 clang11            Pass: 100%/4   | Total:  2h 09m | Avg: 32m 24s | Max: 34m 01s | Hits:  62%/3412  
  🟩 clang12            Pass: 100%/4   | Total:  2h 11m | Avg: 32m 53s | Max: 34m 06s | Hits:  62%/3412  
  🟩 clang13            Pass: 100%/4   | Total:  2h 12m | Avg: 33m 06s | Max: 34m 05s | Hits:  62%/3412  
  🟩 clang14            Pass: 100%/4   | Total:  2h 11m | Avg: 32m 50s | Max: 35m 37s | Hits:  62%/3412  
  🟩 clang15            Pass: 100%/4   | Total:  2h 07m | Avg: 31m 58s | Max: 32m 37s | Hits:  62%/3404  
  🟩 clang16            Pass: 100%/4   | Total:  2h 08m | Avg: 32m 00s | Max: 32m 55s | Hits:  62%/3404  
  🟩 clang17            Pass: 100%/26  | Total: 10h 56m | Avg: 25m 15s | Max: 37m 09s | Hits:  86%/21832 
  🟩 gcc6               Pass: 100%/2   | Total: 50m 02s | Avg: 25m 01s | Max: 26m 14s | Hits:  61%/1550  
  🟩 gcc7               Pass: 100%/6   | Total:  2h 54m | Avg: 29m 06s | Max: 32m 38s | Hits:  61%/4887  
  🟩 gcc8               Pass: 100%/6   | Total:  2h 50m | Avg: 28m 21s | Max: 33m 11s | Hits:  61%/4887  
  🟩 gcc9               Pass: 100%/6   | Total:  2h 55m | Avg: 29m 15s | Max: 33m 19s | Hits:  61%/4887  
  🟩 gcc10              Pass: 100%/4   | Total:  2h 12m | Avg: 33m 09s | Max: 34m 20s | Hits:  61%/3412  
  🟩 gcc11              Pass: 100%/7   | Total:  4h 15m | Avg: 36m 27s | Max: 42m 32s | Hits:  61%/5957  
  🟩 gcc12              Pass: 100%/4   | Total:  2h 18m | Avg: 34m 39s | Max: 37m 29s | Hits:  61%/3404  
  🟩 gcc13              Pass: 100%/28  | Total: 11h 32m | Avg: 24m 43s | Max: 36m 58s | Hits:  83%/23828 
  🟩 Intel2023.2.0      Pass: 100%/3   | Total:  1h 45m | Avg: 35m 16s | Max: 36m 36s | Hits:  62%/2331  
  🟩 MSVC14.16          Pass: 100%/1   | Total: 46m 55s | Avg: 46m 55s | Max: 46m 55s | Hits:  66%/695   
  🟩 MSVC14.29          Pass: 100%/2   | Total:  1h 22m | Avg: 41m 13s | Max: 42m 58s | Hits:  66%/1390  
  🟩 MSVC14.39          Pass: 100%/3   | Total:  2h 04m | Avg: 41m 39s | Max: 42m 48s | Hits:  66%/2085  
🟩 cxx_name
  🟩 clang              Pass: 100%/59  | Total:  1d 04h | Avg: 28m 59s | Max: 37m 09s | Hits:  72%/49731 
  🟩 gcc                Pass: 100%/63  | Total:  1d 05h | Avg: 28m 23s | Max: 42m 32s | Hits:  71%/52812 
  🟩 Intel              Pass: 100%/3   | Total:  1h 45m | Avg: 35m 16s | Max: 36m 36s | Hits:  62%/2331  
  🟩 MSVC               Pass: 100%/6   | Total:  4h 14m | Avg: 42m 23s | Max: 46m 55s | Hits:  66%/4170  
🟩 gpu
  🟩 v100               Pass: 100%/131 | Total:  2d 16h | Avg: 29m 27s | Max: 46m 55s | Hits:  71%/109044
🟩 jobs
  🟩 Build              Pass: 100%/99  | Total:  2d 04h | Avg: 31m 42s | Max: 46m 55s | Hits:  62%/81812 
  🟩 DeviceLaunch       Pass: 100%/8   | Total:  2h 46m | Avg: 20m 48s | Max: 25m 17s | Hits:  99%/6808  
  🟩 GraphCapture       Pass: 100%/8   | Total:  2h 23m | Avg: 17m 54s | Max: 22m 08s | Hits:  99%/6808  
  🟩 HostLaunch         Pass: 100%/8   | Total:  2h 49m | Avg: 21m 09s | Max: 27m 31s | Hits:  99%/6808  
  🟩 TestGPU            Pass: 100%/8   | Total:  4h 01m | Avg: 30m 12s | Max: 37m 09s | Hits:  99%/6808  
🟩 os
  🟩 ubuntu18.04        Pass: 100%/14  | Total:  5h 54m | Avg: 25m 18s | Max: 29m 27s | Hits:  61%/10859 
  🟩 ubuntu20.04        Pass: 100%/35  | Total: 19h 06m | Avg: 32m 46s | Max: 35m 37s | Hits:  61%/29855 
  🟩 ubuntu22.04        Pass: 100%/76  | Total:  1d 11h | Avg: 27m 41s | Max: 42m 32s | Hits:  78%/64160 
  🟩 windows2022        Pass: 100%/6   | Total:  4h 14m | Avg: 42m 23s | Max: 46m 55s | Hits:  66%/4170  
🟩 sm
  🟩 60;70;80;90        Pass: 100%/3   | Total:  2h 06m | Avg: 42m 09s | Max: 42m 32s | Hits:  61%/2553  
  🟩 90a                Pass: 100%/4   | Total:  1h 10m | Avg: 17m 36s | Max: 18m 38s | Hits:  61%/3404  
🟩 std
  🟩 11                 Pass: 100%/34  | Total: 16h 24m | Avg: 28m 56s | Max: 42m 32s | Hits:  70%/28503 
  🟩 14                 Pass: 100%/37  | Total: 18h 40m | Avg: 30m 16s | Max: 46m 55s | Hits:  70%/30588 
  🟩 17                 Pass: 100%/36  | Total: 17h 58m | Avg: 29m 57s | Max: 42m 13s | Hits:  70%/29822 
  🟩 20                 Pass: 100%/24  | Total: 11h 17m | Avg: 28m 13s | Max: 42m 02s | Hits:  75%/20131

🟩 thrust: Pass: 100%/118 | Total: 1d 23h | Avg: 23m 58s | Max: 50m 44s | Hits: 72%/139266

🟩 cpu
  🟩 amd64              Pass: 100%/110 | Total:  1d 19h | Avg: 23m 54s | Max: 50m 44s | Hits:  72%/129822
  🟩 arm64              Pass: 100%/8   | Total:  3h 18m | Avg: 24m 50s | Max: 29m 42s | Hits:  67%/9444  
🟩 ctk
  🟩 11.1               Pass: 100%/15  | Total:  5h 59m | Avg: 23m 58s | Max: 44m 58s | Hits:  67%/17705 
  🟩 11.8               Pass: 100%/3   | Total:  1h 42m | Avg: 34m 06s | Max: 35m 58s | Hits:  67%/3543  
  🟩 12.4               Pass: 100%/100 | Total:  1d 15h | Avg: 23m 40s | Max: 50m 44s | Hits:  73%/118018
🟩 cudacxx_full
  🟩 clang-cuda17       Pass: 100%/2   | Total: 45m 30s | Avg: 22m 45s | Max: 22m 54s | Hits:  66%/2360  
  🟩 nvcc11.1           Pass: 100%/15  | Total:  5h 59m | Avg: 23m 58s | Max: 44m 58s | Hits:  67%/17705 
  🟩 nvcc11.8           Pass: 100%/3   | Total:  1h 42m | Avg: 34m 06s | Max: 35m 58s | Hits:  67%/3543  
  🟩 nvcc12.4           Pass: 100%/98  | Total:  1d 14h | Avg: 23m 41s | Max: 50m 44s | Hits:  73%/115658
🟩 cudacxx_name
  🟩 clang-cuda         Pass: 100%/2   | Total: 45m 30s | Avg: 22m 45s | Max: 22m 54s | Hits:  66%/2360  
  🟩 nvcc               Pass: 100%/116 | Total:  1d 22h | Avg: 23m 59s | Max: 50m 44s | Hits:  72%/136906
🟩 cxx_full
  🟩 clang9             Pass: 100%/6   | Total:  2h 22m | Avg: 23m 45s | Max: 27m 14s | Hits:  67%/7080  
  🟩 clang10            Pass: 100%/3   | Total:  1h 22m | Avg: 27m 29s | Max: 29m 00s | Hits:  67%/3540  
  🟩 clang11            Pass: 100%/4   | Total:  1h 39m | Avg: 24m 45s | Max: 26m 25s | Hits:  67%/4720  
  🟩 clang12            Pass: 100%/4   | Total:  1h 41m | Avg: 25m 16s | Max: 26m 12s | Hits:  67%/4720  
  🟩 clang13            Pass: 100%/4   | Total:  1h 38m | Avg: 24m 38s | Max: 26m 09s | Hits:  67%/4720  
  🟩 clang14            Pass: 100%/4   | Total:  1h 37m | Avg: 24m 27s | Max: 26m 06s | Hits:  67%/4720  
  🟩 clang15            Pass: 100%/4   | Total:  1h 41m | Avg: 25m 23s | Max: 26m 56s | Hits:  67%/4720  
  🟩 clang16            Pass: 100%/4   | Total:  1h 38m | Avg: 24m 40s | Max: 26m 52s | Hits:  67%/4720  
  🟩 clang17            Pass: 100%/18  | Total:  5h 16m | Avg: 17m 34s | Max: 25m 53s | Hits:  81%/21240 
  🟩 gcc6               Pass: 100%/2   | Total: 43m 40s | Avg: 21m 50s | Max: 23m 27s | Hits:  67%/2360  
  🟩 gcc7               Pass: 100%/6   | Total:  2h 22m | Avg: 23m 44s | Max: 27m 11s | Hits:  67%/7086  
  🟩 gcc8               Pass: 100%/6   | Total:  2h 23m | Avg: 23m 55s | Max: 26m 12s | Hits:  67%/7086  
  🟩 gcc9               Pass: 100%/6   | Total:  2h 30m | Avg: 25m 02s | Max: 29m 24s | Hits:  67%/7086  
  🟩 gcc10              Pass: 100%/4   | Total:  1h 49m | Avg: 27m 16s | Max: 29m 19s | Hits:  67%/4724  
  🟩 gcc11              Pass: 100%/7   | Total:  3h 29m | Avg: 29m 57s | Max: 35m 58s | Hits:  67%/8267  
  🟩 gcc12              Pass: 100%/4   | Total:  1h 49m | Avg: 27m 18s | Max: 28m 34s | Hits:  67%/4724  
  🟩 gcc13              Pass: 100%/20  | Total:  5h 43m | Avg: 17m 11s | Max: 29m 42s | Hits:  79%/23620 
  🟩 Intel2023.2.0      Pass: 100%/3   | Total:  1h 35m | Avg: 31m 43s | Max: 34m 26s | Hits:  67%/3549  
  🟩 MSVC14.16          Pass: 100%/1   | Total: 44m 58s | Avg: 44m 58s | Max: 44m 58s | Hits:  65%/1176  
  🟩 MSVC14.29          Pass: 100%/2   | Total:  1h 36m | Avg: 48m 06s | Max: 49m 53s | Hits:  65%/2352  
  🟩 MSVC14.39          Pass: 100%/6   | Total:  3h 23m | Avg: 33m 50s | Max: 50m 44s | Hits:  82%/7056  
🟩 cxx_name
  🟩 clang              Pass: 100%/51  | Total: 18h 58m | Avg: 22m 18s | Max: 29m 00s | Hits:  72%/60180 
  🟩 gcc                Pass: 100%/55  | Total: 20h 51m | Avg: 22m 45s | Max: 35m 58s | Hits:  71%/64953 
  🟩 Intel              Pass: 100%/3   | Total:  1h 35m | Avg: 31m 43s | Max: 34m 26s | Hits:  67%/3549  
  🟩 MSVC               Pass: 100%/9   | Total:  5h 44m | Avg: 38m 14s | Max: 50m 44s | Hits:  76%/10584 
🟩 gpu
  🟩 v100               Pass: 100%/118 | Total:  1d 23h | Avg: 23m 58s | Max: 50m 44s | Hits:  72%/139266
🟩 jobs
  🟩 Build              Pass: 100%/99  | Total:  1d 19h | Avg: 26m 29s | Max: 50m 44s | Hits:  66%/116850
  🟩 TestCPU            Pass: 100%/11  | Total:  1h 39m | Avg:  9m 00s | Max: 18m 19s | Hits:  99%/12972 
  🟩 TestGPU            Pass: 100%/8   | Total:  1h 47m | Avg: 13m 22s | Max: 17m 29s | Hits:  99%/9444  
🟩 os
  🟩 ubuntu18.04        Pass: 100%/14  | Total:  5h 14m | Avg: 22m 28s | Max: 25m 10s | Hits:  67%/16529 
  🟩 ubuntu20.04        Pass: 100%/35  | Total: 14h 55m | Avg: 25m 35s | Max: 29m 24s | Hits:  67%/41313 
  🟩 ubuntu22.04        Pass: 100%/60  | Total: 21h 14m | Avg: 21m 14s | Max: 35m 58s | Hits:  75%/70840 
  🟩 windows2022        Pass: 100%/9   | Total:  5h 44m | Avg: 38m 14s | Max: 50m 44s | Hits:  76%/10584 
🟩 sm
  🟩 60;70;80;90        Pass: 100%/3   | Total:  1h 42m | Avg: 34m 06s | Max: 35m 58s | Hits:  67%/3543  
  🟩 90a                Pass: 100%/4   | Total:  1h 01m | Avg: 15m 16s | Max: 17m 44s | Hits:  61%/4724  
🟩 std
  🟩 11                 Pass: 100%/30  | Total: 10h 15m | Avg: 20m 31s | Max: 32m 04s | Hits:  73%/35418 
  🟩 14                 Pass: 100%/34  | Total: 14h 32m | Avg: 25m 40s | Max: 48m 43s | Hits:  71%/40122 
  🟩 17                 Pass: 100%/33  | Total: 14h 01m | Avg: 25m 30s | Max: 50m 44s | Hits:  70%/38946 
  🟩 20                 Pass: 100%/21  | Total:  8h 19m | Avg: 23m 45s | Max: 49m 42s | Hits:  74%/24780

👃 Inspect Changes

Modifications in project?

	Project
	CCCL Infrastructure
	libcu++
	CUB
+/-	Thrust
	CUDA Experimental

Modifications in project or dependencies?

	Project
	CCCL Infrastructure
	libcu++
+/-	CUB
+/-	Thrust
	CUDA Experimental

🏃‍ Runner counts (total jobs: 249)

#	Runner
178	`linux-amd64-cpu16`
40	`linux-amd64-gpu-v100-latest-1`
16	`linux-arm64-cpu16`
15	`windows-amd64-cpu16`

bernhardmgruber marked this pull request as ready for review June 4, 2024 10:35

bernhardmgruber requested review from a team as code owners June 4, 2024 10:35

bernhardmgruber requested review from miscco and ericniebler June 4, 2024 10:35

bernhardmgruber added the thrust label Jun 4, 2024

bernhardmgruber enabled auto-merge (squash) June 4, 2024 10:41

miscco approved these changes Jun 4, 2024

View reviewed changes

bernhardmgruber merged commit 4808778 into NVIDIA:main Jun 4, 2024
556 checks passed

bernhardmgruber deleted the clean_agent_launcher branch June 4, 2024 12:59

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Use variadic interfaces in agent launcher #1804

Use variadic interfaces in agent launcher #1804

bernhardmgruber commented Jun 4, 2024 •

edited

Loading

miscco Jun 4, 2024

bernhardmgruber Jun 4, 2024

github-actions bot commented Jun 4, 2024

🟩 cub: Pass: 100%/131 | Total: 2d 16h | Avg: 29m 27s | Max: 46m 55s | Hits: 71%/109044

🟩 thrust: Pass: 100%/118 | Total: 1d 23h | Avg: 23m 58s | Max: 50m 44s | Hits: 72%/139266

👃 Inspect Changes

Modifications in project?

Modifications in project or dependencies?

🏃‍ Runner counts (total jobs: 249)

Use variadic interfaces in agent launcher #1804

Use variadic interfaces in agent launcher #1804

Conversation

bernhardmgruber commented Jun 4, 2024 • edited Loading

miscco Jun 4, 2024

Choose a reason for hiding this comment

bernhardmgruber Jun 4, 2024

Choose a reason for hiding this comment

github-actions bot commented Jun 4, 2024

🟩 cub: Pass: 100%/131 | Total: 2d 16h | Avg: 29m 27s | Max: 46m 55s | Hits: 71%/109044

🟩 thrust: Pass: 100%/118 | Total: 1d 23h | Avg: 23m 58s | Max: 50m 44s | Hits: 72%/139266

👃 Inspect Changes

Modifications in project?

Modifications in project or dependencies?

🏃‍ Runner counts (total jobs: 249)

bernhardmgruber commented Jun 4, 2024 •

edited

Loading