Discrepancy between Documentation and Actual Hardware/Compiler Behavior for __vset* Intrinsics

Accelerated Computing GPU-Accelerated Libraries
1

I would like to report a discrepancy between the CUDA Math API documentation and the actual runtime behavior of __vseteq2 (and potentially other __vset* SIMD intrinsics).

According to the CUDA Math API Manual (CUDA_Math_API.pdf):

“Performs per-halfword (un)signed comparison: returns 1 if both parts compare equal. Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding parts function performs comparison ‘a’ part == ‘b’ part. If both equalities are satisfied, function returns 1. Returns 1 if a = b, else returns 0.”

Based on the documentation, this function should perform a reduction comparison and return a scalar boolean/integer value (1 or 0).

However, the actual implementation returns a packed 32-bit result representing per-halfword comparison masks (e.g., 0x00010001 when both halfwords match, or 0x00010000 when only one halfword matches), rather than a scalar 0 or 1.

Minimal Reproducible Example:

#include <stdio.h>
#include <cuda_runtime.h>

__global__ void test_vseteq2() {
    unsigned int a = 0x12345678;
    unsigned int b = 0x12340000; // High 16 bits match, low 16 bits differ

    unsigned int res = __vseteq2(a, b);

    // Documented behavior: Should return 0 (since low halfwords do not match)
    // Actual behavior:     Returns 0x00010000 (packed per-halfword comparison result)
    printf("a: 0x%08X, b: 0x%08X => __vseteq2 output: 0x%08X\n", a, b, res);
}

int main() {
    test_vseteq2<<<1, 1>>>();
    cudaDeviceSynchronize();
    return 0;
}

Observed Output:

a: 0x12345678, b: 0x12340000 => __vseteq2 output: 0x00010000

Questions:

  1. Is this a documentation error in the CUDA Math API Guide, or an unexpected compiler intrinsic implementation?

  2. If this is a documentation error, can we officially rely on the packed per-halfword return pattern (0x00010001, 0x00010000, etc.) in future CUDA Toolkit releases without risking breaking changes?

2

Thanks for your report. This is a documentation issue. The result is intended to be per-component “boolean”, so current behavior is not going to change. Note that we also have __vcmp* SIMD intrinsics, for which the result is per-component bit-mask.