AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[rocprofiler-sdk] Fix Stream ID Error for Attachment (#1142)

Browse Source 
* Changed stream error warning, remove regex search from attach execute test

* Formatting

* Revert accidental change

* Fix stream hang error due to grabbing same lock twice

* Updated add stream code, need to update tests

* Update attachment tests to use streams, threads, and multiple devices

* Update tests and fix stream issues

* Updated error messages to be more explicit, updated json to csv code in conftest to include streams and threads

* Formatting

* Add attachment label to attachment tests and update validation to fix errors

* Fix attach twice conftest

* Disabled thread san tests for attachment since they no longer work with bin file changes

* Updated for comment

* Added null check for getting attach status
This commit is contained in:
itrowbri
2025-10-17 16:34:05 -05:00
committed by GitHub
parent 620ccbeb82
commit e7a26594b7
10 changed files with 219 additions and 90 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rocprofiler-sdk/tests/bin/attachment-test/attachment_test.cpp
+101 -47
View File
@@ -24,10 +24,25 @@
#include <rocprofiler-sdk-roctx/roctx.h>
#include <unistd.h>
#include <chrono>
#include <cstdlib>
#include <iostream>
#include <string>
#include <thread>
#include <vector>
/* Macro for checking GPU API return values */
#define HIP_ASSERT(call) \
do \
{ \
hipError_t gpuErr = call; \
if(hipSuccess != gpuErr) \
{ \
printf( \
"GPU API Error - %s:%d: '%s'\n", __FILE__, __LINE__, hipGetErrorString(gpuErr)); \
exit(1); \
} \
} while(0)
__global__ void
simple_kernel(float* data, int size)
{
@@ -38,44 +53,23 @@ simple_kernel(float* data, int size)
}
}
int
main(int /*argc*/, char** /*argv*/)
void
execute_kernels(const size_t tid,
const hipStream_t stream,
const size_t stream_id,
const size_t device_id)
{
std::cout << "Attachment test app started with PID: " << getpid() << std::endl;
// Initialize HIP
int device_count = 0;
hipError_t err = hipGetDeviceCount(&device_count);
if(err != hipSuccess || device_count == 0)
{
std::cerr << "No HIP devices found or error getting device count" << std::endl;
return 1;
}
std::cout << "After first call " << getpid() << std::endl;
// Set device
err = hipSetDevice(0);
if(err != hipSuccess)
{
std::cerr << "Failed to set device 0" << std::endl;
return 1;
}
HIP_ASSERT(hipSetDevice(device_id));
// Allocate memory
const int size = 1024 * 1024; // 1M elements
const size_t bytes = size * sizeof(float);
float* h_data = new float[size];
float* d_data;
float* d_data = nullptr;
err = hipMalloc(&d_data, bytes);
if(err != hipSuccess)
{
std::cerr << "Failed to allocate device memory" << std::endl;
delete[] h_data;
return 1;
}
HIP_ASSERT(hipMalloc(&d_data, bytes));
// Initialize data
for(int i = 0; i < size; ++i)
@@ -84,7 +78,8 @@ main(int /*argc*/, char** /*argv*/)
}
// Run kernels in a loop for a while
std::cout << "Starting kernel execution loop..." << std::endl;
std::cout << "Starting kernel execution loop for thread " << tid << " with stream " << stream_id
<< " on device " << device_id << "...\n";
const int num_iterations = 30;
for(int iter = 0; iter < num_iterations; ++iter)
@@ -95,10 +90,11 @@ main(int /*argc*/, char** /*argv*/)
// Copy data to device
roctxMark("Start_H2D_Copy");
err = hipMemcpy(d_data, h_data, bytes, hipMemcpyHostToDevice);
auto err = hipMemcpyAsync(d_data, h_data, bytes, hipMemcpyHostToDevice, stream);
if(err != hipSuccess)
{
std::cerr << "Failed to copy data to device" << std::endl;
std::cerr << "Failed to copy data for thread " << tid << " with stream " << stream_id
<< " on device " << device_id << "...\n";
roctxRangePop(); // Removed - ROCTx not linked
break;
}
@@ -109,46 +105,104 @@ main(int /*argc*/, char** /*argv*/)
int blocks_per_grid = (size + threads_per_block - 1) / threads_per_block;
hipLaunchKernelGGL(
simple_kernel, dim3(blocks_per_grid), dim3(threads_per_block), 0, 0, d_data, size);
simple_kernel, dim3(blocks_per_grid), dim3(threads_per_block), 0, stream, d_data, size);
// Copy data back
roctxMark("Start_D2H_Copy");
err = hipMemcpy(h_data, d_data, bytes, hipMemcpyDeviceToHost);
err = hipMemcpyAsync(h_data, d_data, bytes, hipMemcpyDeviceToHost, stream);
if(err != hipSuccess)
{
std::cerr << "Failed to copy data from device" << std::endl;
std::cerr << "Failed to copy data for thread " << tid << " with stream " << stream_id
<< " on device " << device_id << "...\n";
roctxRangePop(); // Removed - ROCTx not linked
break;
}
// Wait for completion
roctxMark("Device_Synchronize");
err = hipDeviceSynchronize();
roctxMark("Stream_Synchronize");
err = hipStreamSynchronize(stream);
if(err != hipSuccess)
{
std::cerr << "Failed to synchronize device" << std::endl;
std::cerr << "Failed to synchronize stream " << stream_id << " with thread " << tid
<< " on device " << device_id << "...\n";
roctxRangePop(); // Removed - ROCTx not linked
break;
}
roctxRangePop(); // Removed - ROCTx not linked
std::cout << "Iteration " << (iter + 1) << "/" << num_iterations << " completed"
<< std::endl;
// Small delay between iterations
std::this_thread::sleep_for(std::chrono::milliseconds(500));
}
std::cout << "Kernel execution loop completed" << std::endl;
std::cout << "Kernel execution loop completed for thread " << tid << " with stream "
<< stream_id << " on device " << device_id << "...\n";
// Cleanup
err = hipFree(d_data);
if(err != hipSuccess)
{
std::cerr << "Warning: Failed to free device memory" << std::endl;
}
HIP_ASSERT(hipFree(d_data));
delete[] h_data;
}
int
main(int argc, char** argv)
{
size_t nthreads{32};
size_t nstreams{8};
int ndevices{0};
for(int i = 1; i < argc; ++i)
{
auto _arg = std::string{argv[i]};
if(_arg == "?" || _arg == "-h" || _arg == "--help")
{
fprintf(stderr,
"usage: attachment-test [NUM_THREADS (%zu)] [NUM_STREAMS (%zu)] "
"[NUM_DEVICES (%d)]\n",
nthreads,
nstreams,
ndevices);
exit(EXIT_SUCCESS);
}
}
if(argc > 1) nthreads = std::atoll(argv[1]);
if(argc > 2) nstreams = std::atoll(argv[2]);
if(argc > 3) ndevices = std::stoi(argv[3]);
std::cout << "Attachment test app started with PID: " << getpid() << std::endl;
// Initialize HIP
int device_count = 0;
HIP_ASSERT(hipGetDeviceCount(&device_count));
if(device_count == 0)
{
std::cerr << "No HIP devices found or error getting device count" << std::endl;
return 1;
}
// Default ndecives to device_count. Ensure that we do not use more devices than are available
ndevices = ndevices == 0 ? device_count : ndevices;
if(ndevices > device_count)
{
std::cout << "Using " << device_count << " HIP devices instead of the requested "
<< ndevices << "\n";
ndevices = device_count;
}
std::cout << "After first call " << getpid() << std::endl;
auto _threads = std::vector<std::thread>{};
auto _streams = std::vector<hipStream_t>(nstreams);
_threads.reserve(nthreads);
for(auto& itr : _streams)
HIP_ASSERT(hipStreamCreate(&itr));
for(size_t i = 0; i < nthreads; ++i)
_threads.emplace_back(
execute_kernels, i, _streams.at(i % nstreams), i % nstreams, i % ndevices);
for(auto& itr : _threads)
itr.join();
// Destroy streams
for(auto itr : _streams)
HIP_ASSERT(hipStreamDestroy(itr));
std::cout << "Attachment test app finished" << std::endl;