[FEA]: Add a utility to programmatically dump PTX/SASS for a kernel/device function
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 486
- Avg merge
- 2d 6h
- Merged PRs (30d)
- 295
Description
### Is this a duplicate?
- [x] I confirmed there appear to be no [duplicate issues](https://github.com/NVIDIA/cccl/issues) for this request and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)
### Area
General CCCL
### Is your feature request related to a problem? Please describe.
As a CUDA C++ developer, I would like the ability to programmatically query the generated PTX/SASS for a specified kernel/device function. This would be useful for a variety of things, like verifying the absence of local memory instructions in CI/testing.
### Describe the solution you'd like
In an ideal world, I would like an API that I can just pass a pointer to a `__global__` or `__device__` function and get the PTX/SASS as a string:
```
__device__ void foo(){...}
template
__global__ void kernel(){...}
std::string ptx = cuda::get_ptx(foo);
std::string sass = cuda::get_sass(kernel);
```
I'm not sure if this is actually possible with the APIs the driver/runtime exposes today.
### Describe alternatives you've considered
_No response_
### Additional context
@gevtushenko built a version of this for our own testing in c.parallel:
https://github.com/NVIDIA/cccl/blob/8b8e2818d8aa3c7dfd1ade62075a96f2b2b626de/c/parallel/test/test_util.h#L32-L80
Contributor guide
Assessment
This issue has not been assessed yet.