Skip to content

Commit 84d487b

Browse files
jrmadsenanushe@amd.com
jrmadsen
authored and
anushe@amd.com
committed
Add reproducer for hipGraphLaunch GPU activity bubbles
1 parent 08d01e4 commit 84d487b

File tree

3 files changed

+182
-0
lines changed

3 files changed

+182
-0
lines changed

‎projects/rocprofiler-sdk/tests/bin/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -45,3 +45,4 @@ add_subdirectory(attachment-test)
4545
add_subdirectory(hip-host)
4646
add_subdirectory(module-loading-test)
4747
add_subdirectory(late-start-tracing)
48+
add_subdirectory(hip-graph-bubbles)

‎projects/rocprofiler-sdk/tests/bin/hip-graph-bubbles/CMakeLists.txt‎

Lines changed: 46 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,46 @@
1+
#
2+
#
3+
#
4+
cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR)
5+
6+
if(NOT CMAKE_HIP_COMPILER)
7+
find_program(
8+
amdclangpp_EXECUTABLE
9+
NAMES amdclang++
10+
HINTS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm
11+
PATHS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm
12+
PATH_SUFFIXES bin llvm/bin NO_CACHE)
13+
mark_as_advanced(amdclangpp_EXECUTABLE)
14+
15+
if(amdclangpp_EXECUTABLE)
16+
set(CMAKE_HIP_COMPILER "${amdclangpp_EXECUTABLE}")
17+
endif()
18+
endif()
19+
20+
project(rocprofiler-sdk-tests-bin-hip-graph-bubbles LANGUAGES CXX HIP)
21+
22+
foreach(_TYPE DEBUG MINSIZEREL RELEASE RELWITHDEBINFO)
23+
if("${CMAKE_HIP_FLAGS_${_TYPE}}" STREQUAL "")
24+
set(CMAKE_HIP_FLAGS_${_TYPE} "${CMAKE_CXX_FLAGS_${_TYPE}}")
25+
endif()
26+
endforeach()
27+
28+
set(CMAKE_CXX_STANDARD 17)
29+
set(CMAKE_CXX_EXTENSIONS OFF)
30+
set(CMAKE_CXX_STANDARD_REQUIRED ON)
31+
set(CMAKE_HIP_STANDARD 17)
32+
set(CMAKE_HIP_EXTENSIONS OFF)
33+
set(CMAKE_HIP_STANDARD_REQUIRED ON)
34+
35+
set_source_files_properties(hip-graph-bubbles.cpp PROPERTIES LANGUAGE HIP)
36+
add_executable(hip-graph-bubbles)
37+
target_sources(hip-graph-bubbles PRIVATE hip-graph-bubbles.cpp)
38+
target_compile_options(hip-graph-bubbles PRIVATE -W -Wall -Wextra -Wpedantic -Wshadow
39+
-Werror)
40+
41+
find_package(Threads REQUIRED)
42+
target_link_libraries(hip-graph-bubbles PRIVATE Threads::Threads)
43+
44+
find_package(rocprofiler-sdk-roctx REQUIRED)
45+
target_link_libraries(hip-graph-bubbles
46+
PRIVATE rocprofiler-sdk-roctx::rocprofiler-sdk-roctx)

‎projects/rocprofiler-sdk/tests/bin/hip-graph-bubbles/hip-graph-bubbles.cpp‎

