NVIDIA / NVIDIA/cuda-python

[BUG]: Strange/complex behavior for ptxas_option -v

Offen
#1,124 3 Kommentare 0 Reaktionen 0 zugewiesene Personen Auf GitHub ansehen

Dieses Issue hat noch niemand übernommen.

awaiting-response bug cuda.core
Vorherrschende Sprache
Cython
Sterne
3.4k
Forks
329
Ø Merge
1 T. 23 Std.
Gemergte PRs (30 T.)
116

Beschreibung

Is this a duplicate?
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 |
+-----------------------------------------+------------------------+----------------------+

Beitragsleitfaden

Beitragsleitfaden öffnen

Erste Schritte

  1. Lies das ganze Issue und danach den Beitragsleitfaden des Projekts.
  2. Schreib ins Issue, dass du es übernimmst — das erspart doppelte Arbeit.
  3. Forke das Repository und arbeite in einem Branch.
  4. Öffne einen Pull Request, der die Issue-Nummer nennt.

Rechercherichtung

Beginne damit, das bereitgestellte Skript mit cuda.core.experimental Device, Program und ProgramOptions auszuführen und die CUDA-to-PTX- sowie die PTX-to-CUBIN-Kompilierung mit den aufgeführten ptxas_options zu vergleichen. Verfolge, wie ptxas_options und logs durch beide Kompilierungspfade weitergegeben werden; die Aufgabe ist abgeschlossen, wenn PTX-to-CUBIN die erwarteten ptxas-Verbose-Informationen ausgibt, während ungültige Optionen weiterhin den gemeldeten nvJitLink-Fehler erzeugen.

Vom Indexierungsmodell aus dem Issue-Text verfasst.

Bewertung

Tech-Stack
cpp, python
Bereich
tooling
Issue-Typ
Bug
Schwierigkeit
4/5
Geschätzter Aufwand
3-5 Tage
Aktivitätsstatus
Veraltet
Klarheit
Größtenteils klar
Anfängerfreundlichkeit
35/100

Neue Issues direkt in Ihr Postfach

Eine kurze Übersicht über anfängerfreundliche GitHub-Issues.