AI/rocm-systems
2
0 Fork 0
Koodi Ongelmat Pull-pyynnöt Actions Paketit Projektit Julkaisut Wiki Toiminta

SWDEV-397576 Adding protection for HIP calls to understand MI failures

Selaa lähdekoodia 
Change-Id: Icd6a2cb3be5520ee57fdd76bfafa06801b7061c0
This commit is contained in:
gobhardw authored and Gopesh Bhardwaj committed 2023-05-11 13:02:31 -04:00
1 vanhempi 55c6a2d4ad
commit 0b401d71c6
5 files changed
+116 -116
Download Patch File Download Diff File Expand all files Collapse all files

No files matched your search

tests/featuretests/profiler/apps/hello_world_omp.cpp
+11 -20
Näytä tiedosto
@@ -31,39 +31,31 @@ THE SOFTWARE.
// HIP header
#include <hip/hip_runtime.h>
#include "utils/test_helper.h"
#define NUM_THREADS 16
#define CHECK(cmd) \
{ \
hipError_t error = cmd; \
if (error != hipSuccess) { \
fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), \
error, __FILE__, __LINE__); \
exit(EXIT_FAILURE); \
} \
}
__global__ void hip_helloworld(unsigned omp_id, int *A_d) {
__global__ void hip_helloworld(unsigned omp_id, int* A_d) {
// Note: the printf command will only work if printf is enabled in your build.
// printf("Hello World... from HIP thread = %u\n", omp_id);
A_d[omp_id] = omp_id;
}
int main(int argc, char *argv[]) {
int main(int argc, char* argv[]) {
int *A_h, *A_d;
size_t Nbytes = NUM_THREADS * sizeof(int);
hipDeviceProp_t props;
CHECK(hipGetDeviceProperties(&props, 0 /*deviceID*/));
HIP_RC(hipGetDeviceProperties(&props, 0 /*deviceID*/));
// printf("info: running on device %s\n", props.name);
A_h = reinterpret_cast<int *>(malloc(Nbytes));
CHECK(hipMalloc(&A_d, Nbytes));
A_h = reinterpret_cast<int*>(malloc(Nbytes));
HIP_RC(hipMalloc(&A_d, Nbytes));
for (int i = 0; i < NUM_THREADS; i++) {
A_h[i] = 0;
}
CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
HIP_RC(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
// Beginning of parallel region
#pragma omp parallel num_threads(NUM_THREADS)
@@ -71,13 +63,12 @@ int main(int argc, char *argv[]) {
// fprintf(stderr, "Hello World... from OMP thread = %d\n",
// omp_get_thread_num());
hipLaunchKernelGGL(hip_helloworld, dim3(1), dim3(1), 0, 0,
omp_get_thread_num(), A_d);
HIP_KL(hipLaunchKernelGGL(hip_helloworld, dim3(1), dim3(1), 0, 0, omp_get_thread_num(), A_d));
}
// Ending of parallel region
hipStreamSynchronize(0);
CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
HIP_RC(hipStreamSynchronize(0));
HIP_RC(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
// printf("Device Results:\n");
for (int i = 0; i < NUM_THREADS; i++) {
// printf(" A_d[%d] = %d\n", i, A_h[i]);
@@ -86,6 +77,6 @@ int main(int argc, char *argv[]) {
printf("PASSED!\n");
free(A_h);
CHECK(hipFree(A_d));
HIP_RC(hipFree(A_d));
return 0;
}