Lines changed: 135 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,135 @@
1+
/*
2+
Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved.
3+
4+
Permission is hereby granted, free of charge, to any person obtaining a copy
5+
of this software and associated documentation files (the "Software"), to deal
6+
in the Software without restriction, including without limitation the rights
7+
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
8+
copies of the Software, and to permit persons to whom the Software is
9+
furnished to do so, subject to the following conditions:
10+
11+
The above copyright notice and this permission notice shall be included in
12+
all copies or substantial portions of the Software.
13+
14+
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
15+
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
16+
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
17+
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
18+
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
19+
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
20+
THE SOFTWARE.
21+
*/
22+
23+
#include <hip/hip_runtime.h>
24+
#include <rocprofiler-sdk-roctx/roctx.h>
25+
26+
#include <chrono>
27+
#include <iomanip>
28+
#include <iostream>
29+
#include <vector>
30+
31+
#define HIP_CHECK(cmd) \
32+
{ \
33+
hipError_t error = cmd; \
34+
if(error != hipSuccess) \
35+
{ \
36+
std::cerr << "HIP error: " << hipGetErrorString(error) << " at " << __FILE__ << ":" \
37+
<< __LINE__ << std::endl; \
38+
exit(EXIT_FAILURE); \
39+
} \
40+
}
41+
42+
// Simple kernel that does minimal work
43+
__global__ void
44+
simpleKernel(int* data, int value)
45+
{
46+
int idx = blockIdx.x * blockDim.x + threadIdx.x;
47+
data[idx] = value + idx;
48+
}
49+
50+
int
51+
main()
52+
{
53+
const int NUM_KERNELS = 2000;
54+
const int NUM_ITERATIONS = 200;
55+
const int ARRAY_SIZE = 256;
56+
57+
std::cout << "Creating HIP graph with " << NUM_KERNELS << " kernel launches" << std::endl;
58+
std::cout << "Will execute graph " << NUM_ITERATIONS << " times" << std::endl;
59+
60+
// Allocate device memory
61+
int* d_data;
62+
HIP_CHECK(hipMalloc(&d_data, ARRAY_SIZE * sizeof(int)));
63+
64+
// Create graph
65+
hipGraph_t graph;
66+
HIP_CHECK(hipGraphCreate(&graph, 0));
67+
68+
// Create stream for graph capture
69+
hipStream_t stream;
70+
HIP_CHECK(hipStreamCreate(&stream));
71+
72+
// Begin graph capture
73+
HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
74+
75+
// Launch many kernels
76+
dim3 blockSize(256);
77+
dim3 gridSize(1);
78+
79+
for(int i = 0; i < NUM_KERNELS; i++)
80+
{
81+
hipLaunchKernelGGL(simpleKernel, gridSize, blockSize, 0, stream, d_data, i);
82+
}
83+
84+
// End graph capture
85+
HIP_CHECK(hipStreamEndCapture(stream, &graph));
86+
87+
// Create executable graph
88+
hipGraphExec_t graphExec;
89+
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0));
90+
91+
std::cout << "Graph created and instantiated successfully" << std::endl;
92+
std::cout << "Starting graph execution loop..." << std::endl;
93+
94+
// Start timing
95+
auto start = std::chrono::high_resolution_clock::now();
96+
97+
// Execute the graph multiple times
98+
for(int iter = 0; iter < NUM_ITERATIONS; iter++)
99+
{
100+
roctxRangePush("graph_launch");
101+
HIP_CHECK(hipGraphLaunch(graphExec, stream));
102+
roctxRangePop();
103+
104+
if((iter + 1) % 50 == 0)
105+
{
106+
std::cout << "Completed " << (iter + 1) << " iterations" << std::endl;
107+
}
108+
}
109+
110+
// Wait for completion
111+
HIP_CHECK(hipStreamSynchronize(stream));
112+
113+
// End timing
114+
auto end = std::chrono::high_resolution_clock::now();
115+
std::chrono::duration<double> elapsed = end - start;
116+
117+
std::cout << "All iterations completed successfully!" << std::endl;
118+
std::cout << std::fixed << std::setprecision(4);
119+
std::cout << "\n=== Timing Results ===" << std::endl;
120+
std::cout << "Total execution time: " << elapsed.count() << " seconds" << std::endl;
121+
std::cout << "Total kernel launches: " << (NUM_KERNELS * NUM_ITERATIONS) << std::endl;
122+
std::cout << "Average time per iteration: " << (elapsed.count() / NUM_ITERATIONS) << " seconds"
123+
<< std::endl;
124+
std::cout << "======================" << std::endl;
125+
126+
// Cleanup
127+
HIP_CHECK(hipGraphExecDestroy(graphExec));
128+
HIP_CHECK(hipGraphDestroy(graph));
129+
HIP_CHECK(hipStreamDestroy(stream));
130+
HIP_CHECK(hipFree(d_data));
131+
132+
std::cout << "Test completed successfully" << std::endl;
133+
134+
return 0;
135+
}

0 commit comments

Comments
 (0)