Benjamin Welton
de685246a7
* 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>
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();
|
|
}
|