Interpret CUDA profiler log file

Question

Here is the log file of running CUDA profiler (nvprof) on some codes, which have a mix of Thrust, cublas and curand. The first is a kernel I wrote, so no problem there. But I'm not sure how to interpret 2nd to 5th lines, which took up substantial run time.

>  Time(%)   Time    Calls   Avg     Min     Max    Name  %      s       ms      ms      ms
>   
>  28.12     6.82    24,543.00   0.28    0.01    0.64   dev_update_dW1(doub....)
>  23.78     5.77    12,272.00   0.47    0.46    0.49   void thrust::system::cud....
>  14.32     3.47    12,272.00   0.28    0.28    0.29   void thrust::system::cud....
>  10.82     2.62    12,272.00   0.21    0.21    0.22   void thrust::system::cud....
>  4.93      1.20    24,544.00   0.05    0.05    0.05   void thrust::system::cud....
>  3.98      0.96    12,272.00   0.08    0.08    0.09   Act_dAct(double*, long, double*, double*)

The 2nd to 5th lines are printed below in full:

2nd line : void thrust::system::cuda::detail::detail::launch_closure_by_value>, thrust::counting_iterator<__int64, thrust::use_default, thrust::use_default, thrust::use_default>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, __int64, thrust::tuple, thrust::detail::normal_iterator, thrust::system::cuda::detail::tag, thrust::use_default, thrust::use_default>>, thrust::system::detail::generic::detail::max_element_reduction>, thrust::system::cuda::detail::detail::blocked_thread_array>>(double)

3rd line : void thrust::system::cuda::detail::detail::launch_closure_by_value>, thrust::detail::normal_iterator>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, unsigned int, thrust::detail::device_unary_transform_functor, thrust::system::cuda::detail::detail::blocked_thread_array>>(double)

4th line : void thrust::system::cuda::detail::detail::launch_closure_by_value>, double, thrust::use_default>, __int64, double, thrust::detail::normal_iterator>, thrust::plus, thrust::system::cuda::detail::detail::blocked_thread_array>>(exp_functor)

5th line : void thrust::system::cuda::detail::detail::launch_closure_by_value, unsigned int, thrust::detail::device_generate_functor>, thrust::system::cuda::detail::detail::blocked_thread_array>>(double)

EDIT :

I have this function (softmax) that uses max_element and transform_reduce

void Softmax_ThrustMatrix(thrust::device_vector<double>& mat, int Nrow, int Ncol, thrust::device_vector<double>& Outmat) {
thrust::device_vector<double> x(Ncol, 0.0);
thrust::device_vector<double> v(Ncol, 0.0);
thrust::device_vector<double>::iterator mx;
double tmp = 0.0, logsm=0.0;
dim3 grid, block;

block.x = 16;
block.y = 1;
grid.x = Ncol / block.x + 1;
grid.y = 1;

for ( int i=0; i < Nrow; i++ ) {
    GetRow<<<grid,block>>>(thrust::raw_pointer_cast(&mat[0]), i, Nrow, Ncol, thrust::raw_pointer_cast(&x[0]));

    mx = thrust::max_element(x.begin(), x.end());

    tmp = thrust::transform_reduce(x.begin(), x.end(), exp_functor(*mx), 0.0, thrust::plus<double>() );
    logsm = *mx + log(tmp);

    thrust::transform(x.begin(), x.end(), v.begin(), exp_functor(logsm));

    SetRow<<<grid,block>>>(thrust::raw_pointer_cast(&v[0]), i, Nrow, Ncol, thrust::raw_pointer_cast(&Outmat[0]));
    }
}

Robert Crovella Robert Crovella · Accepted Answer · 2015-01-05T16:12:04

Thrust code at a low level is not any different than CUDA code (at least for thrust code targetting a GPU). Thrust, as a template library, abstracts away many aspects of CUDA at the source code level, but the profiler doesn't know any difference between thrust code and ordinary cuda code.

Lines 2-5 represent the profiler data on 4 kernel launches. It's evident from their syntax that they are probably not kernels you wrote - they are emanating from within the depths of thrust template functions.

"Launch closure" is thrust-under-the-hood-speak for a kernel launched by thrust to perform some function. Since you have 3 thrust calls in the code you have shown, and are also showing GetRow and SetRow kernels that you wrote, and those kernels don't show up in your profiler output anywhere, it's not evident to me that the profiler output you have shown is related to the code you have shown. You haven't shown the code that calls the kernels that do appear in your output ( dev_update_dW1 and Act_dAct), so it seems fairly clear to me that the code you have shown is not useful for further interpretation of your profiler output.

In any event, lines 2-5 represent CUDA kernels, launched by thrust, that are emanating from thrust calls in your code (somewhere).

Note that it's also possible for thrust to launch kernels for some other non-obvious purposes, such as instantiation of device vectors.

Interpret CUDA profiler log file

1 Answers