Skip to content

Commit a78d285

Browse files
j-stephanj-stephan
committed
Cosmetic changes
Signed-off-by: Jan Stephan <jan.stephan@amd.com>
1 parent bdb8daa commit a78d285

File tree

5 files changed

+24
-19
lines changed

5 files changed

+24
-19
lines changed

‎docs/tools/example_codes/graph_api_tutorial_main_graph_capture.hip‎

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -144,7 +144,7 @@ auto main() -> int
144144
<< volGeom.dim.z << std::endl;
145145

146146
// Initialize per-stream data
147-
for(auto streamIdx = 0; streamIdx < streams.size(); ++streamIdx)
147+
for(auto streamIdx = 0u; streamIdx < streams.size(); ++streamIdx)
148148
{
149149
std::cout << "Initializing stream " << streamIdx << "... " << std::flush;
150150
auto stream = streams.at(streamIdx);
@@ -265,7 +265,7 @@ auto main() -> int
265265

266266
auto start = std::chrono::steady_clock::now();
267267

268-
auto projIdx = 0;
268+
auto projIdx = 0u;
269269
while(projIdx < projGeom.numProj)
270270
{
271271
// [sphinx-begin-capture-start]
@@ -283,7 +283,7 @@ auto main() -> int
283283
auto forkEvent = hipEvent_t{};
284284
hip_check(hipEventCreate(&forkEvent));
285285
hip_check(hipEventRecord(forkEvent, streams.at(0)));
286-
hip_check(hipStreamWaitEvent(streams.at(streamIdx), forkEvent));
286+
hip_check(hipStreamWaitEvent(streams.at(streamIdx), forkEvent, 0));
287287
hip_check(hipEventDestroy(forkEvent)); // Can destroy after wait is recorded
288288
}
289289
// [sphinx-fork-end]
@@ -438,7 +438,7 @@ auto main() -> int
438438
auto joinEvent = hipEvent_t{};
439439
hip_check(hipEventCreate(&joinEvent));
440440
hip_check(hipEventRecord(joinEvent, streams.at(streamIdx)));
441-
hip_check(hipStreamWaitEvent(streams.at(0), joinEvent));
441+
hip_check(hipStreamWaitEvent(streams.at(0), joinEvent, 0));
442442
hip_check(hipEventDestroy(joinEvent)); // Can destroy after wait is recorded
443443
}
444444
// [sphinx-join-end]
@@ -554,7 +554,7 @@ auto main() -> int
554554
std::cout << "Done!" << std::endl;
555555

556556
auto const elapsed = std::chrono::duration<double>{stop - start};
557-
std::cout << "Reconstruction time: " << elapsed << std::endl;
557+
std::cout << "Reconstruction time: " << elapsed.count() << 's' << std::endl;
558558

559559
// Cleanup
560560
if(graphFinalCreated)

‎docs/tools/example_codes/graph_api_tutorial_main_graph_creation.hip‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -265,7 +265,7 @@ auto main() -> int
265265

266266
auto start = std::chrono::steady_clock::now();
267267

