NVIDIA / NVIDIA/cccl

[FEA]: Add a utility to programmatically dump PTX/SASS for a kernel/device function

Open
#4,849 5 comments 1 reaction 0 assignees View on GitHub
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

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.