[BUG]: Strange/complex behavior for ptxas_option -v
Nadie ha tomado este issue todavía.
- Lenguaje dominante
- Cython
- Estrellas
- 3.4k
- Forks
- 329
- Merge medio
- 1 d 23 h
- PR fusionados (30 d)
- 116
Descripción
Is this a duplicate?
- I confirmed there appear to be no duplicate issues for this bug and that I agree to the Code of Conduct
Type of Bug
Silent Failure
Component
cuda.core
Describe the bug
Hello!
It looks like cuda-core JIT respects the caching behavior for nvrtc such that setting the environment variable CUDA_CACHE_DISABLE to 1 causes the ptxas '-v' option to always print register info.
When I set the environment variable and run the script below:
import sys
from cuda.core.experimental import Device, Program, ProgramOptions
cuda_code = r"""
extern "C"
__global__
void
kernel() {
float x = 3.0f;
printf("%f\n", x);
}
"""
dev = Device()
dev.set_current()
capability = dev.compute_capability
arch=f"sm_{capability.major}{capability.minor}"
program_options = \
ProgramOptions(
std="c++11",
arch=arch,
ptxas_options=['-v'],
)
prog = Program(cuda_code, code_type="c++", options=program_options)
ptx = prog.compile("ptx", logs=sys.stdout,).code.decode('utf-8')
print('Compiled PTX')
program_options = \
ProgramOptions(
arch=arch,
ptxas_options=['-v']
)
prog = Program(ptx, code_type="ptx", options=program_options)
cubin = prog.compile("cubin", logs=sys.stdout,)
print('Compiled CUBIN')
The ptxas info always prints, but it only prints when compiling the code from CUDA to PTX:
python ptxas_v.py
ptxas info : 4 bytes gmem, 16 bytes cmem[4]
ptxas info : Compiling entry function 'kernel' for 'sm_89'
ptxas info : Function properties for kernel
ptxas . 8 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 24 registers, used 0 barriers, 8 bytes cumulative stack size, 352 bytes cmem[0]
ptxas info : Compile time = 1.209 ms
Compiled PTX
Compiled CUBIN
I think the ptxas -v information just isn't printing when doing PTX->CUBIN. It doesn't print if I change it from ['-v'] to ['-O1','-v'], or ['-O2','-v'], but setting the ptxas_options to options that don't exist causes a nvJitLink error, i.e. ['-O1a', '-v'].
Any ideas?
Thanks!
How to Reproduce
Run the script above, change the ptxas_options for the PTX -> CUBIN.
Expected behavior
Should see ptxas_info print the verbose message for CUBIN compilation.
Operating System
CachyOS
nvidia-smi output
Fri Oct 10 11:20:36 2025
+-----------------------------------------------------------------------------------------+
| NVIDIA-SMI 580.95.05 Driver Version: 580.95.05 CUDA Version: 13.0 |
+-----------------------------------------+------------------------+----------------------+
| GPU Name Persistence-M | Bus-Id Disp.A | Volatile Uncorr. ECC |
| Fan Temp Perf Pwr:Usage/Cap | Memory-Usage | GPU-Util Compute M. |
| | | MIG M. |
|=========================================+========================+======================|
| 0 NVIDIA GeForce RTX 4090 Off | 00000000:01:00.0 On | Off |
| 0% 47C P5 37W / 450W | 1754MiB / 24564MiB | 1% Default |
| | | N/A |
+-----------------------------------------+------------------------+----------------------+
Guía de contribución
Primeros pasos
- Lee el issue completo y luego la guía de contribución del proyecto.
- Comenta en el issue que vas a ocuparte — evita que dos personas hagan lo mismo.
- Haz un fork del repositorio y trabaja en una rama.
- Abre un pull request que haga referencia al número del issue.
Línea de trabajo
Comienza ejecutando el script proporcionado con cuda.core.experimental Device, Program y ProgramOptions, comparando la compilación de CUDA-to-PTX y PTX-to-CUBIN con las ptxas_options indicadas. Rastrea cómo se pasan ptxas_options y logs a través de ambas rutas de compilación; se considera terminado cuando PTX-to-CUBIN emite la información verbose esperada de ptxas, mientras que las opciones no válidas siguen produciendo el error de nvJitLink indicado.
Escrito por el modelo de indexación a partir del texto del issue.
Evaluación
- Stack tecnológico
- cpp, python
- Área
- tooling
- Tipo de issue
- Error
- Dificultad
- 4/5
- Tiempo estimado
- 3-5 días
- Estado de actividad
- Estancado
- Claridad
- Bastante claro
- Aptitud para principiantes
- 35/100