From 80557d425b014d2af745d1f0893e8b5caa4ec861 Mon Sep 17 00:00:00 2001 From: Adel Johar Date: Mon, 24 Feb 2025 12:51:47 +0100 Subject: [PATCH] Docs: Add page for Complex Math API --- .wordlist.txt | 1 + docs/index.md | 1 + docs/reference/complex_math_api.rst | 361 ++++++++++++++++++++++++++++ docs/sphinx/_toc.yml.in | 1 + 4 files changed, 364 insertions(+) create mode 100644 docs/reference/complex_math_api.rst diff --git a/.wordlist.txt b/.wordlist.txt index de7c91b31a..a97cbc235d 100644 --- a/.wordlist.txt +++ b/.wordlist.txt @@ -33,6 +33,7 @@ deallocate decompositions denormal Dereferencing +DFT dll DirectX EIGEN diff --git a/docs/index.md b/docs/index.md index 247c58e2fd..34d8360df9 100644 --- a/docs/index.md +++ b/docs/index.md @@ -43,6 +43,7 @@ The HIP documentation is organized into the following categories: * [HIP runtime API](./reference/hip_runtime_api_reference) * [HIP math API](./reference/math_api) +* [HIP complex math API](./reference/complex_math_api) * [HIP environment variables](./reference/env_variables) * [CUDA to HIP API Function Comparison](./reference/api_syntax) * [List of deprecated APIs](./reference/deprecated_api_list) diff --git a/docs/reference/complex_math_api.rst b/docs/reference/complex_math_api.rst new file mode 100644 index 0000000000..9b42757ff1 --- /dev/null +++ b/docs/reference/complex_math_api.rst @@ -0,0 +1,361 @@ +.. 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. + +.. tab-set:: + + .. tab-item:: Single Precision + + .. list-table:: + :header-rows: 1 + :widths: 40 60 + + * - Type + - Description + + * - ``hipFloatComplex``/``hipComplex`` + - Complex number using single-precision (float) values + + .. tab-item:: Double Precision + + .. list-table:: + :header-rows: 1 + :widths: 40 60 + + * - Type + - Description + + * - ``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 +----------------- + +.. 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 +---------------- + +.. 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 +------------------ + +.. 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 +--------------- + +.. 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 + #include + #include + #include + #include + + #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 cpuDFT(const std::vector& input) { + const int N = input.size(); + std::vector 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 signal(N); + for (int i = 0; i < N; i++) { + float t = static_cast(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 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<<>>(d_signal, d_output, N); + HIP_CHECK(hipGetLastError()); + + // Get GPU results + std::vector 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; + } diff --git a/docs/sphinx/_toc.yml.in b/docs/sphinx/_toc.yml.in index 34050b2448..1f63d5fdaf 100644 --- a/docs/sphinx/_toc.yml.in +++ b/docs/sphinx/_toc.yml.in @@ -109,6 +109,7 @@ subtrees: - file: doxygen/html/annotated - file: doxygen/html/files - file: reference/math_api + - file: reference/complex_math_api - file: reference/env_variables - file: reference/api_syntax - file: reference/deprecated_api_list