codekingpro/portable-devtools
114k
1/*
2 * Copyright 2022 NVIDIA Corporation. All rights reserved.
3 *
4 * NOTICE TO LICENSEE:
5 *
6 * This source code and/or documentation ("Licensed Deliverables") are
7 * subject to NVIDIA intellectual property rights under U.S. and
8 * international Copyright laws.
9 *
10 * These Licensed Deliverables contained herein is PROPRIETARY and
11 * CONFIDENTIAL to NVIDIA and is being provided under the terms and
12 * conditions of a form of NVIDIA software license agreement by and
13 * between NVIDIA and Licensee ("License Agreement") or electronically
14 * accepted by Licensee. Notwithstanding any terms or conditions to
15 * the contrary in the License Agreement, reproduction or disclosure
16 * of the Licensed Deliverables to any third party without the express
17 * written consent of NVIDIA is prohibited.
18 *
19 * NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE
20 * LICENSE AGREEMENT, NVIDIA MAKES NO REPRESENTATION ABOUT THE
21 * SUITABILITY OF THESE LICENSED DELIVERABLES FOR ANY PURPOSE. IT IS
22 * PROVIDED "AS IS" WITHOUT EXPRESS OR IMPLIED WARRANTY OF ANY KIND.
23 * NVIDIA DISCLAIMS ALL WARRANTIES WITH REGARD TO THESE LICENSED
24 * DELIVERABLES, INCLUDING ALL IMPLIED WARRANTIES OF MERCHANTABILITY,
25 * NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE.
26 * NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE
27 * LICENSE AGREEMENT, IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY
28 * SPECIAL, INDIRECT, INCIDENTAL, OR CONSEQUENTIAL DAMAGES, OR ANY
29 * DAMAGES WHATSOEVER RESULTING FROM LOSS OF USE, DATA OR PROFITS,
30 * WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE OR OTHER TORTIOUS
31 * ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE OR PERFORMANCE
32 * OF THESE LICENSED DELIVERABLES.
33 *
34 * U.S. Government End Users. These Licensed Deliverables are a
35 * "commercial item" as that term is defined at 48 C.F.R. 2.101 (OCT
36 * 1995), consisting of "commercial computer software" and "commercial
37 * computer software documentation" as such terms are used in 48
38 * C.F.R. 12.212 (SEPT 1995) and is provided to the U.S. Government
39 * only as a commercial end item. Consistent with 48 C.F.R.12.212 and
40 * 48 C.F.R. 227.7202-1 through 227.7202-4 (JUNE 1995), all
41 * U.S. Government End Users acquire the Licensed Deliverables with
42 * only those rights set forth herein.
43 *
44 * Any use of the Licensed Deliverables in individual and commercial
45 * software must include, in the user documentation and internal
46 * comments to the code, the above Disclaimer and U.S. Government End
47 * Users Notice.
48 */
49
50#ifndef __CUDA_FP8_H__
51#define __CUDA_FP8_H__
52
53/* Set up function decorations */
54#if defined(__CUDACC__)
55#define __CUDA_FP8_DECL__ static __device__ __inline__
56#define __CUDA_HOSTDEVICE_FP8__ __host__ __device__
57#define __CUDA_HOSTDEVICE_FP8_DECL__ static __host__ __device__ __inline__
58#else /* !defined(__CUDACC__) */
59#if defined(__GNUC__)
60#define __CUDA_HOSTDEVICE_FP8_DECL__ static __attribute__((unused))
61#else
62#define __CUDA_HOSTDEVICE_FP8_DECL__ static
63#endif /* defined(__GNUC__) */
64#define __CUDA_HOSTDEVICE_FP8__
65#endif /* defined(__CUDACC_) */
66
67#if !defined(_MSC_VER) && __cplusplus >= 201103L
68#define __CPP_VERSION_AT_LEAST_11_FP8
69#elif _MSC_FULL_VER >= 190024210 && _MSVC_LANG >= 201103L
70#define __CPP_VERSION_AT_LEAST_11_FP8
71#endif
72
73/* bring in __half_raw data type */
74#include "cuda_fp16.h"
75/* bring in __nv_bfloat16_raw data type */
76#include "cuda_bf16.h"
77/* bring in float2, double4, etc vector types */
78#include "vector_types.h"
79
80/**
81 * \defgroup CUDA_MATH_INTRINSIC_FP8 FP8 Intrinsics
82 * This section describes fp8 intrinsic functions.
83 * To use these functions, include the header file \p cuda_fp8.h in your
84 * program.
85 * The following macros are available to help users selectively enable/disable
86 * various definitions present in the header file:
87 * - \p __CUDA_NO_FP8_CONVERSIONS__ - If defined, this macro will prevent any
88 * use of the C++ type conversions (converting constructors and conversion
89 * operators) defined in the header.
90 * - \p __CUDA_NO_FP8_CONVERSION_OPERATORS__ - If defined, this macro will
91 * prevent any use of the C++ conversion operators from \p fp8 to other types.
92 */
93
94/**
95 * \defgroup CUDA_MATH_FP8_MISC FP8 Conversion and Data Movement
96 * \ingroup CUDA_MATH_INTRINSIC_FP8
97 * To use these functions, include the header file \p cuda_fp8.h in your
98 * program.
99 */
100
101/**
102 * \ingroup CUDA_MATH_FP8_MISC
103 * \brief 8-bit \p unsigned \p integer
104 * type abstraction used to for \p fp8 floating-point
105 * numbers storage.
106 */
107typedef unsigned char __nv_fp8_storage_t;
108
109/**
110 * \ingroup CUDA_MATH_FP8_MISC
111 * \brief 16-bit \p unsigned \p integer
112 * type abstraction used to for storage of pairs of
113 * \p fp8 floating-point numbers.
114 */
115typedef unsigned short int __nv_fp8x2_storage_t;
116
117/**
118 * \ingroup CUDA_MATH_FP8_MISC
119 * \brief 32-bit \p unsigned \p integer
120 * type abstraction used to for storage of tetrads of
121 * \p fp8 floating-point numbers.
122 */
123typedef unsigned int __nv_fp8x4_storage_t;
124
125/**
126 * \ingroup CUDA_MATH_FP8_MISC
127 * \brief Enumerates the modes applicable when
128 * performing a narrowing conversion to \p fp8 destination types.
129 */
130typedef enum __nv_saturation_t {
131 /**
132 * Means no saturation to finite is performed when conversion
133 * results in rounding values outside the range of destination
134 * type.
135 * NOTE: for fp8 type of e4m3 kind, the results that are larger
136 * than the maximum representable finite number of the target
137 * format become NaN.
138 */
139 __NV_NOSAT,
140 /**
141 * Means input larger than the maximum representable
142 * finite number MAXNORM of the target format round to the
143 * MAXNORM of the same sign as input.
144 */
145 __NV_SATFINITE,
146} __nv_saturation_t;
147
148/**
149 * \ingroup CUDA_MATH_FP8_MISC
150 * \brief Enumerates the possible
151 * interpretations of the 8-bit values when referring to them as
152 * \p fp8 types.
153 */
154typedef enum __nv_fp8_interpretation_t {
155 __NV_E4M3, /**< Stands for \p fp8 numbers of \p e4m3 kind. */
156 __NV_E5M2, /**< Stands for \p fp8 numbers of \p e5m2 kind. */
157} __nv_fp8_interpretation_t;
158
159/* Forward-declaration of C-style APIs */
160
161/**
162 * \ingroup CUDA_MATH_FP8_MISC
163 * \brief Converts input \p double precision \p x to \p fp8 type of the
164 * requested kind using round-to-nearest-even rounding and requested saturation
165 * mode.
166 *
167 * \details Converts input \p x to \p fp8 type of the kind specified by
168 * \p fp8_interpretation parameter,
169 * using round-to-nearest-even rounding and
170 * saturation mode specified by \p saturate parameter.
171 *
172 * \returns
173 * - The \p __nv_fp8_storage_t value holds the result of conversion.
174 */
175__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8_storage_t
176__nv_cvt_double_to_fp8(const double x, const __nv_saturation_t saturate,
177 const __nv_fp8_interpretation_t fp8_interpretation);
178
179/**
180 * \ingroup CUDA_MATH_FP8_MISC
181 * \brief Converts input vector of two \p double precision numbers packed
182 * in \p double2 \p x into a vector of two values of \p fp8 type of
183 * the requested kind using round-to-nearest-even rounding and requested
184 * saturation mode.
185 *
186 * \details Converts input vector \p x to a vector of two \p fp8 values of the
187 * kind specified by \p fp8_interpretation parameter, using
188 * round-to-nearest-even rounding and saturation mode specified by \p saturate
189 * parameter.
190 *
191 * \returns
192 * - The \p __nv_fp8x2_storage_t value holds the result of conversion.
193 */
194__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8x2_storage_t
195__nv_cvt_double2_to_fp8x2(const double2 x, const __nv_saturation_t saturate,
196 const __nv_fp8_interpretation_t fp8_interpretation);
197
198/**
199 * \ingroup CUDA_MATH_FP8_MISC
200 * \brief Converts input \p single precision \p x to \p fp8 type of the
201 * requested kind using round-to-nearest-even rounding and requested saturation
202 * mode.
203 *
204 * \details Converts input \p x to \p fp8 type of the kind specified by
205 * \p fp8_interpretation parameter,
206 * using round-to-nearest-even rounding and
207 * saturation mode specified by \p saturate parameter.
208 *
209 * \returns
210 * - The \p __nv_fp8_storage_t value holds the result of conversion.
211 */
212__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8_storage_t
213__nv_cvt_float_to_fp8(const float x, const __nv_saturation_t saturate,
214 const __nv_fp8_interpretation_t fp8_interpretation);
215
216/**
217 * \ingroup CUDA_MATH_FP8_MISC
218 * \brief Converts input vector of two \p single precision numbers packed
219 * in \p float2 \p x into a vector of two values of \p fp8 type of
220 * the requested kind using round-to-nearest-even rounding and requested
221 * saturation mode.
222 *
223 * \details Converts input vector \p x to a vector of two \p fp8 values of the
224 * kind specified by \p fp8_interpretation parameter, using
225 * round-to-nearest-even rounding and saturation mode specified by \p saturate
226 * parameter.
227 *
228 * \returns
229 * - The \p __nv_fp8x2_storage_t value holds the result of conversion.
230 */
231__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8x2_storage_t
232__nv_cvt_float2_to_fp8x2(const float2 x, const __nv_saturation_t saturate,
233 const __nv_fp8_interpretation_t fp8_interpretation);
234
235/**
236 * \ingroup CUDA_MATH_FP8_MISC
237 * \brief Converts input \p half precision \p x to \p fp8 type of the requested
238 * kind using round-to-nearest-even rounding and requested saturation mode.
239 *
240 * \details Converts input \p x to \p fp8 type of the kind specified by
241 * \p fp8_interpretation parameter,
242 * using round-to-nearest-even rounding and
243 * saturation mode specified by \p saturate parameter.
244 *
245 * \returns
246 * - The \p __nv_fp8_storage_t value holds the result of conversion.
247 */
248__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8_storage_t
249__nv_cvt_halfraw_to_fp8(const __half_raw x, const __nv_saturation_t saturate,
250 const __nv_fp8_interpretation_t fp8_interpretation);
251
252/**
253 * \ingroup CUDA_MATH_FP8_MISC
254 * \brief Converts input vector of two \p half precision numbers packed
255 * in \p __half2_raw \p x into a vector of two values of \p fp8 type of
256 * the requested kind using round-to-nearest-even rounding and requested
257 * saturation mode.
258 *
259 * \details Converts input vector \p x to a vector of two \p fp8 values of the
260 * kind specified by \p fp8_interpretation parameter, using
261 * round-to-nearest-even rounding and saturation mode specified by \p saturate
262 * parameter.
263 *
264 * \returns
265 * - The \p __nv_fp8x2_storage_t value holds the result of conversion.
266 */
267__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8x2_storage_t __nv_cvt_halfraw2_to_fp8x2(
268 const __half2_raw x, const __nv_saturation_t saturate,
269 const __nv_fp8_interpretation_t fp8_interpretation);
270
271/**
272 * \ingroup CUDA_MATH_FP8_MISC
273 * \brief Converts input \p nv_bfloat16 precision \p x to \p fp8 type of the
274 * requested kind using round-to-nearest-even rounding and requested saturation
275 * mode.
276 *
277 * \details Converts input \p x to \p fp8 type of the kind specified by
278 * \p fp8_interpretation parameter,
279 * using round-to-nearest-even rounding and
280 * saturation mode specified by \p saturate parameter.
281 *
282 * \returns
283 * - The \p __nv_fp8_storage_t value holds the result of conversion.
284 */
285__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8_storage_t __nv_cvt_bfloat16raw_to_fp8(
286 const __nv_bfloat16_raw x, const __nv_saturation_t saturate,
287 const __nv_fp8_interpretation_t fp8_interpretation);
288
289/**
290 * \ingroup CUDA_MATH_FP8_MISC
291 * \brief Converts input vector of two \p nv_bfloat16 precision numbers packed
292 * in \p __nv_bfloat162_raw \p x into a vector of two values of \p fp8 type of
293 * the requested kind using round-to-nearest-even rounding and requested
294 * saturation mode.
295 *
296 * \details Converts input vector \p x to a vector of two \p fp8 values of the
297 * kind specified by \p fp8_interpretation parameter, using
298 * round-to-nearest-even rounding and saturation mode specified by \p saturate
299 * parameter.
300 *
301 * \returns
302 * - The \p __nv_fp8x2_storage_t value holds the result of conversion.
303 */
304__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8x2_storage_t
305__nv_cvt_bfloat16raw2_to_fp8x2(
306 const __nv_bfloat162_raw x, const __nv_saturation_t saturate,
307 const __nv_fp8_interpretation_t fp8_interpretation);
308
309/**
310 * \ingroup CUDA_MATH_FP8_MISC
311 * \brief Converts input \p fp8 \p x of the specified kind
312 * to \p half precision.
313 *
314 * \details Converts input \p x of \p fp8 type of the kind specified by
315 * \p fp8_interpretation parameter
316 * to \p half precision.
317 *
318 * \returns
319 * - The \p __half_raw value holds the result of conversion.
320 */
321__CUDA_HOSTDEVICE_FP8_DECL__ __half_raw
322__nv_cvt_fp8_to_halfraw(const __nv_fp8_storage_t x,
323 const __nv_fp8_interpretation_t fp8_interpretation);
324/**
325 * \ingroup CUDA_MATH_FP8_MISC
326 * \brief Converts input vector of two \p fp8 values of the specified kind
327 * to a vector of two \p half precision values packed in \p __half2_raw
328 * structure.
329 *
330 * \details Converts input vector \p x of \p fp8 type of the kind specified by
331 * \p fp8_interpretation parameter
332 * to a vector of two \p half precision values and returns as \p __half2_raw
333 * structure.
334 *
335 * \returns
336 * - The \p __half2_raw value holds the result of conversion.
337 */
338__CUDA_HOSTDEVICE_FP8_DECL__ __half2_raw
339__nv_cvt_fp8x2_to_halfraw2(const __nv_fp8x2_storage_t x,
340 const __nv_fp8_interpretation_t fp8_interpretation);
341
342#if defined(__cplusplus)
343
344#define __CUDA_FP8_TYPES_EXIST__
345
346/* Forward-declaration of structures defined in "cuda_fp8.hpp" */
347struct __nv_fp8_e5m2;
348struct __nv_fp8x2_e5m2;
349struct __nv_fp8x4_e5m2;
350
351struct __nv_fp8_e4m3;
352struct __nv_fp8x2_e4m3;
353struct __nv_fp8x4_e4m3;
354
355#endif /* defined(__cplusplus) */
356
357#include "cuda_fp8.hpp"
358
359#undef __CUDA_FP8_DECL__
360#undef __CUDA_HOSTDEVICE_FP8__
361#undef __CUDA_HOSTDEVICE_FP8_DECL__
362
363#if defined(__CPP_VERSION_AT_LEAST_11_FP8)
364#undef __CPP_VERSION_AT_LEAST_11_FP8
365#endif /* defined(__CPP_VERSION_AT_LEAST_11_FP8) */
366
367#endif /* end of include guard: __CUDA_FP8_H__ */
368