cuda - Detecting ptx kernel of Thrust transform

Question

I have following thrust::transform call.

my_functor *f_1 = new my_functor();
thrust::transform(data.begin(), data.end(), data.begin(),*f_1);

I want to detect it's corresponding kernel in PTX file. But there are many kernels containing my_functor in their mangled names.

For example-

_ZN6thrust6system4cuda6detail6detail23launch_closure_by_valueINS2_17for_each_n_detail18for_each_n_closureINS_12zip_iteratorINS_5tupleINS_6detail15normal_iteratorINS_10device_ptrIiEEEESD_NS_9null_typeESE_SE_SE_SE_SE_SE_SE_EEEEjNS9_30device_unary_transform_functorI10my_functorEENS3_20blocked_thread_arrayEEEEEvT_

_ZN6thrust6system4cuda6detail6detail23launch_closure_by_valueINS2_17for_each_n_detail18for_each_n_closureINS_12zip_iteratorINS_5tupleINS_6detail15normal_iteratorINS_10device_ptrIiEEEESD_NS_9null_typeESE_SE_SE_SE_SE_SE_SE_EEEElNS9_30device_unary_transform_functorI10my_functorEENS3_20blocked_thread_arrayEEEEEvT_

_ZN6thrust6detail15device_functionINS0_30device_unary_transform_functorI10my_functorEEvEC1ERKS4_

Which kernel is launched and what are these other kernels?

Answer 1

If you are using Visual Studio, use Nvidia NSIGHT Visual Studio Edition which comes with the CUDA Toolkit.

Go to the "Nsight" menu, click on the "Start Performance Analysis..." entry.

• In "Activity type", select "Profile CUDA Application"
• In "Experiment settings", tick "Collect Information for CUDA Source View"
• Choose "All" in the "Experiments to Run" listbox
• In "Capture Control", tick "Open Report on Stop" and select "CUDA Source View" in the listbox

Then, click on "Launch" and wait for your application to be fully executed. You will see additional output in the console from Nsight.

After the execution, the "CUDA Source View" window will open. - Select "Source and PTX" in the "View" listbox You will be able to find the correspondance between source code and generated PTX. When you click on a line in the source code, one or more lines are highlighted in green in the PTX code.