HIP complex math API

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:

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.

Type

Description

hip​Float​Complex
Complex number using single-precision (float) values
(note: hip​Complex is an alias of hip​Float​Complex)

hip​Double​Complex

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.

Function

Description
hip​Float​Complex
make_​hip​Float​Complex(
float a,
float b
)
Creates a complex number
(note: make_​hip​Complex is an alias of make_​hip​Float​Complex)
z=a+bi
float
hip​Crealf(
hip​Float​Complex z
)
Returns real part of z
ℜ(z)=x
float
hip​Cimagf(
hip​Float​Complex z
)
Returns imaginary part of z
ℑ(z)=y

Function

Description
hip​Double​Complex
make_​hip​Double​Complex(
double a,
double b
)
Creates a complex number
z=a+bi
double
hip​Creal(
hip​Double​Complex z
)
Returns real part of z
ℜ(z)=x
double
hip​Cimag(
hip​Double​Complex z
)
Returns imaginary part of z
ℑ(z)=y

Basic Arithmetic#

Operations for performing standard arithmetic with complex numbers, including addition, subtraction, multiplication, division, and fused multiply-add.

Function

Description
hip​Float​Complex
hip​Caddf(
hip​Float​Complex p,
hip​Float​Complex q
)
Addition of two single-precision complex values
(a+bi)+(c+di)=(a+c)+(b+d)i
hip​Float​Complex
hip​Csubf(
hip​Float​Complex p,
hip​Float​Complex q
)
Subtraction of two single-precision complex values
(a+bi)−(c+di)=(a−c)+(b−d)i
hip​Float​Complex
hip​Cmulf(
hip​Float​Complex p,
hip​Float​Complex q
)
Multiplication of two single-precision complex values
(a+bi)(c+di)=(ac−bd)+(bc+ad)i
hip​Float​Complex
hip​Cdivf(
hip​Float​Complex p,
hip​Float​Complex q
)
Division of two single-precision complex values
a+bic+di=(ac+bd)+(bc−ad)ic2+d2
hip​Float​Complex
hip​Cfmaf(
hip​Complex p,
hip​Complex q,
hip​Complex r
)
Fused multiply-add of three single-precision complex values
(a+bi)(c+di)+(e+fi)

Function

Description
hip​Double​Complex
hip​Cadd(
hip​Double​Complex p,
hip​Double​Complex q
)
Addition of two double-precision complex values
(a+bi)+(c+di)=(a+c)+(b+d)i
hip​Double​Complex
hip​Csub(
hip​Double​Complex p,
hip​Double​Complex q
)
Subtraction of two double-precision complex values
(a+bi)−(c+di)=(a−c)+(b−d)i
hip​Double​Complex
hip​Cmul(
hip​Double​Complex p,
hip​Double​Complex q
)
Multiplication of two double-precision complex values
(a+bi)(c+di)=(ac−bd)+(bc+ad)i
hip​Double​Complex
hip​Cdiv(
hip​Double​Complex p,
hip​Double​Complex q
)
Division of two double-precision complex values
a+bic+di=(ac+bd)+(bc−ad)ic2+d2
hip​Double​Complex
hip​Cfma(
hip​Double​Complex p,
hip​Double​Complex q,
hip​Double​Complex r
)
Fused multiply-add of three double-precision complex values
(a+bi)(c+di)+(e+fi)

Complex Operations#

Functions for complex-specific calculations, including conjugate determination and magnitude (absolute value) computation.

Function

Description
hip​Float​Complex
hip​Conjf(
hip​Float​Complex z
)
Complex conjugate
a+bi―=a−bi
float
hip​Cabsf(
hip​Float​Complex z
)
Absolute value (magnitude)
|a+bi|=a2+b2
float
hip​Csqabsf(
hip​Float​Complex z
)
Squared absolute value
|a+bi|2=a2+b2

Function

Description
hip​Double​Complex
hip​Conj(
hip​Double​Complex z
)
Complex conjugate
a+bi―=a−bi
double
hip​Cabs(
hip​Double​Complex z
)
Absolute value (magnitude)
|a+bi|=a2+b2
double
hip​Csqabs(
hip​Double​Complex z
)
Squared absolute value
|a+bi|2=a2+b2

Type Conversion#

Utility functions for conversion between single-precision and double-precision complex number formats.

Function

Description
hip​Float​Complex
hip​Complex​Double​To​Float(
hip​Double​Complex z
)

Converts double-precision to single-precision complex
hip​Double​Complex
hip​Complex​Float​To​Double(
hip​Float​Complex 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.

#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;
}