Summary
cpp/include/nvforest/detail/gpu_introspection.hpp does not classify architecture 1070 as supporting 1,024 resident threads per SM. It therefore falls through to:
MAX_THREADS_PER_SM = 2048
Impact
MIN_BLOCKS_PER_SM is derived from MAX_THREADS_PER_SM / MAX_THREADS_PER_BLOCK. For architecture 1070, this produces 8 instead of 4.
This value is used by infer_kernel in its __launch_bounds__ declaration. As a result, downstream cuML source builds on architecture 1070 require a local patch adding __CUDA_ARCH__ == 1070 to the 1,024-thread branch.
Expected behavior
Architecture 1070 uses:
MAX_THREADS_PER_SM = 1024
MIN_BLOCKS_PER_SM = 4
Relevant code
cpp/include/nvforest/detail/gpu_introspection.hpp
cpp/include/nvforest/detail/infer_kernel/gpu.cuh
Summary
cpp/include/nvforest/detail/gpu_introspection.hppdoes not classify architecture1070as supporting 1,024 resident threads per SM. It therefore falls through to:Impact
MIN_BLOCKS_PER_SMis derived fromMAX_THREADS_PER_SM / MAX_THREADS_PER_BLOCK. For architecture1070, this produces8instead of4.This value is used by
infer_kernelin its__launch_bounds__declaration. As a result, downstream cuML source builds on architecture1070require a local patch adding__CUDA_ARCH__ == 1070to the 1,024-thread branch.Expected behavior
Architecture
1070uses:MAX_THREADS_PER_SM = 1024MIN_BLOCKS_PER_SM = 4Relevant code
cpp/include/nvforest/detail/gpu_introspection.hppcpp/include/nvforest/detail/infer_kernel/gpu.cuh