The Thrust function below can get the maximum blocks of for a CUDA launch CUDA 5.0, which is used by Sparse Matrix Vector multiplication(SpMV) in CUSP, and it is a technique for setting up execution for persistent threads. The first line is the header file.
#include <thrust/detail/backend/cuda/arch.h>
thrust::detail::backend::cuda::arch::max_active_blocks(kernel<float,int,VECTORS_PER_BLOCK,TH READS_PER_VECTOR>,THREADS_PER_BLOCK,(size_t)0)
But the function is not supported by CUDA 5.5. Was this technique not supported by CUDA 5.5, or should I use some other function instead?
There was never any supported way to perform this computation in any version of Thrust. Headers inside
thrust/detail
and identifiers inside adetail
namespace are part of Thrust's implementation -- they are not public features. Using them will break your code.That said, there's some standalone code implementing the occupancy calculator in this repository:
https://github.com/jaredhoberock/cuda_launch_config