Update docs 2025 04 14 (#54)
* Update docs 2025 03 31
- Docs: remove virtual_rocr.rst
- Fix documentation warnings
- Reformat HIP RTC
- Docs: Refactor HIP porting guide
- Docs: Expand HIP porting guide and CUDA driver porting guide
- Minor fix
- Docs: Update environment variables file
- Bump rocm-docs-core[api_reference] from 1.15.0 to 1.17.0 in /docs/sphinx
- Docs: Update FP8 page to show both FP8 and FP16 types
- Bump sphinxcontrib-doxylink from 1.12.4 to 1.13.0 in /docs/sphinx
- Bumps [rocm-docs-core[api_reference]](https://github.com/ROCm/rocm-docs-core) from 1.17.0 to 1.17.1.
- Remove external link
- Update programming model
- Bump rocm-docs-core[api_reference] from 1.17.1 to 1.18.1 in /docs/sphinx
- Docs: Add page for Complex Math API
- Docs: Add page about HIP error codes
- Update docs: the compilation cache is enabled by default
- Fix fns32 function mask type in doc
* Bump rocm-docs-core[api_reference] from 1.18.1 to 1.18.2 in /docs/sphinx
Bumps [rocm-docs-core[api_reference]](https://github.com/ROCm/rocm-docs-core) from 1.18.1 to 1.18.2.
- [Release notes](https://github.com/ROCm/rocm-docs-core/releases)
- [Changelog](https://github.com/ROCm/rocm-docs-core/blob/develop/CHANGELOG.md)
- [Commits](https://github.com/ROCm/rocm-docs-core/compare/v1.18.1...v1.18.2)
---
updated-dependencies:
- dependency-name: rocm-docs-core[api_reference]
dependency-version: 1.18.2
dependency-type: direct:production
update-type: version-update:semver-patch
* Fix readme link
* Docs: Fix verbose paths generated by doxygen
* Handle git ssh in docs conf.py
[ROCm/hip commit: d0cf32a63a]
This commit is contained in:
Kiss, Istvan
GitHub
@@ -0,0 +1,446 @@
|
||||
.. meta::
|
||||
:description: This chapter describes the complex math functions that are accessible in HIP.
|
||||
:keywords: AMD, ROCm, HIP, CUDA, complex math functions, HIP complex math functions
|
||||
|
||||
.. _complex_math_api_reference:
|
||||
|
||||
********************************************************************************
|
||||
HIP complex math API
|
||||
********************************************************************************
|
||||
|
||||
HIP provides built-in support for complex number operations through specialized types and functions,
|
||||
available for both single-precision (float) and double-precision (double) calculations. All complex types
|
||||
and functions are available on both host and device.
|
||||
|
||||
For any complex number ``z``, the form is:
|
||||
|
||||
.. math::
|
||||
|
||||
z = x + yi
|
||||
|
||||
where ``x`` is the real part and ``y`` is the imaginary part.
|
||||
|
||||
Complex Number Types
|
||||
====================
|
||||
|
||||
A brief overview of the specialized data types used to represent complex numbers in HIP, available
|
||||
in both single and double precision formats.
|
||||
|
||||
.. list-table::
|
||||
:header-rows: 1
|
||||
:widths: 40 60
|
||||
|
||||
* - Type
|
||||
- Description
|
||||
|
||||
* - ``hipFloatComplex``
|
||||
- | Complex number using single-precision (float) values
|
||||
| (note: ``hipComplex`` is an alias of ``hipFloatComplex``)
|
||||
|
||||
* - ``hipDoubleComplex``
|
||||
- Complex number using double-precision (double) values
|
||||
|
||||
Complex Number Functions
|
||||
========================
|
||||
|
||||
A comprehensive collection of functions for creating and manipulating complex numbers, organized by
|
||||
functional categories for easy reference.
|
||||
|
||||
Type Construction
|
||||
-----------------
|
||||
|
||||
Functions for creating complex number objects and extracting their real and imaginary components.
|
||||
|
||||
.. tab-set::
|
||||
|
||||
.. tab-item:: Single Precision
|
||||
|
||||
.. list-table::
|
||||
:header-rows: 1
|
||||
:widths: 40 60
|
||||
|
||||
* - Function
|
||||
- Description
|
||||
|
||||
* - | ``hipFloatComplex``
|
||||
| ``make_hipFloatComplex(``
|
||||
| ``float a,``
|
||||
| ``float b``
|
||||
| ``)``
|
||||
- | Creates a complex number
|
||||
| (note: ``make_hipComplex`` is an alias of ``make_hipFloatComplex``)
|
||||
| :math:`z = a + bi`
|
||||
|
||||
* - | ``float``
|
||||
| ``hipCrealf(``
|
||||
| ``hipFloatComplex z``
|
||||
| ``)``
|
||||
- | Returns real part of z
|
||||
| :math:`\Re(z) = x`
|
||||
|
||||
* - | ``float``
|
||||
| ``hipCimagf(``
|
||||
| ``hipFloatComplex z``
|
||||
| ``)``
|
||||
- | Returns imaginary part of z
|
||||
| :math:`\Im(z) = y`
|
||||
|
||||
.. tab-item:: Double Precision
|
||||
|
||||
.. list-table::
|
||||
:header-rows: 1
|
||||
:widths: 40 60
|
||||
|
||||
* - Function
|
||||
- Description
|
||||
|
||||
* - | ``hipDoubleComplex``
|
||||
| ``make_hipDoubleComplex(``
|
||||
| ``double a,``
|
||||
| ``double b``
|
||||
| ``)``
|
||||
- | Creates a complex number
|
||||
| :math:`z = a + bi`
|
||||
|
||||
* - | ``double``
|
||||
| ``hipCreal(``
|
||||
| ``hipDoubleComplex z``
|
||||
| ``)``
|
||||
- | Returns real part of z
|
||||
| :math:`\Re(z) = x`
|
||||
|
||||
* - | ``double``
|
||||
| ``hipCimag(``
|
||||
| ``hipDoubleComplex z``
|
||||
| ``)``
|
||||
- | Returns imaginary part of z
|
||||
| :math:`\Im(z) = y`
|
||||
|
||||
Basic Arithmetic
|
||||
----------------
|
||||
|
||||
Operations for performing standard arithmetic with complex numbers, including addition,
|
||||
subtraction, multiplication, division, and fused multiply-add.
|
||||
|
||||
.. tab-set::
|
||||
|
||||
.. tab-item:: Single Precision
|
||||
|
||||
.. list-table::
|
||||
:header-rows: 1
|
||||
:widths: 40 60
|
||||
|
||||
* - Function
|
||||
- Description
|
||||
|
||||
* - | ``hipFloatComplex``
|
||||
| ``hipCaddf(``
|
||||
| ``hipFloatComplex p,``
|
||||
| ``hipFloatComplex q``
|
||||
| ``)``
|
||||
- | Addition of two single-precision complex values
|
||||
| :math:`(a + bi) + (c + di) = (a + c) + (b + d)i`
|
||||
|
||||
* - | ``hipFloatComplex``
|
||||
| ``hipCsubf(``
|
||||
| ``hipFloatComplex p,``
|
||||
| ``hipFloatComplex q``
|
||||
| ``)``
|
||||
- | Subtraction of two single-precision complex values
|
||||
| :math:`(a + bi) - (c + di) = (a - c) + (b - d)i`
|
||||
|
||||
* - | ``hipFloatComplex``
|
||||
| ``hipCmulf(``
|
||||
| ``hipFloatComplex p,``
|
||||
| ``hipFloatComplex q``
|
||||
| ``)``
|
||||
- | Multiplication of two single-precision complex values
|
||||
| :math:`(a + bi)(c + di) = (ac - bd) + (bc + ad)i`
|
||||
|
||||
* - | ``hipFloatComplex``
|
||||
| ``hipCdivf(``
|
||||
| ``hipFloatComplex p,``
|
||||
| ``hipFloatComplex q``
|
||||
| ``)``
|
||||
- | Division of two single-precision complex values
|
||||
| :math:`\frac{a + bi}{c + di} = \frac{(ac + bd) + (bc - ad)i}{c^2 + d^2}`
|
||||
|
||||
* - | ``hipFloatComplex``
|
||||
| ``hipCfmaf(``
|
||||
| ``hipComplex p,``
|
||||
| ``hipComplex q,``
|
||||
| ``hipComplex r``
|
||||
| ``)``
|
||||
- | Fused multiply-add of three single-precision complex values
|
||||
| :math:`(a + bi)(c + di) + (e + fi)`
|
||||
|
||||
.. tab-item:: Double Precision
|
||||
|
||||
.. list-table::
|
||||
:header-rows: 1
|
||||
:widths: 40 60
|
||||
|
||||
* - Function
|
||||
- Description
|
||||
|
||||
* - | ``hipDoubleComplex``
|
||||
| ``hipCadd(``
|
||||
| ``hipDoubleComplex p,``
|
||||
| ``hipDoubleComplex q``
|
||||
| ``)``
|
||||
- | Addition of two double-precision complex values
|
||||
| :math:`(a + bi) + (c + di) = (a + c) + (b + d)i`
|
||||
|
||||
* - | ``hipDoubleComplex``
|
||||
| ``hipCsub(``
|
||||
| ``hipDoubleComplex p,``
|
||||
| ``hipDoubleComplex q``
|
||||
| ``)``
|
||||
- | Subtraction of two double-precision complex values
|
||||
| :math:`(a + bi) - (c + di) = (a - c) + (b - d)i`
|
||||
|
||||
* - | ``hipDoubleComplex``
|
||||
| ``hipCmul(``
|
||||
| ``hipDoubleComplex p,``
|
||||
| ``hipDoubleComplex q``
|
||||
| ``)``
|
||||
- | Multiplication of two double-precision complex values
|
||||
| :math:`(a + bi)(c + di) = (ac - bd) + (bc + ad)i`
|
||||
|
||||
* - | ``hipDoubleComplex``
|
||||
| ``hipCdiv(``
|
||||
| ``hipDoubleComplex p,``
|
||||
| ``hipDoubleComplex q``
|
||||
| ``)``
|
||||
- | Division of two double-precision complex values
|
||||
| :math:`\frac{a + bi}{c + di} = \frac{(ac + bd) + (bc - ad)i}{c^2 + d^2}`
|
||||
|
||||
* - | ``hipDoubleComplex``
|
||||
| ``hipCfma(``
|
||||
| ``hipDoubleComplex p,``
|
||||
| ``hipDoubleComplex q,``
|
||||
| ``hipDoubleComplex r``
|
||||
| ``)``
|
||||
- | Fused multiply-add of three double-precision complex values
|
||||
| :math:`(a + bi)(c + di) + (e + fi)`
|
||||
|
||||
Complex Operations
|
||||
------------------
|
||||
|
||||
Functions for complex-specific calculations, including conjugate determination and magnitude
|
||||
(absolute value) computation.
|
||||
|
||||
.. tab-set::
|
||||
|
||||
.. tab-item:: Single Precision
|
||||
|
||||
.. list-table::
|
||||
:header-rows: 1
|
||||
:widths: 40 60
|
||||
|
||||
* - Function
|
||||
- Description
|
||||
|
||||
* - | ``hipFloatComplex``
|
||||
| ``hipConjf(``
|
||||
| ``hipFloatComplex z``
|
||||
| ``)``
|
||||
- | Complex conjugate
|
||||
| :math:`\overline{a + bi} = a - bi`
|
||||
|
||||
* - | ``float``
|
||||
| ``hipCabsf(``
|
||||
| ``hipFloatComplex z``
|
||||
| ``)``
|
||||
- | Absolute value (magnitude)
|
||||
| :math:`|a + bi| = \sqrt{a^2 + b^2}`
|
||||
|
||||
* - | ``float``
|
||||
| ``hipCsqabsf(``
|
||||
| ``hipFloatComplex z``
|
||||
| ``)``
|
||||
- | Squared absolute value
|
||||
| :math:`|a + bi|^2 = a^2 + b^2`
|
||||
|
||||
.. tab-item:: Double Precision
|
||||
|
||||
.. list-table::
|
||||
:header-rows: 1
|
||||
:widths: 40 60
|
||||
|
||||
* - Function
|
||||
- Description
|
||||
|
||||
* - | ``hipDoubleComplex``
|
||||
| ``hipConj(``
|
||||
| ``hipDoubleComplex z``
|
||||
| ``)``
|
||||
- | Complex conjugate
|
||||
| :math:`\overline{a + bi} = a - bi`
|
||||
|
||||
* - | ``double``
|
||||
| ``hipCabs(``
|
||||
| ``hipDoubleComplex z``
|
||||
| ``)``
|
||||
- | Absolute value (magnitude)
|
||||
| :math:`|a + bi| = \sqrt{a^2 + b^2}`
|
||||
|
||||
* - | ``double``
|
||||
| ``hipCsqabs(``
|
||||
| ``hipDoubleComplex z``
|
||||
| ``)``
|
||||
- | Squared absolute value
|
||||
| :math:`|a + bi|^2 = a^2 + b^2`
|
||||
|
||||
Type Conversion
|
||||
---------------
|
||||
|
||||
Utility functions for conversion between single-precision and double-precision complex number formats.
|
||||
|
||||
.. list-table::
|
||||
:header-rows: 1
|
||||
:widths: 40 60
|
||||
|
||||
* - Function
|
||||
- Description
|
||||
|
||||
* - | ``hipFloatComplex``
|
||||
| ``hipComplexDoubleToFloat(``
|
||||
| ``hipDoubleComplex z``
|
||||
| ``)``
|
||||
- Converts double-precision to single-precision complex
|
||||
|
||||
* - | ``hipDoubleComplex``
|
||||
| ``hipComplexFloatToDouble(``
|
||||
| ``hipFloatComplex z``
|
||||
| ``)``
|
||||
- Converts single-precision to double-precision complex
|
||||
|
||||
Example Usage
|
||||
=============
|
||||
|
||||
The following example demonstrates using complex numbers to compute the Discrete Fourier Transform (DFT)
|
||||
of a simple signal on the GPU. The DFT converts a signal from the time domain to the frequency domain.
|
||||
The kernel function ``computeDFT`` shows various HIP complex math operations in action:
|
||||
|
||||
* Creating complex numbers with ``make_hipFloatComplex``
|
||||
* Performing complex multiplication with ``hipCmulf``
|
||||
* Accumulating complex values with ``hipCaddf``
|
||||
|
||||
The example also demonstrates proper use of complex number handling on both host and device, including
|
||||
memory allocation, transfer, and validation of results between CPU and GPU implementations.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <hip/hip_complex.h>
|
||||
#include <iostream>
|
||||
#include <vector>
|
||||
#include <cmath>
|
||||
|
||||
#define HIP_CHECK(expression) \
|
||||
{ \
|
||||
const hipError_t err = expression; \
|
||||
if (err != hipSuccess) { \
|
||||
std::cerr << "HIP error: " \
|
||||
<< hipGetErrorString(err) \
|
||||
<< " at " << __LINE__ << "\n"; \
|
||||
exit(EXIT_FAILURE); \
|
||||
} \
|
||||
}
|
||||
|
||||
// Kernel to compute DFT
|
||||
__global__ void computeDFT(const float* input,
|
||||
hipFloatComplex* output,
|
||||
const int N)
|
||||
{
|
||||
int k = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (k >= N) return;
|
||||
|
||||
hipFloatComplex sum = make_hipFloatComplex(0.0f, 0.0f);
|
||||
|
||||
for (int n = 0; n < N; n++) {
|
||||
float angle = -2.0f * M_PI * k * n / N;
|
||||
hipFloatComplex w = make_hipFloatComplex(cosf(angle), sinf(angle));
|
||||
hipFloatComplex x = make_hipFloatComplex(input[n], 0.0f);
|
||||
sum = hipCaddf(sum, hipCmulf(x, w));
|
||||
}
|
||||
|
||||
output[k] = sum;
|
||||
}
|
||||
|
||||
// CPU implementation of DFT for verification
|
||||
std::vector<hipFloatComplex> cpuDFT(const std::vector<float>& input) {
|
||||
const int N = input.size();
|
||||
std::vector<hipFloatComplex> result(N);
|
||||
|
||||
for (int k = 0; k < N; k++) {
|
||||
hipFloatComplex sum = make_hipFloatComplex(0.0f, 0.0f);
|
||||
for (int n = 0; n < N; n++) {
|
||||
float angle = -2.0f * M_PI * k * n / N;
|
||||
hipFloatComplex w = make_hipFloatComplex(cosf(angle), sinf(angle));
|
||||
hipFloatComplex x = make_hipFloatComplex(input[n], 0.0f);
|
||||
sum = hipCaddf(sum, hipCmulf(x, w));
|
||||
}
|
||||
result[k] = sum;
|
||||
}
|
||||
return result;
|
||||
}
|
||||
|
||||
int main() {
|
||||
const int N = 256; // Signal length
|
||||
const int blockSize = 256;
|
||||
|
||||
// Generate input signal: sum of two sine waves
|
||||
std::vector<float> signal(N);
|
||||
for (int i = 0; i < N; i++) {
|
||||
float t = static_cast<float>(i) / N;
|
||||
signal[i] = sinf(2.0f * M_PI * 10.0f * t) + // 10 Hz component
|
||||
0.5f * sinf(2.0f * M_PI * 20.0f * t); // 20 Hz component
|
||||
}
|
||||
|
||||
// Compute reference solution on CPU
|
||||
std::vector<hipFloatComplex> cpu_output = cpuDFT(signal);
|
||||
|
||||
// Allocate device memory
|
||||
float* d_signal;
|
||||
hipFloatComplex* d_output;
|
||||
HIP_CHECK(hipMalloc(&d_signal, N * sizeof(float)));
|
||||
HIP_CHECK(hipMalloc(&d_output, N * sizeof(hipFloatComplex)));
|
||||
|
||||
// Copy input to device
|
||||
HIP_CHECK(hipMemcpy(d_signal, signal.data(), N * sizeof(float),
|
||||
hipMemcpyHostToDevice));
|
||||
|
||||
// Launch kernel
|
||||
dim3 grid((N + blockSize - 1) / blockSize);
|
||||
dim3 block(blockSize);
|
||||
computeDFT<<<grid, block>>>(d_signal, d_output, N);
|
||||
HIP_CHECK(hipGetLastError());
|
||||
|
||||
// Get GPU results
|
||||
std::vector<hipFloatComplex> gpu_output(N);
|
||||
HIP_CHECK(hipMemcpy(gpu_output.data(), d_output, N * sizeof(hipFloatComplex),
|
||||
hipMemcpyDeviceToHost));
|
||||
|
||||
// Verify results
|
||||
bool passed = true;
|
||||
const float tolerance = 1e-5f; // Adjust based on precision requirements
|
||||
|
||||
for (int i = 0; i < N; i++) {
|
||||
float diff_real = std::abs(hipCrealf(gpu_output[i]) - hipCrealf(cpu_output[i]));
|
||||
float diff_imag = std::abs(hipCimagf(gpu_output[i]) - hipCimagf(cpu_output[i]));
|
||||
|
||||
if (diff_real > tolerance || diff_imag > tolerance) {
|
||||
passed = false;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
std::cout << "DFT Verification: " << (passed ? "PASSED" : "FAILED") << "\n";
|
||||
|
||||
// Cleanup
|
||||
HIP_CHECK(hipFree(d_signal));
|
||||
HIP_CHECK(hipFree(d_output));
|
||||
return passed ? 0 : 1;
|
||||
}
|
||||
File diff suppressed because it is too large
Load Diff
@@ -1,230 +0,0 @@
|
||||
.. meta::
|
||||
:description: This page describes FP8 numbers present in HIP.
|
||||
:keywords: AMD, ROCm, HIP, fp8, fnuz, ocp
|
||||
|
||||
*******************************************************************************
|
||||
FP8 Numbers
|
||||
*******************************************************************************
|
||||
|
||||
`FP8 numbers <https://arxiv.org/pdf/2209.05433>`_ were introduced to accelerate deep learning inferencing. They provide higher throughput of matrix operations because the smaller size allows more of them in the available fixed memory.
|
||||
|
||||
HIP has two FP8 number representations called *FP8-OCP* and *FP8-FNUZ*.
|
||||
|
||||
Open Compute Project(OCP) number definition can be found `here <https://www.opencompute.org/documents/ocp-8-bit-floating-point-specification-ofp8-revision-1-0-2023-12-01-pdf-1>`_.
|
||||
|
||||
Definition of FNUZ: fnuz suffix means only finite and NaN values are supported. Unlike other types, Inf are not supported.
|
||||
NaN is when sign bit is set and all other exponent and mantissa bits are 0. All other values are finite.
|
||||
This provides one extra value of exponent and adds to the range of supported FP8 numbers.
|
||||
|
||||
FP8 Definition
|
||||
==============
|
||||
|
||||
FP8 numbers are composed of a sign, an exponent and a mantissa. Their sizes are dependent on the format.
|
||||
There are two formats of FP8 numbers, E4M3 and E5M2.
|
||||
|
||||
- E4M3: 1 bit sign, 4 bit exponent, 3 bit mantissa
|
||||
- E5M2: 1 bit sign, 5 bit exponent, 2 bit mantissa
|
||||
|
||||
HIP Header
|
||||
==========
|
||||
|
||||
The `HIP header <https://github.com/ROCm/clr/blob/develop/hipamd/include/hip/amd_detail/amd_hip_fp8.h>`_ defines the FP8 ocp/fnuz numbers.
|
||||
|
||||
Supported Devices
|
||||
=================
|
||||
|
||||
.. list-table:: Supported devices for fp8 numbers
|
||||
:header-rows: 1
|
||||
|
||||
* - Device Type
|
||||
- FNUZ FP8
|
||||
- OCP FP8
|
||||
* - Host
|
||||
- Yes
|
||||
- Yes
|
||||
* - gfx942
|
||||
- Yes
|
||||
- No
|
||||
* - gfx1200/gfx1201
|
||||
- No
|
||||
- Yes
|
||||
|
||||
Usage
|
||||
=====
|
||||
|
||||
To use the FP8 numbers inside HIP programs.
|
||||
|
||||
.. code-block:: c
|
||||
|
||||
#include <hip/hip_fp8.h>
|
||||
|
||||
FP8 numbers can be used on CPU side:
|
||||
|
||||
.. code-block:: c
|
||||
|
||||
__hip_fp8_storage_t convert_float_to_fp8(
|
||||
float in, /* Input val */
|
||||
__hip_fp8_interpretation_t interpret, /* interpretation of number E4M3/E5M2 */
|
||||
__hip_saturation_t sat /* Saturation behavior */
|
||||
) {
|
||||
return __hip_cvt_float_to_fp8(in, sat, interpret);
|
||||
}
|
||||
|
||||
The same can be done in kernels as well.
|
||||
|
||||
.. code-block:: c
|
||||
|
||||
__device__ __hip_fp8_storage_t d_convert_float_to_fp8(
|
||||
float in,
|
||||
__hip_fp8_interpretation_t interpret,
|
||||
__hip_saturation_t sat) {
|
||||
return __hip_cvt_float_to_fp8(in, sat, interpret);
|
||||
}
|
||||
|
||||
An important thing to note here is if you use this on gfx94x GPU, it will be fnuz number but on any other GPU it will be an OCP number.
|
||||
|
||||
The following code example does roundtrip FP8 conversions on both the CPU and GPU and compares the results.
|
||||
|
||||
.. code-block:: c
|
||||
|
||||
#include <hip/hip_fp8.h>
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <iostream>
|
||||
#include <vector>
|
||||
|
||||
#define hip_check(hip_call) \
|
||||
{ \
|
||||
auto hip_res = hip_call; \
|
||||
if (hip_res != hipSuccess) { \
|
||||
std::cerr << "Failed in hip call: " << #hip_call \
|
||||
<< " with error: " << hipGetErrorName(hip_res) << std::endl; \
|
||||
std::abort(); \
|
||||
} \
|
||||
}
|
||||
|
||||
__device__ __hip_fp8_storage_t d_convert_float_to_fp8(
|
||||
float in, __hip_fp8_interpretation_t interpret, __hip_saturation_t sat) {
|
||||
return __hip_cvt_float_to_fp8(in, sat, interpret);
|
||||
}
|
||||
|
||||
__device__ float d_convert_fp8_to_float(float in,
|
||||
__hip_fp8_interpretation_t interpret) {
|
||||
__half hf = __hip_cvt_fp8_to_halfraw(in, interpret);
|
||||
return hf;
|
||||
}
|
||||
|
||||
__global__ void float_to_fp8_to_float(float *in,
|
||||
__hip_fp8_interpretation_t interpret,
|
||||
__hip_saturation_t sat, float *out,
|
||||
size_t size) {
|
||||
int i = threadIdx.x;
|
||||
if (i < size) {
|
||||
auto fp8 = d_convert_float_to_fp8(in[i], interpret, sat);
|
||||
out[i] = d_convert_fp8_to_float(fp8, interpret);
|
||||
}
|
||||
}
|
||||
|
||||
__hip_fp8_storage_t
|
||||
convert_float_to_fp8(float in, /* Input val */
|
||||
__hip_fp8_interpretation_t
|
||||
interpret, /* interpretation of number E4M3/E5M2 */
|
||||
__hip_saturation_t sat /* Saturation behavior */
|
||||
) {
|
||||
return __hip_cvt_float_to_fp8(in, sat, interpret);
|
||||
}
|
||||
|
||||
float convert_fp8_to_float(
|
||||
__hip_fp8_storage_t in, /* Input val */
|
||||
__hip_fp8_interpretation_t
|
||||
interpret /* interpretation of number E4M3/E5M2 */
|
||||
) {
|
||||
__half hf = __hip_cvt_fp8_to_halfraw(in, interpret);
|
||||
return hf;
|
||||
}
|
||||
|
||||
int main() {
|
||||
constexpr size_t size = 32;
|
||||
hipDeviceProp_t prop;
|
||||
hip_check(hipGetDeviceProperties(&prop, 0));
|
||||
bool is_supported = (std::string(prop.gcnArchName).find("gfx94") != std::string::npos) || // gfx94x
|
||||
(std::string(prop.gcnArchName).find("gfx120") != std::string::npos); // gfx120x
|
||||
if(!is_supported) {
|
||||
std::cerr << "Need a gfx94x or gfx120x, but found: " << prop.gcnArchName << std::endl;
|
||||
std::cerr << "No device conversions are supported, only host conversions are supported." << std::endl;
|
||||
return -1;
|
||||
}
|
||||
|
||||
const __hip_fp8_interpretation_t interpret = (std::string(prop.gcnArchName).find("gfx94") != std::string::npos)
|
||||
? __HIP_E4M3_FNUZ // gfx94x
|
||||
: __HIP_E4M3; // gfx120x
|
||||
constexpr __hip_saturation_t sat = __HIP_SATFINITE;
|
||||
|
||||
std::vector<float> in;
|
||||
in.reserve(size);
|
||||
for (size_t i = 0; i < size; i++) {
|
||||
in.push_back(i + 1.1f);
|
||||
}
|
||||
|
||||
std::cout << "Converting float to fp8 and back..." << std::endl;
|
||||
// CPU convert
|
||||
std::vector<float> cpu_out;
|
||||
cpu_out.reserve(size);
|
||||
for (const auto &fval : in) {
|
||||
auto fp8 = convert_float_to_fp8(fval, interpret, sat);
|
||||
cpu_out.push_back(convert_fp8_to_float(fp8, interpret));
|
||||
}
|
||||
|
||||
// GPU convert
|
||||
float *d_in, *d_out;
|
||||
hip_check(hipMalloc(&d_in, sizeof(float) * size));
|
||||
hip_check(hipMalloc(&d_out, sizeof(float) * size));
|
||||
|
||||
hip_check(hipMemcpy(d_in, in.data(), sizeof(float) * in.size(),
|
||||
hipMemcpyHostToDevice));
|
||||
|
||||
float_to_fp8_to_float<<<1, size>>>(d_in, interpret, sat, d_out, size);
|
||||
|
||||
std::vector<float> gpu_out(size, 0.0f);
|
||||
hip_check(hipMemcpy(gpu_out.data(), d_out, sizeof(float) * gpu_out.size(),
|
||||
hipMemcpyDeviceToHost));
|
||||
|
||||
hip_check(hipFree(d_in));
|
||||
hip_check(hipFree(d_out));
|
||||
|
||||
// Validation
|
||||
for (size_t i = 0; i < size; i++) {
|
||||
if (cpu_out[i] != gpu_out[i]) {
|
||||
std::cerr << "cpu round trip result: " << cpu_out[i]
|
||||
<< " - gpu round trip result: " << gpu_out[i] << std::endl;
|
||||
std::abort();
|
||||
}
|
||||
}
|
||||
std::cout << "...CPU and GPU round trip convert matches." << std::endl;
|
||||
}
|
||||
|
||||
There are C++ style classes available as well.
|
||||
|
||||
.. code-block:: c
|
||||
|
||||
__hip_fp8_e4m3_fnuz fp8_val(1.1f); // gfx94x
|
||||
__hip_fp8_e4m3 fp8_val(1.1f); // gfx120x
|
||||
|
||||
Each type of FP8 number has its own class:
|
||||
|
||||
- __hip_fp8_e4m3
|
||||
- __hip_fp8_e5m2
|
||||
- __hip_fp8_e4m3_fnuz
|
||||
- __hip_fp8_e5m2_fnuz
|
||||
|
||||
There is support of vector of FP8 types.
|
||||
|
||||
- __hip_fp8x2_e4m3: holds 2 values of OCP FP8 e4m3 numbers
|
||||
- __hip_fp8x4_e4m3: holds 4 values of OCP FP8 e4m3 numbers
|
||||
- __hip_fp8x2_e5m2: holds 2 values of OCP FP8 e5m2 numbers
|
||||
- __hip_fp8x4_e5m2: holds 4 values of OCP FP8 e5m2 numbers
|
||||
- __hip_fp8x2_e4m3_fnuz: holds 2 values of FP8 fnuz e4m3 numbers
|
||||
- __hip_fp8x4_e4m3_fnuz: holds 4 values of FP8 fnuz e4m3 numbers
|
||||
- __hip_fp8x2_e5m2_fnuz: holds 2 values of FP8 fnuz e5m2 numbers
|
||||
- __hip_fp8x4_e5m2_fnuz: holds 4 values of FP8 fnuz e5m2 numbers
|
||||
|
||||
FNUZ extensions will be available on gfx94x only.
|
||||
@@ -0,0 +1,470 @@
|
||||
.. meta::
|
||||
:description: This page describes the FP8 and FP16 types present in HIP.
|
||||
:keywords: AMD, ROCm, HIP, fp8, fnuz, ocp
|
||||
|
||||
*******************************************************************************
|
||||
Low precision floating point types
|
||||
*******************************************************************************
|
||||
|
||||
Modern computing tasks often require balancing numerical precision against hardware resources
|
||||
and processing speed. Low precision floating point number formats in HIP include FP8 (Quarter Precision)
|
||||
and FP16 (Half Precision), which reduce memory and bandwidth requirements compared to traditional
|
||||
32-bit or 64-bit formats. The following sections detail their specifications, variants, and provide
|
||||
practical guidance for implementation in HIP.
|
||||
|
||||
FP8 (Quarter Precision)
|
||||
=======================
|
||||
|
||||
`FP8 (Floating Point 8-bit) numbers <https://arxiv.org/pdf/2209.05433>`_ were introduced
|
||||
as a compact numerical format specifically tailored for deep learning inference. By reducing
|
||||
precision while maintaining computational effectiveness, FP8 allows for significant memory
|
||||
savings and improved processing speed. This makes it particularly beneficial for deploying
|
||||
large-scale models with strict efficiency constraints.
|
||||
|
||||
Unlike traditional floating-point formats such as FP32 or even FP16, FP8 further optimizes
|
||||
performance by enabling a higher volume of matrix operations per second. Its reduced bit-width
|
||||
minimizes bandwidth requirements, making it an attractive choice for hardware accelerators
|
||||
in deep learning applications.
|
||||
|
||||
There are two primary FP8 formats:
|
||||
|
||||
- **E4M3 Format**
|
||||
|
||||
- Sign: 1 bit
|
||||
- Exponent: 4 bits
|
||||
- Mantissa: 3 bits
|
||||
|
||||
- **E5M2 Format**
|
||||
|
||||
- Sign: 1 bit
|
||||
- Exponent: 5 bits
|
||||
- Mantissa: 2 bits
|
||||
|
||||
The E4M3 format offers higher precision with a narrower range, while the E5M2 format provides
|
||||
a wider range at the cost of some precision.
|
||||
|
||||
Additionally, FP8 numbers have two representations:
|
||||
|
||||
- **FP8-OCP (Open Compute Project)**
|
||||
|
||||
- `This <https://www.opencompute.org/documents/ocp-8-bit-floating-point-specification-ofp8-revision-1-0-2023-12-01-pdf-1>`_
|
||||
is a standardized format developed by the Open Compute Project to ensure compatibility
|
||||
across various hardware and software implementations.
|
||||
|
||||
- **FP8-FNUZ (Finite and NaN Only)**
|
||||
|
||||
- A specialized format optimized for specific computations, supporting only finite and NaN values
|
||||
(no Inf support).
|
||||
- This provides one extra value of exponent and adds to the range of supported FP8 numbers.
|
||||
- **NaN Definition**: When the sign bit is set, and all other exponent and mantissa bits are zero.
|
||||
|
||||
The FNUZ representation provides an extra exponent value, expanding the range of representable
|
||||
numbers compared to standard FP8 formats.
|
||||
|
||||
|
||||
HIP Header
|
||||
----------
|
||||
|
||||
The `HIP FP8 header <https://github.com/ROCm/clr/blob/develop/hipamd/include/hip/amd_detail/amd_hip_fp8.h>`_
|
||||
defines the FP8 ocp/fnuz numbers.
|
||||
|
||||
Supported Devices
|
||||
-----------------
|
||||
|
||||
Different GPU models support different FP8 formats. Here's a breakdown:
|
||||
|
||||
.. list-table:: Supported devices for fp8 numbers
|
||||
:header-rows: 1
|
||||
|
||||
* - Device Type
|
||||
- FNUZ FP8
|
||||
- OCP FP8
|
||||
* - Host
|
||||
- Yes
|
||||
- Yes
|
||||
* - CDNA1
|
||||
- No
|
||||
- No
|
||||
* - CDNA2
|
||||
- No
|
||||
- No
|
||||
* - CDNA3
|
||||
- Yes
|
||||
- No
|
||||
* - RDNA2
|
||||
- No
|
||||
- No
|
||||
* - RDNA3
|
||||
- No
|
||||
- No
|
||||
|
||||
Using FP8 Numbers in HIP Programs
|
||||
---------------------------------
|
||||
|
||||
To use the FP8 numbers inside HIP programs.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <hip/hip_fp8.h>
|
||||
|
||||
FP8 numbers can be used on CPU side:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
__hip_fp8_storage_t convert_float_to_fp8(
|
||||
float in, /* Input val */
|
||||
__hip_fp8_interpretation_t interpret, /* interpretation of number E4M3/E5M2 */
|
||||
__hip_saturation_t sat /* Saturation behavior */
|
||||
) {
|
||||
return __hip_cvt_float_to_fp8(in, sat, interpret);
|
||||
}
|
||||
|
||||
The same can be done in kernels as well.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
__device__ __hip_fp8_storage_t d_convert_float_to_fp8(
|
||||
float in,
|
||||
__hip_fp8_interpretation_t interpret,
|
||||
__hip_saturation_t sat) {
|
||||
return __hip_cvt_float_to_fp8(in, sat, interpret);
|
||||
}
|
||||
|
||||
Note: On a gfx94x GPU, the type will default to the fnuz type.
|
||||
|
||||
The following code example does roundtrip FP8 conversions on both the CPU and GPU and compares the results.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <hip/hip_fp8.h>
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <iostream>
|
||||
#include <vector>
|
||||
|
||||
#define hip_check(hip_call) \
|
||||
{ \
|
||||
auto hip_res = hip_call; \
|
||||
if (hip_res != hipSuccess) { \
|
||||
std::cerr << "Failed in HIP call: " << #hip_call \
|
||||
<< " at " << __FILE__ << ":" << __LINE__ \
|
||||
<< " with error: " << hipGetErrorString(hip_res) << std::endl; \
|
||||
std::abort(); \
|
||||
} \
|
||||
}
|
||||
|
||||
__device__ __hip_fp8_storage_t d_convert_float_to_fp8(
|
||||
float in, __hip_fp8_interpretation_t interpret, __hip_saturation_t sat) {
|
||||
return __hip_cvt_float_to_fp8(in, sat, interpret);
|
||||
}
|
||||
|
||||
__device__ float d_convert_fp8_to_float(float in,
|
||||
__hip_fp8_interpretation_t interpret) {
|
||||
__half hf = __hip_cvt_fp8_to_halfraw(in, interpret);
|
||||
return hf;
|
||||
}
|
||||
|
||||
__global__ void float_to_fp8_to_float(float *in,
|
||||
__hip_fp8_interpretation_t interpret,
|
||||
__hip_saturation_t sat, float *out,
|
||||
size_t size) {
|
||||
int i = threadIdx.x;
|
||||
if (i < size) {
|
||||
auto fp8 = d_convert_float_to_fp8(in[i], interpret, sat);
|
||||
out[i] = d_convert_fp8_to_float(fp8, interpret);
|
||||
}
|
||||
}
|
||||
|
||||
__hip_fp8_storage_t
|
||||
convert_float_to_fp8(float in, /* Input val */
|
||||
__hip_fp8_interpretation_t
|
||||
interpret, /* interpretation of number E4M3/E5M2 */
|
||||
__hip_saturation_t sat /* Saturation behavior */
|
||||
) {
|
||||
return __hip_cvt_float_to_fp8(in, sat, interpret);
|
||||
}
|
||||
|
||||
float convert_fp8_to_float(
|
||||
__hip_fp8_storage_t in, /* Input val */
|
||||
__hip_fp8_interpretation_t
|
||||
interpret /* interpretation of number E4M3/E5M2 */
|
||||
) {
|
||||
__half hf = __hip_cvt_fp8_to_halfraw(in, interpret);
|
||||
return hf;
|
||||
}
|
||||
|
||||
int main() {
|
||||
constexpr size_t size = 32;
|
||||
hipDeviceProp_t prop;
|
||||
hip_check(hipGetDeviceProperties(&prop, 0));
|
||||
bool is_supported = (std::string(prop.gcnArchName).find("gfx94") != std::string::npos); // gfx94x
|
||||
if(!is_supported) {
|
||||
std::cerr << "Need a gfx94x, but found: " << prop.gcnArchName << std::endl;
|
||||
std::cerr << "No device conversions are supported, only host conversions are supported." << std::endl;
|
||||
return -1;
|
||||
}
|
||||
|
||||
const __hip_fp8_interpretation_t interpret = (std::string(prop.gcnArchName).find("gfx94") != std::string::npos)
|
||||
? __HIP_E4M3_FNUZ // gfx94x
|
||||
: __HIP_E4M3;
|
||||
constexpr __hip_saturation_t sat = __HIP_SATFINITE;
|
||||
|
||||
std::vector<float> in;
|
||||
in.reserve(size);
|
||||
for (size_t i = 0; i < size; i++) {
|
||||
in.push_back(i + 1.1f);
|
||||
}
|
||||
|
||||
std::cout << "Converting float to fp8 and back..." << std::endl;
|
||||
// CPU convert
|
||||
std::vector<float> cpu_out;
|
||||
cpu_out.reserve(size);
|
||||
for (const auto &fval : in) {
|
||||
auto fp8 = convert_float_to_fp8(fval, interpret, sat);
|
||||
cpu_out.push_back(convert_fp8_to_float(fp8, interpret));
|
||||
}
|
||||
|
||||
// GPU convert
|
||||
float *d_in, *d_out;
|
||||
hip_check(hipMalloc(&d_in, sizeof(float) * size));
|
||||
hip_check(hipMalloc(&d_out, sizeof(float) * size));
|
||||
|
||||
hip_check(hipMemcpy(d_in, in.data(), sizeof(float) * in.size(),
|
||||
hipMemcpyHostToDevice));
|
||||
|
||||
float_to_fp8_to_float<<<1, size>>>(d_in, interpret, sat, d_out, size);
|
||||
|
||||
std::vector<float> gpu_out(size, 0.0f);
|
||||
hip_check(hipMemcpy(gpu_out.data(), d_out, sizeof(float) * gpu_out.size(),
|
||||
hipMemcpyDeviceToHost));
|
||||
|
||||
hip_check(hipFree(d_in));
|
||||
hip_check(hipFree(d_out));
|
||||
|
||||
// Validation
|
||||
for (size_t i = 0; i < size; i++) {
|
||||
if (cpu_out[i] != gpu_out[i]) {
|
||||
std::cerr << "cpu round trip result: " << cpu_out[i]
|
||||
<< " - gpu round trip result: " << gpu_out[i] << std::endl;
|
||||
std::abort();
|
||||
}
|
||||
}
|
||||
std::cout << "...CPU and GPU round trip convert matches." << std::endl;
|
||||
}
|
||||
|
||||
There are C++ style classes available as well.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
__hip_fp8_e4m3_fnuz fp8_val(1.1f); // gfx94x
|
||||
__hip_fp8_e4m3 fp8_val(1.1f);
|
||||
|
||||
Each type of FP8 number has its own class:
|
||||
|
||||
- __hip_fp8_e4m3
|
||||
- __hip_fp8_e5m2
|
||||
- __hip_fp8_e4m3_fnuz
|
||||
- __hip_fp8_e5m2_fnuz
|
||||
|
||||
There is support of vector of FP8 types.
|
||||
|
||||
- __hip_fp8x2_e4m3: holds 2 values of OCP FP8 e4m3 numbers
|
||||
- __hip_fp8x4_e4m3: holds 4 values of OCP FP8 e4m3 numbers
|
||||
- __hip_fp8x2_e5m2: holds 2 values of OCP FP8 e5m2 numbers
|
||||
- __hip_fp8x4_e5m2: holds 4 values of OCP FP8 e5m2 numbers
|
||||
- __hip_fp8x2_e4m3_fnuz: holds 2 values of FP8 fnuz e4m3 numbers
|
||||
- __hip_fp8x4_e4m3_fnuz: holds 4 values of FP8 fnuz e4m3 numbers
|
||||
- __hip_fp8x2_e5m2_fnuz: holds 2 values of FP8 fnuz e5m2 numbers
|
||||
- __hip_fp8x4_e5m2_fnuz: holds 4 values of FP8 fnuz e5m2 numbers
|
||||
|
||||
FNUZ extensions will be available on gfx94x only.
|
||||
|
||||
FP16 (Half Precision)
|
||||
=====================
|
||||
|
||||
FP16 (Floating Point 16-bit) numbers offer a balance between precision and
|
||||
efficiency, making them a widely adopted standard for accelerating deep learning
|
||||
inference. With higher precision than FP8 but lower memory requirements than FP32,
|
||||
FP16 enables faster computations while preserving model accuracy.
|
||||
|
||||
Deep learning workloads often involve massive datasets and complex calculations,
|
||||
making FP32 computationally expensive. FP16 helps mitigate these costs by reducing
|
||||
storage and bandwidth demands, allowing for increased throughput without significant
|
||||
loss of numerical stability. This format is particularly useful for training and
|
||||
inference in GPUs and TPUs optimized for half-precision arithmetic.
|
||||
|
||||
There are two primary FP16 formats:
|
||||
|
||||
- **float16 Format**
|
||||
|
||||
- Sign: 1 bit
|
||||
- Exponent: 5 bits
|
||||
- Mantissa: 10 bits
|
||||
|
||||
- **bfloat16 Format**
|
||||
|
||||
- Sign: 1 bit
|
||||
- Exponent: 8 bits
|
||||
- Mantissa: 7 bits
|
||||
|
||||
The float16 format offers higher precision with a narrower range, while the bfloat16
|
||||
format provides a wider range at the cost of some precision.
|
||||
|
||||
Additionally, FP16 numbers have standardized representations developed by industry
|
||||
initiatives to ensure compatibility across various hardware and software implementations.
|
||||
Unlike FP8, which has specific representations like OCP and FNUZ, FP16 is more uniformly
|
||||
supported with its two main formats, float16 and bfloat16.
|
||||
|
||||
HIP Header
|
||||
----------
|
||||
|
||||
The `HIP FP16 header <https://github.com/ROCm/clr/blob/develop/hipamd/include/hip/amd_detail/amd_hip_fp16.h>`_
|
||||
defines the float16 format.
|
||||
|
||||
The `HIP BF16 header <https://github.com/ROCm/clr/blob/develop/hipamd/include/hip/amd_detail/amd_hip_bf16.h>`_
|
||||
defines the bfloat16 format.
|
||||
|
||||
Supported Devices
|
||||
-----------------
|
||||
|
||||
Different GPU models support different FP16 formats. Here's a breakdown:
|
||||
|
||||
.. list-table:: Supported devices for fp16 numbers
|
||||
:header-rows: 1
|
||||
|
||||
* - Device Type
|
||||
- float16
|
||||
- bfloat16
|
||||
* - Host
|
||||
- Yes
|
||||
- Yes
|
||||
* - CDNA1
|
||||
- Yes
|
||||
- Yes
|
||||
* - CDNA2
|
||||
- Yes
|
||||
- Yes
|
||||
* - CDNA3
|
||||
- Yes
|
||||
- Yes
|
||||
* - RDNA2
|
||||
- Yes
|
||||
- Yes
|
||||
* - RDNA3
|
||||
- Yes
|
||||
- Yes
|
||||
|
||||
Using FP16 Numbers in HIP Programs
|
||||
----------------------------------
|
||||
|
||||
To use the FP16 numbers inside HIP programs.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <hip/hip_fp16.h> // for float16
|
||||
#include <hip/hip_bf16.h> // for bfloat16
|
||||
|
||||
The following code example adds two float16 values on the GPU and compares the results
|
||||
against summed float values on the CPU.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <hip/hip_fp16.h>
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <iostream>
|
||||
#include <vector>
|
||||
|
||||
#define hip_check(hip_call) \
|
||||
{ \
|
||||
auto hip_res = hip_call; \
|
||||
if (hip_res != hipSuccess) { \
|
||||
std::cerr << "Failed in HIP call: " << #hip_call \
|
||||
<< " at " << __FILE__ << ":" << __LINE__ \
|
||||
<< " with error: " << hipGetErrorString(hip_res) << std::endl; \
|
||||
std::abort(); \
|
||||
} \
|
||||
}
|
||||
|
||||
__global__ void add_half_precision(__half* in1, __half* in2, float* out, size_t size) {
|
||||
int idx = threadIdx.x;
|
||||
if (idx < size) {
|
||||
// Load as half, perform addition in float, store as float
|
||||
float sum = __half2float(in1[idx] + in2[idx]);
|
||||
out[idx] = sum;
|
||||
}
|
||||
}
|
||||
|
||||
int main() {
|
||||
constexpr size_t size = 32;
|
||||
constexpr float tolerance = 1e-1f; // Allowable numerical difference
|
||||
|
||||
// Initialize input vectors as floats
|
||||
std::vector<float> in1(size), in2(size);
|
||||
for (size_t i = 0; i < size; i++) {
|
||||
in1[i] = i + 1.1f;
|
||||
in2[i] = i + 2.2f;
|
||||
}
|
||||
|
||||
// Compute expected results in full precision on CPU
|
||||
std::vector<float> cpu_out(size);
|
||||
for (size_t i = 0; i < size; i++) {
|
||||
cpu_out[i] = in1[i] + in2[i]; // Direct float addition
|
||||
}
|
||||
|
||||
// Allocate device memory (store input as half, output as float)
|
||||
__half *d_in1, *d_in2;
|
||||
float *d_out;
|
||||
hip_check(hipMalloc(&d_in1, sizeof(__half) * size));
|
||||
hip_check(hipMalloc(&d_in2, sizeof(__half) * size));
|
||||
hip_check(hipMalloc(&d_out, sizeof(float) * size));
|
||||
|
||||
// Convert input to half and copy to device
|
||||
std::vector<__half> in1_half(size), in2_half(size);
|
||||
for (size_t i = 0; i < size; i++) {
|
||||
in1_half[i] = __float2half(in1[i]);
|
||||
in2_half[i] = __float2half(in2[i]);
|
||||
}
|
||||
|
||||
hip_check(hipMemcpy(d_in1, in1_half.data(), sizeof(__half) * size, hipMemcpyHostToDevice));
|
||||
hip_check(hipMemcpy(d_in2, in2_half.data(), sizeof(__half) * size, hipMemcpyHostToDevice));
|
||||
|
||||
// Launch kernel
|
||||
add_half_precision<<<1, size>>>(d_in1, d_in2, d_out, size);
|
||||
|
||||
// Copy result back to host
|
||||
std::vector<float> gpu_out(size, 0.0f);
|
||||
hip_check(hipMemcpy(gpu_out.data(), d_out, sizeof(float) * size, hipMemcpyDeviceToHost));
|
||||
|
||||
// Free device memory
|
||||
hip_check(hipFree(d_in1));
|
||||
hip_check(hipFree(d_in2));
|
||||
hip_check(hipFree(d_out));
|
||||
|
||||
// Validation with tolerance
|
||||
for (size_t i = 0; i < size; i++) {
|
||||
if (std::fabs(cpu_out[i] - gpu_out[i]) > tolerance) {
|
||||
std::cerr << "Mismatch at index " << i << ": CPU result = " << cpu_out[i]
|
||||
<< ", GPU result = " << gpu_out[i] << std::endl;
|
||||
std::abort();
|
||||
}
|
||||
}
|
||||
|
||||
std::cout << "Success: CPU and GPU half-precision addition match within tolerance!" << std::endl;
|
||||
}
|
||||
|
||||
|
||||
There are C++ style classes available as well.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
__half fp16_val(1.1f); // float16
|
||||
__hip_bfloat16 fp16_val(1.1f); // bfloat16
|
||||
|
||||
Each type of FP16 number has its own class:
|
||||
|
||||
- __half
|
||||
- __hip_bfloat16
|
||||
|
||||
There is support of vector of FP16 types.
|
||||
|
||||
- __half2: holds 2 values of float16 numbers
|
||||
- __hip_bfloat162: holds 2 values of bfloat16 numbers
|
||||
+1449
-1019
File diff suppressed because it is too large
Load Diff
@@ -1,35 +0,0 @@
|
||||
.. meta::
|
||||
:description: This chapter lists user-mode API interfaces and libraries
|
||||
necessary for host applications to launch compute kernels to
|
||||
available HSA ROCm kernel agents.
|
||||
:keywords: AMD, ROCm, HIP, HSA, ROCR runtime, virtual memory management
|
||||
|
||||
*******************************************************************************
|
||||
HSA runtime API for ROCm
|
||||
*******************************************************************************
|
||||
|
||||
The following functions are located in the https://github.com/ROCm/ROCR-Runtime repository.
|
||||
|
||||
.. doxygenfunction:: hsa_amd_vmem_address_reserve
|
||||
|
||||
.. doxygenfunction:: hsa_amd_vmem_address_free
|
||||
|
||||
.. doxygenfunction:: hsa_amd_vmem_handle_create
|
||||
|
||||
.. doxygenfunction:: hsa_amd_vmem_handle_release
|
||||
|
||||
.. doxygenfunction:: hsa_amd_vmem_map
|
||||
|
||||
.. doxygenfunction:: hsa_amd_vmem_unmap
|
||||
|
||||
.. doxygenfunction:: hsa_amd_vmem_set_access
|
||||
|
||||
.. doxygenfunction:: hsa_amd_vmem_get_access
|
||||
|
||||
.. doxygenfunction:: hsa_amd_vmem_export_shareable_handle
|
||||
|
||||
.. doxygenfunction:: hsa_amd_vmem_import_shareable_handle
|
||||
|
||||
.. doxygenfunction:: hsa_amd_vmem_retain_alloc_handle
|
||||
|
||||
.. doxygenfunction:: hsa_amd_vmem_get_alloc_properties_from_handle
|
||||
Reference in New Issue
Block a user