[Enhancement] Save compilation time from cute templates

[LeiWang1999]

NVIDIA's CUDA Toolkit 12.8 recently published a blog post about optimizing compile times with nvcc: [Optimizing Compile Times for CUDA C++](https://developer.nvidia.com/blog/optimizing-compile-times-for-cuda-c).

We are encountering a similar issue by utilizing cute as our backend, which can lead to significant template compilation overhead. We might benefit from adopting their approach to reduce our compilation times through performance trace optimizations.

To reproduce the compilation time:

import time
import subprocess
import os
start = time.time()


os.system("nvcc -std=c++17 -w -Xcudafe --diag_suppress=177 --compiler-options '-fPIC' -lineinfo --shared /tmp/tmp5hql14m2.cu -lcuda -gencode arch=compute_89,code=sm_89 -I/root/tilelang/tilelang/../src -I/root/tilelang/tilelang/../3rdparty/cutlass/include -diag-suppress=20013 -o /tmp/tmp5hql14m2.so")


end = time.time()
print(f"Time taken: {end - start} seconds")

Need some volunteers who have stable machine with cuda 12.8 installed to help :)

[chenghuaWang]

Which TileLang example generated the file tmp5hql14m2.cu?

[LeiWang1999]

@chenghuaWang Sorry for the confusion. To see the compile command, insert the following line before https://github.com/tile-ai/tilelang/blob/main/tilelang/jit/adapter/libgen.py#L91

print(" ".join(command))

You can observe the compile command when compiling a kernel. The Quick Start guide is also a good reference for obtaining the command.

[zcnrex]

import time
import subprocess
import os
start = time.time()


os.system("nvcc -std=c++17 -w -Xcudafe --diag_suppress=177 --compiler-options '-fPIC' -lineinfo --shared /tmp/tmpqlzpupii.cu -lcuda -gencode arch=compute_89,code=sm_89 -I/home/czeng/tilelang/tilelang/../src -I/home/czeng/tilelang/tilelang/../3rdparty/cutlass/include -diag-suppress=20013 -o /tmp/tmpqlzpupii.so --fdevice-time-trace tilelang-compile")


end = time.time()
print(f"Time taken: {end - start} seconds")

Above without --fdevice-time-trace
Time taken: 17.49080181121826 seconds

with --fdevice-time-trace
Time taken: 25.499202966690063 seconds

$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2025 NVIDIA Corporation
Built on Wed_Jan_15_19:20:09_PST_2025
Cuda compilation tools, release 12.8, V12.8.61
Build cuda_12.8.r12.8/compiler.35404655_0

Shared tracing output offline

[cherichy]

This option is used to generate a trace file of the compilation process, which can be used to analysis the bottleneck rather than accelerating the compilation directly.

[LeiWang1999]

pr #461 may help resolve this problem.