268-
auto projIdx = 0;
268+
auto projIdx = 0u;
269269
while(projIdx < projGeom.numProj)
270270
{
271271
auto batchSize = std::min(numBranches, static_cast<int>(projGeom.numProj - projIdx));
@@ -752,7 +752,7 @@ auto main() -> int
752752
std::cout << "Done!" << std::endl;
753753

754754
auto const elapsed = std::chrono::duration<double>{stop - start};
755-
std::cout << "Reconstruction time: " << elapsed << std::endl;
755+
std::cout << "Reconstruction time: " << elapsed.count() << 's' << std::endl;
756756

757757
// Cleanup
758758
if(graphFinalCreated)

‎docs/tools/example_codes/graph_api_tutorial_main_streams.hip‎

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -146,7 +146,7 @@ auto main() -> int
146146
<< volGeom.dim.z << std::endl;
147147

148148
// Initialize per-stream data
149-
for(auto streamIdx = 0; streamIdx < streams.size(); ++streamIdx)
149+
for(auto streamIdx = 0u; streamIdx < streams.size(); ++streamIdx)
150150
{
151151
std::cout << "Initializing stream " << streamIdx << "... " << std::flush;
152152
auto stream = streams.at(streamIdx);
@@ -259,7 +259,7 @@ auto main() -> int
259259
auto start = std::chrono::steady_clock::now();
260260

261261
// [sphinx-batch-start]
262-
auto projIdx = 0;
262+
auto projIdx = 0u;
263263
while(projIdx < projGeom.numProj)
264264
{
265265
auto batchSize = std::min(numStreams, static_cast<int>(projGeom.numProj - projIdx));
@@ -427,7 +427,7 @@ auto main() -> int
427427
// [sphinx-sync-start]
428428
// First stream waits for other streams to complete
429429
auto completionEvents = std::vector<hipEvent_t>{};
430-
for(auto streamIdx = 1; streamIdx < streams.size(); ++streamIdx)
430+
for(auto streamIdx = 1u; streamIdx < streams.size(); ++streamIdx)
431431
{
432432
auto event = hipEvent_t{};
433433
hip_check(hipEventCreate(&event));
@@ -436,7 +436,7 @@ auto main() -> int
436436
}
437437

438438
for(auto&& event : completionEvents)
439-
hip_check(hipStreamWaitEvent(streams.at(0), event));
439+
hip_check(hipStreamWaitEvent(streams.at(0), event, 0));
440440
// [sphinx-sync-end]
441441

442442
// Obtain reconstruction time before copying back the result
@@ -472,7 +472,7 @@ auto main() -> int
472472
std::cout << "Done!" << std::endl;
473473

474474
auto const elapsed = std::chrono::duration<double>{stop - start};
475-
std::cout << "Reconstruction time: " << elapsed << std::endl;
475+
std::cout << "Reconstruction time: " << elapsed.count() << 's' << std::endl;
476476

477477
for(auto&& event : completionEvents)
478478
hip_check(hipEventDestroy(event));

‎docs/tools/update_example_codes.py‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -323,14 +323,14 @@
323323

324324
# graph_api
325325
urllib.request.urlretrieve(
326-
"https://raw.githubusercontent.com/ROCm/rocm-examples/amd-staging/HIP-Doc/Tutorials/graph_api/src/main_streams.hip"
326+
"https://raw.githubusercontent.com/ROCm/rocm-examples/amd-staging/HIP-Doc/Tutorials/graph_api/src/main_streams.hip",
327327
"docs/tools/example_codes/graph_api_tutorial_main_streams.hip"
328328
)
329329
urllib.request.urlretrieve(
330-
"https://raw.githubusercontent.com/ROCm/rocm-examples/amd-staging/HIP-Doc/Tutorials/graph_api/src/main_graph_capture.hip"
330+
"https://raw.githubusercontent.com/ROCm/rocm-examples/amd-staging/HIP-Doc/Tutorials/graph_api/src/main_graph_capture.hip",
331331
"docs/tools/example_codes/graph_api_tutorial_main_graph_capture.hip"
332332
)
333333
urllib.request.urlretrieve(
334-
"https://raw.githubusercontent.com/ROCm/rocm-examples/amd-staging/HIP-Doc/Tutorials/graph_api/src/main_graph_creation.hip"
334+
"https://raw.githubusercontent.com/ROCm/rocm-examples/amd-staging/HIP-Doc/Tutorials/graph_api/src/main_graph_creation.hip",
335335
"docs/tools/example_codes/graph_api_tutorial_main_graph_creation.hip"
336336
)

‎docs/tutorial/graph_api.rst‎

Lines changed: 9 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -2,6 +2,8 @@
22
:description: HIP graph API tutorial
33
:keywords: AMD, ROCm, HIP, graph API, tutorial
44

5+
.. _hip_graph_api_tutorial:
6+
57
*******************************************************************************
68
HIP Graph API Tutorial
79
*******************************************************************************
@@ -108,12 +110,14 @@ The reconstruction pipeline consists of:
108110

109111
1. **Load** projection data into GPU memory
110112
2. **Preprocess** the projection through six stages:
113+
111114
a. Logarithmic transformation (convert X-ray intensities)
112115
b. Pixel weighting (correct for cone-beam geometry)
113116
c. Forward FFT (transform to frequency domain)
114117
d. Shepp-Logan filtering (enhance edges and improve contrast)
115118
e. Inverse FFT (return to spatial domain)
116119
f. Normalization (account for unnormalized FFT)
120+
117121
3. **Reconstruct** the 3D volume using the Feldkamp-Davis-Kress (FDK) algorithm [FeDK84]_
118122

119123
**Why HIP graphs?** CT scanners process hundreds of projections per scan. By capturing this fixed workflow as a graph,
@@ -175,7 +179,7 @@ concepts.
175179
prefixed with ``main_``.
176180

177181
Step 1: Build the tutorial code
178-
==========================
182+
===============================
179183

180184
The full code for this tutorial is part of the `ROCm examples repository <https://github.com/ROCm/rocm-examples>`__.
181185
Check out the repository:
@@ -332,7 +336,7 @@ Inside the ``build`` directory you will now generate a trace:
332336
333337
.. note::
334338
For more information on the ``rocprofv3`` tool, please refer to its
335-
:ref:`documentation <rocprofiler-sdk:using-rocprov3>`.
339+
:ref:`documentation <rocprofiler-sdk:using-rocprofv3>`.
336340

337341
Analyzing the trace
338342
^^^^^^^^^^^^^^^^^^^
@@ -381,6 +385,7 @@ Inside the main loop, activate capture mode on the first stream:
381385
:dedent:
382386

383387
.. admonition:: What happens during capture?
388+
384389
When :cpp:func:`hipStreamBeginCapture` is called, the stream stops executing operations immediately. Instead, it
385390
records operations into a graph template (``graph`` in the code shown here).
386391

@@ -658,5 +663,5 @@ Resources
658663

659664
.. rubric:: References
660665

661-
.. [FeDK84] L. A. Feldkamp, L. C. Davis and J. W. Kress: "Practical cone-beam algorithm". In *Journal of the Optical Society of America A*, vol. 1, no. 6, pp. 612-619, June 1984, DOI `10.1364/JOSAA.1.000612 <https://dx.doi.org/10.1364/JOSAA.1.000612>`__.
662-
.. [ShLo74] L. A. Shepp and B. F. Logan: "The Fourier reconstruction of a head section". In *IEEE Transactions on Nuclear Science*, vol. 21, no. 3, pp. 21-43, June 1974, DOI `10.1109/TNS.1974.6499235 <https://dx.doi.org/10.1109/TNS.1974.6499235>`__.
666+
.. [FeDK84] L.A. Feldkamp, L.C. Davis and J.W. Kress: "Practical cone-beam algorithm". In *Journal of the Optical Society of America A*, vol. 1, no. 6, pp. 612-619, June 1984, DOI `10.1364/JOSAA.1.000612 <https://dx.doi.org/10.1364/JOSAA.1.000612>`__.
667+
.. [ShLo74] L.A. Shepp and B.F. Logan: "The Fourier reconstruction of a head section". In *IEEE Transactions on Nuclear Science*, vol. 21, no. 3, pp. 21-43, June 1974, DOI `10.1109/TNS.1974.6499235 <https://dx.doi.org/10.1109/TNS.1974.6499235>`__.

0 commit comments

Comments
 (0)