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

[capybara-club]

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 |
+-----------------------------------------+------------------------+----------------------+

[leofang]

Hi @capybara-club, thanks for reaching out, and happy to learn that cuda.core is useful to your project! 🙂

I think you hit this issue: #950. Under the hood, Program invokes Linker to compile PTX to CUBIN, and all underlying backends share the same compiler cache (from the CUDA driver), which led to what you observed.

The correct fix is to expose no_cache to ProgramOptions, similar to LinkerOptions.no_cache. It makes sense now because NVRTC also exposes the same compiler flag to control the caching behavior, and we do want to expose it anyway (#670). This will allow you to pass this flag to Program, and if an invocation of Linker is needed, it can be further passed down.

Question: How urgent is this to your use case? Would it be OK if we get it done by the next minor release (v0.5.0)?

[capybara-club]

@leofang Yeah, its super useful. I'm hoping to have something cool to show soon!

It's not really urgent, it was just really strange to me that changing the flags

1. Didn't seem to invalidate the cache (forcing a print for ptxas_info) or
2. The ptxas_info was being thrown away and the flag change did actually cause a recompile.

Would you expect different flags would cause a recompile and ptxas_info would print correctly for PTX to CUBIN? Such as --use_fast_math?

[leofang]

I think changing the input flags should cause a recompilation if it is not already recompiled. Best to clean up the CUDA cache before doing this experiment, in case the prior runs affect the cache.