Skip to content

Commit 3344621

Browse files
adeljo-amdadeljo-amd
committed
Docs: Add page for Complex Math API
1 parent c9d719f commit 3344621

File tree

4 files changed

+349
-0
lines changed

4 files changed

+349
-0
lines changed

‎.wordlist.txt

+1
Original file line numberDiff line numberDiff line change
@@ -33,6 +33,7 @@ deallocate
3333
decompositions
3434
denormal
3535
Dereferencing
36+
DFT
3637
dll
3738
DirectX
3839
EIGEN

‎docs/index.md

+1
Original file line numberDiff line numberDiff line change
@@ -43,6 +43,7 @@ The HIP documentation is organized into the following categories:
4343

4444
* [HIP runtime API](./reference/hip_runtime_api_reference)
4545
* [HIP math API](./reference/math_api)
46+
* [HIP complex math API](./reference/complex_math_api)
4647
* [HIP environment variables](./reference/env_variables)
4748
* [CUDA to HIP API Function Comparison](./reference/api_syntax)
4849
* [List of deprecated APIs](./reference/deprecated_api_list)

‎docs/reference/complex_math_api.rst

+346
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,346 @@
1+
.. meta::
2+
:description: This chapter describes the complex math functions that are accessible in HIP.
3+
:keywords: AMD, ROCm, HIP, CUDA, complex math functions, HIP complex math functions
4+
5+
.. _complex_math_api_reference:
6+
7+
********************************************************************************
8+
HIP complex math API
9+
********************************************************************************
10+
11+
HIP provides built-in support for complex number operations through specialized types and functions,
12+
available for both single-precision (float) and double-precision (double) calculations. All complex types
13+
and functions are available on both host and device.
14+
15+
For any complex number ``z``, the form is:
16+
17+
.. math::
18+
19+
z = x + yi
20+
21+
where ``x`` is the real part and ``y`` is the imaginary part.
22+
23+
Complex Number Types
24+
====================
25+
26+
.. tab-set::
27+
28+
.. tab-item:: Single Precision
29+
30+
.. list-table::
31+
:header-rows: 1
32+
:widths: 40 60
33+
34+
* - Type
35+
- Description
36+
37+
* - ``hipFloatComplex``/``hipComplex``
38+
- Complex number using single-precision (float) values
39+
40+
.. tab-item:: Double Precision
41+
42+
.. list-table::
43+
:header-rows: 1
44+
:widths: 40 60
45+
46+
* - Type
47+
- Description
48+
49+
* - ``hipDoubleComplex``
50+
- Complex number using double-precision (double) values
51+
52+
Complex Number Functions
53+
========================
54+
55+
Type Construction
56+
-----------------
57+
58+
.. tab-set::
59+
60+
.. tab-item:: Single Precision
61+
62+
.. list-table::
63+
:header-rows: 1
64+
:widths: 40 60
65+
66+
* - Function
67+
- Description
68+
69+
* - ``hipFloatComplex make_hipFloatComplex(float a, float b)``
70+
- | Creates a complex number (note: ``make_hipComplex`` is an alias of ``make_hipFloatComplex``)
71+
| :math:`z = a + bi`
72+
73+
* - ``float hipCrealf(hipFloatComplex z)``
74+
- | Returns real part of z
75+
| :math:`\Re(z) = x`
76+
77+
* - ``float hipCimagf(hipFloatComplex z)``
78+
- | Returns imaginary part of z
79+
| :math:`\Im(z) = y`
80+
81+
.. tab-item:: Double Precision
82+
83+
.. list-table::
84+
:header-rows: 1
85+
:widths: 40 60
86+
87+
* - Function
88+
- Description
89+
90+
* - ``hipDoubleComplex make_hipDoubleComplex(double a, double b)``
91+
- | Creates a complex number
92+
| :math:`z = a + bi`
93+
94+
* - ``double hipCreal(hipDoubleComplex z)``
95+
- | Returns real part of z
96+
| :math:`\Re(z) = x`
97+
98+
* - ``double hipCimag(hipDoubleComplex z)``
99+
- | Returns imaginary part of z
100+
| :math:`\Im(z) = y`
101+
102+
Basic Arithmetic
103+
----------------
104+
105+
.. tab-set::
106+
107+
.. tab-item:: Single Precision
108+
109+
.. list-table::
110+
:header-rows: 1
111+
:widths: 40 60
112+
113+
* - Function
114+
- Description
115+
116+
* - ``hipFloatComplex hipCaddf(hipFloatComplex p, hipFloatComplex q)``
117+
- | Addition of two single-precision complex values
118+
| :math:`(a + bi) + (c + di) = (a + c) + (b + d)i`
119+
120+
* - ``hipFloatComplex hipCsubf(hipFloatComplex p, hipFloatComplex q)``
121+
- | Subtraction of two single-precision complex values
122+
| :math:`(a + bi) - (c + di) = (a - c) + (b - d)i`
123+
124+
* - ``hipFloatComplex hipCmulf(hipFloatComplex p, hipFloatComplex q)``
125+
- | Multiplication of two single-precision complex values
126+
|:math:`(a + bi)(c + di) = (ac - bd) + (bc + ad)i`
127+
128+
* - ``hipFloatComplex hipCdivf(hipFloatComplex p, hipFloatComplex q)``
129+
- | Division of two single-precision complex values
130+
| :math:`\frac{a + bi}{c + di} = \frac{(ac + bd) + (bc - ad)i}{c^2 + d^2}`
131+
132+
* - ``hipFloatComplex hipCfmaf(hipComplex p, hipComplex q, hipComplex r)``
133+
- | Fused multiply-add of three single-precision complex values
134+
| :math:`(a + bi)(c + di) + (e + fi)`
135+
136+
.. tab-item:: Double Precision
137+
138+
.. list-table::
139+
:header-rows: 1
140+
:widths: 40 60
141+
142+
* - Function
143+
- Description
144+
145+
* - ``hipDoubleComplex hipCadd(hipDoubleComplex p, hipDoubleComplex q)``
146+
- | Addition of two double-precision complex values
147+
| :math:`(a + bi) + (c + di) = (a + c) + (b + d)i`
148+
149+
* - ``hipDoubleComplex hipCsub(hipDoubleComplex p, hipDoubleComplex q)``
150+
- | Subtraction of two double-precision complex values
151+
| :math:`(a + bi) - (c + di) = (a - c) + (b - d)i`
152+
153+
* - ``hipDoubleComplex hipCmul(hipDoubleComplex p, hipDoubleComplex q)``
154+
- | Multiplication of two double-precision complex values
155+
| :math:`(a + bi)(c + di) = (ac - bd) + (bc + ad)i`
156+
157+
* - ``hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleComplex q)``
158+
- | Division of two double-precision complex values
159+
| :math:`\frac{a + bi}{c + di} = \frac{(ac + bd) + (bc - ad)i}{c^2 + d^2}`
160+
161+
* - ``hipDoubleComplex hipCfma(hipDoubleComplex p, hipDoubleComplex q, hipDoubleComplex r)``
162+
- | Fused multiply-add of three double-precision complex values
163+
| :math:`(a + bi)(c + di) + (e + fi)`
164+
165+
Complex Operations
166+
------------------
167+
168+
.. tab-set::
169+
170+
.. tab-item:: Single Precision
171+
172+
.. list-table::
173+
:header-rows: 1
174+
:widths: 40 60
175+
176+
* - Function
177+
- Description
178+
179+
* - ``hipFloatComplex hipConjf(hipFloatComplex z)``
180+
- | Complex conjugate
181+
| :math:`\overline{a + bi} = a - bi`
182+
183+
* - ``float hipCabsf(hipFloatComplex z)``
184+
- | Absolute value (magnitude)
185+
| :math:`|a + bi| = \sqrt{a^2 + b^2}`
186+
187+
* - ``float hipCsqabsf(hipFloatComplex z)``
188+
- | Squared absolute value
189+
| :math:`|a + bi|^2 = a^2 + b^2`
190+
191+
.. tab-item:: Double Precision
192+
193+
.. list-table::
194+
:header-rows: 1
195+
:widths: 40 60
196+
197+
* - Function
198+
- Description
199+
200+
* - ``hipDoubleComplex hipConj(hipDoubleComplex z)``
201+
- | Complex conjugate
202+
| :math:`\overline{a + bi} = a - bi`
203+
204+
* - ``double hipCabs(hipDoubleComplex z)``
205+
- | Absolute value (magnitude)
206+
| :math:`|a + bi| = \sqrt{a^2 + b^2}`
207+
208+
* - ``double hipCsqabs(hipDoubleComplex z)``
209+
- | Squared absolute value
210+
| :math:`|a + bi|^2 = a^2 + b^2`
211+
212+
Type Conversion
213+
---------------
214+
215+
.. list-table::
216+
:header-rows: 1
217+
:widths: 40 60
218+
219+
* - Function
220+
- Description
221+
222+
* - ``hipFloatComplex hipComplexDoubleToFloat(hipDoubleComplex z)``
223+
- Converts double-precision to single-precision complex
224+
225+
* - ``hipDoubleComplex hipComplexFloatToDouble(hipFloatComplex z)``
226+
- Converts single-precision to double-precision complex
227+
228+
Example Usage
229+
=============
230+
231+
The following example demonstrates using complex numbers to compute the DFT of a simple signal on the GPU:
232+
233+
.. code-block:: cpp
234+
235+
#include <hip/hip_runtime.h>
236+
#include <hip/hip_complex.h>
237+
#include <iostream>
238+
#include <vector>
239+
#include <cmath>
240+
241+
#define HIP_CHECK(expression) \
242+
{ \
243+
const hipError_t err = expression; \
244+
if (err != hipSuccess) { \
245+
std::cerr << "HIP error: " \
246+
<< hipGetErrorString(err) \
247+
<< " at " << __LINE__ << "\n"; \
248+
exit(EXIT_FAILURE); \
249+
} \
250+
}
251+
252+
// Kernel to compute DFT
253+
__global__ void computeDFT(const float* input,
254+
hipFloatComplex* output,
255+
const int N)
256+
{
257+
int k = blockIdx.x * blockDim.x + threadIdx.x;
258+
if (k >= N) return;
259+
260+
hipFloatComplex sum = make_hipFloatComplex(0.0f, 0.0f);
261+
262+
for (int n = 0; n < N; n++) {
263+
float angle = -2.0f * M_PI * k * n / N;
264+
hipFloatComplex w = make_hipFloatComplex(cosf(angle), sinf(angle));
265+
hipFloatComplex x = make_hipFloatComplex(input[n], 0.0f);
266+
sum = hipCaddf(sum, hipCmulf(x, w));
267+
}
268+
269+
output[k] = sum;
270+
}
271+
272+
// CPU implementation of DFT for verification
273+
std::vector<hipFloatComplex> cpuDFT(const std::vector<float>& input) {
274+
const int N = input.size();
275+
std::vector<hipFloatComplex> result(N);
276+
277+
for (int k = 0; k < N; k++) {
278+
hipFloatComplex sum = make_hipFloatComplex(0.0f, 0.0f);
279+
for (int n = 0; n < N; n++) {
280+
float angle = -2.0f * M_PI * k * n / N;
281+
hipFloatComplex w = make_hipFloatComplex(cosf(angle), sinf(angle));
282+
hipFloatComplex x = make_hipFloatComplex(input[n], 0.0f);
283+
sum = hipCaddf(sum, hipCmulf(x, w));
284+
}
285+
result[k] = sum;
286+
}
287+
return result;
288+
}
289+
290+
int main() {
291+
const int N = 256; // Signal length
292+
const int blockSize = 256;
293+
294+
// Generate input signal: sum of two sine waves
295+
std::vector<float> signal(N);
296+
for (int i = 0; i < N; i++) {
297+
float t = static_cast<float>(i) / N;
298+
signal[i] = sinf(2.0f * M_PI * 10.0f * t) + // 10 Hz component
299+
0.5f * sinf(2.0f * M_PI * 20.0f * t); // 20 Hz component
300+
}
301+
302+
// Compute reference solution on CPU
303+
std::vector<hipFloatComplex> cpu_output = cpuDFT(signal);
304+
305+
// Allocate device memory
306+
float* d_signal;
307+
hipFloatComplex* d_output;
308+
HIP_CHECK(hipMalloc(&d_signal, N * sizeof(float)));
309+
HIP_CHECK(hipMalloc(&d_output, N * sizeof(hipFloatComplex)));
310+
311+
// Copy input to device
312+
HIP_CHECK(hipMemcpy(d_signal, signal.data(), N * sizeof(float),
313+
hipMemcpyHostToDevice));
314+
315+
// Launch kernel
316+
dim3 grid((N + blockSize - 1) / blockSize);
317+
dim3 block(blockSize);
318+
computeDFT<<<grid, block>>>(d_signal, d_output, N);
319+
HIP_CHECK(hipGetLastError());
320+
321+
// Get GPU results
322+
std::vector<hipFloatComplex> gpu_output(N);
323+
HIP_CHECK(hipMemcpy(gpu_output.data(), d_output, N * sizeof(hipFloatComplex),
324+
hipMemcpyDeviceToHost));
325+
326+
// Verify results
327+
bool passed = true;
328+
const float tolerance = 1e-5f; // Adjust based on precision requirements
329+
330+
for (int i = 0; i < N; i++) {
331+
float diff_real = std::abs(hipCrealf(gpu_output[i]) - hipCrealf(cpu_output[i]));
332+
float diff_imag = std::abs(hipCimagf(gpu_output[i]) - hipCimagf(cpu_output[i]));
333+
334+
if (diff_real > tolerance || diff_imag > tolerance) {
335+
passed = false;
336+
break;
337+
}
338+
}
339+
340+
std::cout << "DFT Verification: " << (passed ? "PASSED" : "FAILED") << "\n";
341+
342+
// Cleanup
343+
HIP_CHECK(hipFree(d_signal));
344+
HIP_CHECK(hipFree(d_output));
345+
return passed ? 0 : 1;
346+
}

‎docs/sphinx/_toc.yml.in

+1
Original file line numberDiff line numberDiff line change
@@ -109,6 +109,7 @@ subtrees:
109109
- file: doxygen/html/annotated
110110
- file: doxygen/html/files
111111
- file: reference/math_api
112+
- file: reference/complex_math_api
112113
- file: reference/env_variables
113114
- file: reference/api_syntax
114115
- file: reference/deprecated_api_list

0 commit comments

Comments
 (0)