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

[jrhemstad]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this request and that I agree to the Code of Conduct

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 <typename T>
__global__ void kernel(){...}

std::string ptx = cuda::get_ptx(foo);
std::string sass = cuda::get_sass(kernel<int>);

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:

cccl/c/parallel/test/test_util.h

Lines 32 to 80 in 8b8e281

	inline std::string inspect_sass(const void* cubin, size_t cubin_size)
	{
	namespace fs = std::filesystem;
	fs::path temp_dir = fs::temp_directory_path();
	fs::path temp_in_filename = temp_dir / "temp_in_file.cubin";
	fs::path temp_out_filename = temp_dir / "temp_out_file.sass";
	std::ofstream temp_in_file(temp_in_filename, std::ios::binary);
	if (!temp_in_file)
	{
	throw std::runtime_error("Failed to create temporary file.");
	}
	temp_in_file.write(static_cast<const char*>(cubin), cubin_size);
	temp_in_file.close();
	std::string command = "nvdisasm -gi ";
	command += temp_in_filename;
	command += " > ";
	command += temp_out_filename;
	int exec_code = std::system(command.c_str());
	if (!fs::remove(temp_in_filename))
	{
	throw std::runtime_error("Failed to remove temporary file.");
	}
	if (exec_code != 0)
	{
	throw std::runtime_error("Failed to execute command.");
	}
	std::ifstream temp_out_file(temp_out_filename, std::ios::binary);
	if (!temp_out_file)
	{
	throw std::runtime_error("Failed to create temporary file.");
	}
	const std::string sass{std::istreambuf_iterator<char>(temp_out_file), std::istreambuf_iterator<char>()};
	if (!fs::remove(temp_out_filename))
	{
	throw std::runtime_error("Failed to remove temporary file.");
	}
	return sass;
	}

[gevtushenko]

@pciolkosz can we retreive cubin part for a given kernel by it's address?

[davebayer]

We can! We need to get the CUkernel handle via cudaGetKernel and then retrieve the name via cuKernelGetName.

But I believe we will need that mangled version name for the given architecture.

To get that one we will have to get the kernel symbol address for the current context via cudaGetSymbolAddress, get the CUfuntion handle via cudaGetFuncBySymbol and get the name via cuFuncGetName.

But I will have to test it out.

[leofang]

You can get a handle to the kernel (CUkernel or CUfunction), but there is no public way of retrieving the underlying cubin.

[davebayer]

Aaaah, I misunderstood the question a bit 😅

[davebayer]

In #4996, I'd like to add cuda::demangle function to libcu++. I've put it in (internal) <cuda/binutils> header. I think the header would be a good place for these functions, too