fp16 / bf16 surrogate#

fp16 / bf16 surrogate (sim only) — stands in for <cuda_fp16.h> / <cuda_bf16.h> Values are STORED as float and rounded through the real 16-bit format on every assignment and every arithmetic result, so precision loss is modelled exactly while the arithmetic itself runs in float. Under nvcc the genuine headers are used instead (a ported test guards its own #include with #ifdef __CUDACC__). NOT modelled: the fast-math (“__h*_rn” rounding-mode) variants beyond round-to-nearest, and the __half2_raw/__nv_bfloat162_raw pair unions.

Defined in code/simulator.h.

This page documents 41 symbols: 3 macros, 38 functions.

Macros#

Signature

Description

Availability

Location

Example

SIM_HALF2_OPS(H2, H)

Expands to inline H2 operator+(H2 a, H2 b) { return H2(H(float(a.x) + float(b.x)), H(float(a.y) + float(b.y))); } inline H2 operator-(H2 a, H2 b) { return H2(H(float(a.x) - float(b.x)), H(flo

sim only

simulator.h:3381

SIM_HALF_FN(H, H2, SUF)

Expands to inline H __hadd##SUF(H a, H b) { return H(float(a) + float(b)); } inline H __hsub##SUF(H a, H b) { return H(float(a) - float(b)); } inline H __hmul##SUF(H a, H b) { return H(float(

sim only

simulator.h:3419

SIM_HALF_OPS(H)

Expands to inline H operator+(H a, H b) { return H(float(a) + float(b)); } inline H operator-(H a, H b) { return H(float(a) - float(b)); } inline H operator*(H a, H b) { return H(float(a) * f

sim only

simulator.h:3346

Functions#

Signature

Description

Availability

Location

Example

inline float2 __bfloat1622float2(__nv_bfloat162 b)

sim only

simulator.h:3415

inline float __bfloat162float(__nv_bfloat16 b)

sim only

simulator.h:3412

inline __half __double2half(double d)

sim only

simulator.h:3395

inline __nv_bfloat16 __float2bfloat16(float f)

sim only

simulator.h:3410

inline __nv_bfloat162 __float2bfloat162_rn(float a)

sim only

simulator.h:3414

inline __nv_bfloat16 __float2bfloat16_rn(float f)

sim only

simulator.h:3411

inline __half __float2half(float f)

sim only

simulator.h:3392

View

inline __half2 __float2half2_rn(float a)

sim only

simulator.h:3399

View

inline __half __float2half_rn(float f)

sim only

simulator.h:3393

inline __nv_bfloat162 __floats2bfloat162_rn(float a, float b)

sim only

simulator.h:3413

inline __half2 __floats2half2_rn(float a, float b)

sim only

simulator.h:3398

inline __nv_bfloat16 __hadd(__nv_bfloat16 a, __nv_bfloat16 b)

sim only

simulator.h:3441

inline float2 __half22float2(__half2 h)

sim only

simulator.h:3404

View

inline float __half2float(__half h)

sim only

simulator.h:3394

View

inline int __half2int_rn(__half h)

sim only

simulator.h:3396

operator __half_raw() const

sim only

simulator.h:3331

inline __nv_bfloat162 __halves2bfloat162(__nv_bfloat16 a, __nv_bfloat16 b)

sim only

simulator.h:3401

View

inline __half2 __halves2half2(__half a, __half b)

sim only

simulator.h:3400

View

inline __nv_bfloat16 __hfma(__nv_bfloat16 a, __nv_bfloat16 b, __nv_bfloat16 c)

sim only

simulator.h:3444

View

inline __nv_bfloat16 __high2bfloat16(__nv_bfloat162 b)

sim only

simulator.h:3417

inline float __high2float(__half2 h)

sim only

simulator.h:3408

View

inline __half __high2half(__half2 h)

sim only

simulator.h:3406

inline __nv_bfloat16 __hmul(__nv_bfloat16 a, __nv_bfloat16 b)

sim only

simulator.h:3443

inline __nv_bfloat16 __hsub(__nv_bfloat16 a, __nv_bfloat16 b)

sim only

simulator.h:3442

inline __half __int2half_rn(int i)

sim only

simulator.h:3397

inline __nv_bfloat16 __low2bfloat16(__nv_bfloat162 b)

sim only

simulator.h:3416

inline float __low2float(__half2 h)

sim only

simulator.h:3407

View

inline __half __low2half(__half2 h)

sim only

simulator.h:3405

operator __nv_bfloat16_raw() const

sim only

simulator.h:3341

inline float bf16_to_f32(unsigned short b)

sim only

simulator.h:3306

inline float f16_to_f32(unsigned short h)

sim only

simulator.h:3278

inline unsigned short f32_to_bf16(float f)

bfloat16 is the top 16 bits of a float, round-to-nearest-even.

sim only

simulator.h:3300

operator float() const

sim only

simulator.h:3330

View

inline __half hexp (__half a)

sim only

simulator.h:3448

inline __half hlog (__half a)

sim only

simulator.h:3449

inline __half hsqrt(__half a)

sim only

simulator.h:3447

inline float roundBf16(float f)

sim only

simulator.h:3310

inline float roundHalf(float f)

sim only

simulator.h:3298

Examples#

Code from the repository’s example corpus (code\*.cu); the highlighted line is the call site. These blocks are what the View links above open.

__float2half (test_web_ndb_haxpy.cu:72)
62    half2 *x2 = (half2*)x, *y2 = (half2*)y;
63
64    for (int i = start; i < n2; i += stride)
65        y2[i] = __hfma2(__halves2half2(a, a), x2[i], y2[i]);
66
67    // first thread handles singleton for odd arrays
68    if (start == 0 && (n % 2))
69        y[n - 1] = __hfma(a, x[n - 1], y[n - 1]);
70#else
71    for (int i = start; i < n; i += stride) {
72        y[i] = __float2half(__half2float(a) * __half2float(x[i]) + __half2float(y[i]));
73    }
74#endif
75}
76
77int main() {
78    int deviceCount = 0;
79    cudaError_t e = cudaGetDeviceCount(&deviceCount);
80    if (e != cudaSuccess || deviceCount == 0) {
81        std::printf("No CUDA device available (error %d: %s).\n"
82                    "Compilation and toolkit are fine; kernels need an NVIDIA GPU.\n",
83                    static_cast<int>(e), cudaGetErrorString(e));
__float2half2_rn (test_web_cs_fp16dot.cu:80)
70    __syncthreads();
71    if (threadIdx.x < 1)  v[threadIdx.x] = v[threadIdx.x] + v[threadIdx.x + 1];
72    __syncthreads();
73}
74
75__global__ void scalarProductKernel_intrinsics(half2 const* const a, half2 const* const b,
76                                               float* const results, size_t const size) {
77    const int        stride = gridDim.x * blockDim.x;
78    __shared__ half2 shArray[NUM_OF_THREADS];
79
80    shArray[threadIdx.x] = __float2half2_rn(0.f);
81    half2 value          = __float2half2_rn(0.f);
82
83    for (int i = threadIdx.x + blockDim.x * blockIdx.x; i < (int)size; i += stride)
84        value = __hfma2(a[i], b[i], value);
85
86    shArray[threadIdx.x] = value;
87    __syncthreads();
88    reduceInShared_intrinsics(shArray);
89
90    if (threadIdx.x == 0) {
91        half2 result        = shArray[0];
__half22float2 (test_web_pk_mixed_precision.cu:88)
78    half2*  dual_y = reinterpret_cast<half2*>(d_y);
79    float2* dual_z = reinterpret_cast<float2*>(d_z);
80
81    SHARED_EXTERN(float2, s_data);     // [SIM] was extern __shared__ float2 s_data[]
82    (void)s_data;                      // declared by the sample, never read
83
84#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 530
85    // [SIM] `!defined(__CUDA_ARCH__)` added so the CPU sim takes the half2 path
86    // too — otherwise the sim would only ever exercise the fallback.
87    for (int i = idx_x; i < size; i += stride) {
88        dual_z[i] = __half22float2(__hmul2(dual_y[i], dual_x[i]));
89    }
90#else
91    for (int i = idx_x; i < size; i += stride) {
92        float2 a = __half22float2(dual_x[i]);
93        float2 b = __half22float2(dual_y[i]);
94        dual_z[i] = make_float2(a.x * b.x, a.y * b.y);   // [SIM] was helper_math's operator*
95    }
96#endif
97}
98
99// Host reference, shared by both paths (the sample had one per program).
__half2float (test_web_ndb_haxpy.cu:72)
62    half2 *x2 = (half2*)x, *y2 = (half2*)y;
63
64    for (int i = start; i < n2; i += stride)
65        y2[i] = __hfma2(__halves2half2(a, a), x2[i], y2[i]);
66
67    // first thread handles singleton for odd arrays
68    if (start == 0 && (n % 2))
69        y[n - 1] = __hfma(a, x[n - 1], y[n - 1]);
70#else
71    for (int i = start; i < n; i += stride) {
72        y[i] = __float2half(__half2float(a) * __half2float(x[i]) + __half2float(y[i]));
73    }
74#endif
75}
76
77int main() {
78    int deviceCount = 0;
79    cudaError_t e = cudaGetDeviceCount(&deviceCount);
80    if (e != cudaSuccess || deviceCount == 0) {
81        std::printf("No CUDA device available (error %d: %s).\n"
82                    "Compilation and toolkit are fine; kernels need an NVIDIA GPU.\n",
83                    static_cast<int>(e), cudaGetErrorString(e));
__halves2bfloat162 (test_llmc_layernorm_backward.cu:161)
151// ----------------------------------------------------------------------------
152// GPU kernels (verbatim except the [SIM] shared-memory spelling)
153
154// GPU helper functions for atomicAdd on smaller than 32-bit types
155#ifdef ENABLE_BF16
156__device__ void atomicAddX(__nv_bfloat16* addr, __nv_bfloat16 val) {
157    uintptr_t ptr_val = reinterpret_cast<uintptr_t>(addr);
158    __nv_bfloat162* ptr_bf16 = reinterpret_cast<__nv_bfloat162*>(ptr_val & ~uintptr_t(0x3));
159
160    // Prepare the value to add, setting the other half to zero
161    __nv_bfloat162 add_val = (ptr_val & 0x3) ? __halves2bfloat162(__ushort_as_bfloat16(0), val)
162                                             : __halves2bfloat162(val, __ushort_as_bfloat16(0));
163    atomicAdd(ptr_bf16, add_val);
164}
165#endif
166#ifdef ENABLE_FP16
167__device__ void atomicAddX(half* addr, half val) {
168    uintptr_t ptr_val = reinterpret_cast<uintptr_t>(addr);
169    half2* ptr_fp16 = reinterpret_cast<half2*>(ptr_val & ~uintptr_t(0x3));
170
171    // Prepare the value to add, setting the other half to zero
172    half2 add_val = (ptr_val & 0x3) ? __halves2half2(__ushort_as_half(0), val)
__halves2half2 (test_web_ndb_haxpy.cu:65)
55__global__ void haxpy(int n, __half a, const __half* x, __half* y) {
56    int start  = threadIdx.x + blockDim.x * blockIdx.x;
57    int stride = blockDim.x * gridDim.x;
58
59#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 530
60    int n2 = n / 2;
61    half2 *x2 = (half2*)x, *y2 = (half2*)y;
62
63    for (int i = start; i < n2; i += stride)
64        y2[i] = __hfma2(__halves2half2(a, a), x2[i], y2[i]);
65
66    // first thread handles singleton for odd arrays
67    if (start == 0 && (n % 2))
68        y[n - 1] = __hfma(a, x[n - 1], y[n - 1]);
69#else
70    for (int i = start; i < n; i += stride) {
71        y[i] = __float2half(__half2float(a) * __half2float(x[i]) + __half2float(y[i]));
72    }
73#endif
74}
__hfma (test_web_ndb_haxpy.cu:69)
59#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 530
60    int n2 = n / 2;
61    half2 *x2 = (half2*)x, *y2 = (half2*)y;
62
63    for (int i = start; i < n2; i += stride)
64        y2[i] = __hfma2(__halves2half2(a, a), x2[i], y2[i]);
65
66    // first thread handles singleton for odd arrays
67    if (start == 0 && (n % 2))
68        y[n - 1] = __hfma(a, x[n - 1], y[n - 1]);
69#else
70    for (int i = start; i < n; i += stride) {
71        y[i] = __float2half(__half2float(a) * __half2float(x[i]) + __half2float(y[i]));
72    }
73#endif
74}
75
76int main() {
77    int deviceCount = 0;
78    cudaError_t e = cudaGetDeviceCount(&deviceCount);
79    if (e != cudaSuccess || deviceCount == 0) {
__high2float (test_web_cs_fp16dot.cu:92)
 82    for (int i = threadIdx.x + blockDim.x * blockIdx.x; i < (int)size; i += stride)
 83        value = __hfma2(a[i], b[i], value);
 84
 85    shArray[threadIdx.x] = value;
 86    __syncthreads();
 87    reduceInShared_intrinsics(shArray);
 88
 89    if (threadIdx.x == 0) {
 90        half2 result        = shArray[0];
 91        float f_result      = __low2float(result) + __high2float(result);
 92        results[blockIdx.x] = f_result;
 93    }
 94}
 95
 96__global__ void scalarProductKernel_native(half2 const* const a, half2 const* const b,
 97                                           float* const results, size_t const size) {
 98    const int        stride = gridDim.x * blockDim.x;
 99    __shared__ half2 shArray[NUM_OF_THREADS];
100
101    half2 value(0.f, 0.f);
102    shArray[threadIdx.x] = value;
__low2float (test_web_cs_fp16dot.cu:92)
 82    for (int i = threadIdx.x + blockDim.x * blockIdx.x; i < (int)size; i += stride)
 83        value = __hfma2(a[i], b[i], value);
 84
 85    shArray[threadIdx.x] = value;
 86    __syncthreads();
 87    reduceInShared_intrinsics(shArray);
 88
 89    if (threadIdx.x == 0) {
 90        half2 result        = shArray[0];
 91        float f_result      = __low2float(result) + __high2float(result);
 92        results[blockIdx.x] = f_result;
 93    }
 94}
 95
 96__global__ void scalarProductKernel_native(half2 const* const a, half2 const* const b,
 97                                           float* const results, size_t const size) {
 98    const int        stride = gridDim.x * blockDim.x;
 99    __shared__ half2 shArray[NUM_OF_THREADS];
100
101    half2 value(0.f, 0.f);
102    shArray[threadIdx.x] = value;
float (test_web_cs_matmul_cublas.cu:59)
49} sMatrixSize;
50
51// Reference C = A * B on the host, row-major, verbatim from the sample.
52static void matrixMulCPU(float* C, const float* A, const float* B,
53                         unsigned int hA, unsigned int wA, unsigned int wB) {
54    for (unsigned int i = 0; i < hA; ++i)
55        for (unsigned int j = 0; j < wB; ++j) {
56            double sum = 0;
57            for (unsigned int k = 0; k < wA; ++k)
58                sum += double(A[i * wA + k]) * double(B[k * wB + j]);
59            C[i * wB + j] = float(sum);
60        }
61}
62
63// [SIM] stand-in for helper_functions.h sdkCompareL2fe: relative L2 error.
64static bool compareL2fe(const float* ref, const float* got, unsigned int n, float eps) {
65    double normRef = 0.0, normErr = 0.0;
66    for (unsigned int i = 0; i < n; ++i) {
67        const double d = double(ref[i]) - double(got[i]);
68        normRef += double(ref[i]) * double(ref[i]);
69        normErr += d * d;
70    }
On this page