Benjamin Welton
d05031db89
* Limit the number of HSA signals that are active
There is a hard limit currently to the number of
signals that HSA allows to be created (before weird stuff
happens such as hangs or straight up crashes in HSA). While
there is some work going on to fix this in HSA/AQL. Lets limit the
number we create.
Increased the counter colleciton example to 200K launches, which
with this change no longer hangs/crashes randomly in HSA.
* source formatting (clang-format v11) (#142)
Co-authored-by: bwelton <bwelton@users.noreply.github.com>
* Up timout
---------
Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>
Co-authored-by: bwelton <bwelton@users.noreply.github.com>
[ROCm/rocprofiler-sdk commit: de685246a7]
93 строки
2.8 KiB
C++
93 строки
2.8 KiB
C++
#include <hip/hip_runtime.h>
|
|
|
|
#include "client.hpp"
|
|
|
|
#define HIP_CALL(call) \
|
|
do \
|
|
{ \
|
|
hipError_t err = call; \
|
|
if(err != hipSuccess) \
|
|
{ \
|
|
fprintf(stderr, "%s\n", hipGetErrorString(err)); \
|
|
abort(); \
|
|
} \
|
|
} while(0)
|
|
|
|
__global__ void
|
|
kernelA(int x, int y)
|
|
{
|
|
x = x + y;
|
|
}
|
|
|
|
__global__ void
|
|
kernelB(int x, int y)
|
|
{
|
|
x = x + y;
|
|
}
|
|
|
|
template <typename T>
|
|
__global__ void
|
|
kernelC(T* C_d, const T* A_d, size_t N)
|
|
{
|
|
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
|
|
size_t stride = blockDim.x * gridDim.x;
|
|
for(size_t i = offset; i < N; i += stride)
|
|
{
|
|
C_d[i] = A_d[i] * A_d[i];
|
|
}
|
|
}
|
|
|
|
void
|
|
launchKernals()
|
|
{
|
|
const int NUM_LAUNCH = 200000;
|
|
// Normal HIP Calls
|
|
int* gpuMem;
|
|
[[maybe_unused]] hipDeviceProp_t devProp;
|
|
HIP_CALL(hipGetDeviceProperties(&devProp, 0));
|
|
HIP_CALL(hipMalloc((void**) &gpuMem, 1 * sizeof(int)));
|
|
|
|
for(int i = 0; i < NUM_LAUNCH; i++)
|
|
{
|
|
// KernelA and KernelB to be profiled as part of the session
|
|
hipLaunchKernelGGL(kernelA, dim3(1), dim3(1), 0, 0, 1, 2);
|
|
hipLaunchKernelGGL(kernelB, dim3(1), dim3(1), 0, 0, 1, 2);
|
|
}
|
|
|
|
const int NElems = 512 * 512;
|
|
const int Nbytes = NElems * 2;
|
|
int * A_d, *C_d;
|
|
int A_h[NElems], C_h[NElems];
|
|
|
|
for(int i = 0; i < NElems; i++)
|
|
{
|
|
A_h[i] = i;
|
|
}
|
|
|
|
HIP_CALL(hipDeviceSynchronize());
|
|
|
|
HIP_CALL(hipMalloc(&A_d, Nbytes));
|
|
HIP_CALL(hipMalloc(&C_d, Nbytes));
|
|
HIP_CALL(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
|
|
HIP_CALL(hipDeviceSynchronize());
|
|
const unsigned blocks = 512;
|
|
const unsigned threadsPerBlock = 256;
|
|
for(int i = 0; i < NUM_LAUNCH; i++)
|
|
{
|
|
hipLaunchKernelGGL(kernelC, dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, NElems);
|
|
}
|
|
HIP_CALL(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
|
|
HIP_CALL(hipDeviceSynchronize());
|
|
HIP_CALL(hipFree(gpuMem));
|
|
HIP_CALL(hipFree(A_d));
|
|
HIP_CALL(hipFree(C_d));
|
|
std::cerr << "Run complete\n";
|
|
}
|
|
|
|
int
|
|
main()
|
|
{
|
|
start();
|
|
launchKernals();
|
|
}
|