Replace term "PTX version" by "CUDA architecture" or something else

CUB uses the term "PTX version" in many places, like the PtxVersion* functions to retrieve it, but also amply throughout the GPU architecture dispatching. I think the term is misleading though. In my understanding, the "PTX version" refers to the GPU virtual architecture and not to any version of the PTX ISA.

I think the term "PTX version" comes from the cudaFuncAttributes::ptxVersion member when querying a kernel for its function attributes (using cudaFuncGetAttributes). The documentation says:

The PTX virtual architecture version for which the function was compiled.

So this version number, IIUC, is the GPU virtual architecture that a kernel was compiled for (into PTX) and the currently running binary has PTX assembly for that virtual architecture. The term "compute capability" seems to be a synonym. Both are different from a GPU real architecture, which corresponds to SASS in my understanding (reading some nvcc CLI documentation).

The PTX version btw corresponds to the value of __CUDA_ARCH__ during device compilation, and is part of __CUDA_ARCH_LIST__ and NV_TARGET_SM_INTEGER_LIST (the latter is confusing as well, because it's not an SM version, but a virtual architecture).

CUB also offers a function SmVersion that yields the GPU real architecture for a GPU device. This may be different than the "PTX version" in case we compiled for an older GPU virtual architecture, but run on a newer real architecture. Example: we compiled to PTX for Ampere (compute_80) but run SASS (jitted from Ampere PTX) on Hopper now (sm_90).

I propose we rename all occurences of the term "PTX version" by something like: "[virtual]|[GPU|CUDA] arch[itecture]". As an exercise, search and highlight all occurences of "PTX version" in the CUB documentation on policies and think about whether the new term fits. I think the term "GPU architecture" or "CUDA architecture" would fit here.

I'd suggest going with compute_capability as this is the standard term for this concept.

(I removed a mistaken comment of mine, sorry for noise; after playing around with the ptxVersion attribute in Python, I discovered that it is really just the compute capability set at compile time, not the PTX ISA version, so this term is really misleading. Also, ptxVersion sometimes can be zero 🤷‍♂)

See https://developer.nvidia.com/blog/understanding-ptx-the-assembly-language-of-cuda-gpu-computing/#compute_capability_and_nvidia_gpu_hardware_isas

I'd suggest going with compute_capability as this is the standard term for this concept.

I agree.

after playing around with the ptxVersion attribute in Python, I discovered that it is really just the compute capability set at compile time, not the PTX ISA version, so this term is really misleading.

I had similar confusions and I know I am not alone. We really shouldn't use ptxVersion anymore.

So, I have been revisiting this discussion as part of #6544. Tuning policies are currently written for compute capabilities. For example, we have different tunings for DeviceTransform when running on Ampere (8.0), Hopper (9.0) and Blackwell (10.0). On the host, we query the compute capability for which we have PTX or SASS available to run on the current device. This compute capability is also used on the host side of a device algorithm to select a tuning.

On the device side however, there is more. We can also have arch-specific code paths and family-secific code paths. There are corresponding macros to detect this during device compilation. We could also imagine different tunings being used depending on whether we compile for sm_100 or sm100a for example. So in the kernel, we are dealing with architectures.

Interestingly, if we compile for sm_100a, the architecture-specificness is also embedded into the cubin. cub::PtxVersion will report 1000, but cuobjdump will show:

Fatbin elf code:
================
arch = sm_100a
...
	code for sm_100a
	.target	sm_100a

Fatbin ptx code:
================
arch = sm_100
...

So the cubin does contain information beyond a compute capability, even though our APIs don't report this at the moment. So from the cubin, we could actually query an architecture and not only a compute capability.

In CCCL.C, we don't have a cubin, so we query the current device, which only has a compute capability.

This analysis shifts me towards expressing tuning policies for SM architectures instead of compute capabilities. Host-side APIs currently deal in compute capabilities. So even though we don't make a difference now, we could in the future, and we are about to put the tuning API into the public, so we may speculate a bit into the future.

Fortunately, the newly added abstractions from #5279 give us both, a cuda::arch_id and a cuda::compute_capability, which can be converted between each other (constructor and cuda::to_arch_id). So converting from a compute capability in host code to an arch_id to select a tuning is trivial.

We indeed replaced PTX version in most places by compute capability now.

Navigation Menu

Repository navigation

Replace term "PTX version" by "CUDA architecture" or something else #4081

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

Replace term "PTX version" by "CUDA architecture" or something else #4081

Description

Activity

jrhemstad commented on Mar 11, 2025

leofang commented on Mar 11, 2025

jrhemstad commented on Mar 12, 2025

bernhardmgruber commented on Jun 24, 2025

bernhardmgruber commented on Nov 17, 2025

bernhardmgruber commented on May 18, 2026

Metadata

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

Issue actions