diff --git a/gdb/testsuite/gdb.perf/rocm-break-cond-false.cpp b/gdb/testsuite/gdb.perf/rocm-break-cond-false.cpp index 773ae1a391d..004ff2ce5c0 100644 --- a/gdb/testsuite/gdb.perf/rocm-break-cond-false.cpp +++ b/gdb/testsuite/gdb.perf/rocm-break-cond-false.cpp @@ -18,6 +18,17 @@ #include +#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 kernel () { @@ -31,7 +42,7 @@ main (int argc, char **argv) 0 /*dynamicShared*/, 0 /*stream*/); /* Wait until kernel finishes. */ - hipDeviceSynchronize (); + CHECK (hipDeviceSynchronize ()); return 0; } diff --git a/gdb/testsuite/gdb.rocm/aspace-watchpoint.cpp b/gdb/testsuite/gdb.rocm/aspace-watchpoint.cpp index 04a610cbab7..c32217c57a7 100644 --- a/gdb/testsuite/gdb.rocm/aspace-watchpoint.cpp +++ b/gdb/testsuite/gdb.rocm/aspace-watchpoint.cpp @@ -26,6 +26,17 @@ #include +#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); \ + } \ + } + __device__ void change_memory (char *private_ptr, int *global_ptr) { @@ -74,11 +85,12 @@ int main (int argc, char* argv[]) { int *global_ptr; - hipMalloc (&global_ptr, 4); + CHECK (hipMalloc (&global_ptr, 4)); int init_value_h = 0; - hipMemcpy (global_ptr, &init_value_h, sizeof (init_value_h), hipMemcpyHostToDevice); + CHECK (hipMemcpy (global_ptr, &init_value_h, sizeof (init_value_h), + hipMemcpyHostToDevice)); kernel<<<1, 1, 0>>> (global_ptr); - hipDeviceSynchronize (); + CHECK (hipDeviceSynchronize ()); return 0; } diff --git a/gdb/testsuite/gdb.rocm/branch-fault.cpp b/gdb/testsuite/gdb.rocm/branch-fault.cpp index cf2586c2c21..d169fc867e2 100644 --- a/gdb/testsuite/gdb.rocm/branch-fault.cpp +++ b/gdb/testsuite/gdb.rocm/branch-fault.cpp @@ -17,6 +17,17 @@ #include +#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); \ + } \ + } + typedef void (*func_ptr) (); __global__ void @@ -30,6 +41,6 @@ int main (int argc, char* argv[]) { hipLaunchKernelGGL (kernel, dim3 (1), dim3 (1), 0, 0); - hipDeviceSynchronize (); + CHECK (hipDeviceSynchronize ()); return 0; } diff --git a/gdb/testsuite/gdb.rocm/breakpoint-after-exit.cpp b/gdb/testsuite/gdb.rocm/breakpoint-after-exit.cpp index 2a667460aa0..0d8bf65dbfa 100644 --- a/gdb/testsuite/gdb.rocm/breakpoint-after-exit.cpp +++ b/gdb/testsuite/gdb.rocm/breakpoint-after-exit.cpp @@ -19,6 +19,17 @@ #include #include +#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 the_kernel () {} @@ -27,7 +38,7 @@ int main () { hipLaunchKernelGGL (the_kernel, dim3 (1), dim3 (1), 0, 0); - hipDeviceSynchronize (); /* set breakpoint here */ + CHECK (hipDeviceSynchronize ()); /* set breakpoint here */ return 0; } diff --git a/gdb/testsuite/gdb.rocm/deep-stack.cpp b/gdb/testsuite/gdb.rocm/deep-stack.cpp index ff8065276a9..f1c6edabf7e 100644 --- a/gdb/testsuite/gdb.rocm/deep-stack.cpp +++ b/gdb/testsuite/gdb.rocm/deep-stack.cpp @@ -19,6 +19,17 @@ #include #include +#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); \ + } \ + } + __device__ void base_case () { @@ -52,7 +63,7 @@ int main () { hipLaunchKernelGGL (HIP_KERNEL_NAME (hip_deep), dim3 (1), dim3 (1), 0, 0); - hipDeviceSynchronize (); + CHECK (hipDeviceSynchronize ()); return EXIT_SUCCESS; } diff --git a/gdb/testsuite/gdb.rocm/device-interrupt.cpp b/gdb/testsuite/gdb.rocm/device-interrupt.cpp index d941446e860..4b4534cf2c7 100644 --- a/gdb/testsuite/gdb.rocm/device-interrupt.cpp +++ b/gdb/testsuite/gdb.rocm/device-interrupt.cpp @@ -23,6 +23,17 @@ #include #include +#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); \ + } \ + } + /* Kernel entry point. Loop forever, while avoiding to peg the GPU. */ @@ -52,7 +63,7 @@ main (int argc, char **argv) /* Wait until kernel finishes. In this case, this blocks "forever". */ - hipDeviceSynchronize (); + CHECK (hipDeviceSynchronize ()); return 0; } diff --git a/gdb/testsuite/gdb.rocm/disassemble.cpp b/gdb/testsuite/gdb.rocm/disassemble.cpp index ab3a35a2a49..055bc461717 100644 --- a/gdb/testsuite/gdb.rocm/disassemble.cpp +++ b/gdb/testsuite/gdb.rocm/disassemble.cpp @@ -19,6 +19,17 @@ #include +#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); \ + } \ + } + /* Kernel entry point. */ __global__ void kernel () { @@ -33,7 +44,7 @@ main (int argc, char* argv[]) hipLaunchKernelGGL (kernel, dim3 (1), dim3 (1), 0 /*dynamicShared*/, 0 /*stream*/); - hipDeviceSynchronize (); + CHECK (hipDeviceSynchronize ()); return 0; } diff --git a/gdb/testsuite/gdb.rocm/displaced-stepping.cpp b/gdb/testsuite/gdb.rocm/displaced-stepping.cpp index b7482046dad..e2e19a8f274 100644 --- a/gdb/testsuite/gdb.rocm/displaced-stepping.cpp +++ b/gdb/testsuite/gdb.rocm/displaced-stepping.cpp @@ -19,6 +19,17 @@ #include +#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 kernel () { @@ -49,7 +60,7 @@ main (int argc, char **argv) 0 /*stream*/); /* Wait until kernel finishes. */ - hipDeviceSynchronize (); + CHECK (hipDeviceSynchronize ()); return 0; } diff --git a/gdb/testsuite/gdb.rocm/fork-exec-gpu-to-non-gpu-execer.cpp b/gdb/testsuite/gdb.rocm/fork-exec-gpu-to-non-gpu-execer.cpp index 4be5e49d6d8..b7ab040568d 100644 --- a/gdb/testsuite/gdb.rocm/fork-exec-gpu-to-non-gpu-execer.cpp +++ b/gdb/testsuite/gdb.rocm/fork-exec-gpu-to-non-gpu-execer.cpp @@ -18,6 +18,17 @@ #include #include +#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__ static void kernel1 () {} @@ -50,6 +61,6 @@ main (int argc, char* argv[]) hipLaunchKernelGGL (kernel2, dim3 (1), dim3 (1), 0, 0); - hipDeviceSynchronize (); + CHECK (hipDeviceSynchronize ()); return 0; } diff --git a/gdb/testsuite/gdb.rocm/fork-exec-non-gpu-to-gpu-execee.cpp b/gdb/testsuite/gdb.rocm/fork-exec-non-gpu-to-gpu-execee.cpp index eda280143c4..532cede59b5 100644 --- a/gdb/testsuite/gdb.rocm/fork-exec-non-gpu-to-gpu-execee.cpp +++ b/gdb/testsuite/gdb.rocm/fork-exec-non-gpu-to-gpu-execee.cpp @@ -17,6 +17,17 @@ #include +#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); \ + } \ + } + __device__ static void break_here_execee () {} @@ -30,7 +41,7 @@ kernel () int main (int argc, char* argv[]) { - hipLaunchKernelGGL (kernel, dim3 (1), dim3 (1), 0, 0); - hipDeviceSynchronize (); + kernel<<<1, 1>>> (); + CHECK (hipDeviceSynchronize ()); return 0; } diff --git a/gdb/testsuite/gdb.rocm/lane-execution.cpp b/gdb/testsuite/gdb.rocm/lane-execution.cpp index 33c0f6420af..52bf8708038 100644 --- a/gdb/testsuite/gdb.rocm/lane-execution.cpp +++ b/gdb/testsuite/gdb.rocm/lane-execution.cpp @@ -137,7 +137,7 @@ main () hipLaunchKernelGGL (kernel, grid_dim, block_dim, 0, 0, sInBuff, sOutBuff); - hipDeviceSynchronize (); + CHECK (hipDeviceSynchronize ()); return 0; } diff --git a/gdb/testsuite/gdb.rocm/lane-info.cpp b/gdb/testsuite/gdb.rocm/lane-info.cpp index 6ab1df2abc8..3ed789e5314 100644 --- a/gdb/testsuite/gdb.rocm/lane-info.cpp +++ b/gdb/testsuite/gdb.rocm/lane-info.cpp @@ -19,6 +19,17 @@ #include #include +#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); \ + } \ + } + /* The kernel never returns, via this sleep, so that the .exp file can test background execution (cont&). */ @@ -63,7 +74,7 @@ main () hipLaunchKernelGGL (kernel, dim3 (1), dim3 (64 + 5), 0 /*dynamicShared*/, 0 /*stream*/); - hipDeviceSynchronize (); + CHECK (hipDeviceSynchronize ()); return 0; } diff --git a/gdb/testsuite/gdb.rocm/line-breakpoint-in-kernel.cpp b/gdb/testsuite/gdb.rocm/line-breakpoint-in-kernel.cpp index 9caac71a057..c57b0256bce 100644 --- a/gdb/testsuite/gdb.rocm/line-breakpoint-in-kernel.cpp +++ b/gdb/testsuite/gdb.rocm/line-breakpoint-in-kernel.cpp @@ -18,6 +18,17 @@ #include +#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 kernel () { @@ -31,7 +42,7 @@ main (int argc, char **argv) hipLaunchKernelGGL (kernel, dim3 (1), dim3 (1), 0, 0); /* Wait until kernel finishes. */ - hipDeviceSynchronize (); + CHECK (hipDeviceSynchronize ()); return 0; } diff --git a/gdb/testsuite/gdb.rocm/mi-aspace.cpp b/gdb/testsuite/gdb.rocm/mi-aspace.cpp index d466d8b13df..4db35e7f660 100644 --- a/gdb/testsuite/gdb.rocm/mi-aspace.cpp +++ b/gdb/testsuite/gdb.rocm/mi-aspace.cpp @@ -18,6 +18,17 @@ #include +#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); \ + } \ + } + int global_var; __device__ void @@ -41,7 +52,7 @@ main () local#0. */ hipLaunchKernelGGL (kernel, grid_dim, block_dim, 1024, 0); - hipDeviceSynchronize (); + CHECK (hipDeviceSynchronize ()); return 0; } diff --git a/gdb/testsuite/gdb.rocm/multi-GPU.cpp b/gdb/testsuite/gdb.rocm/multi-GPU.cpp deleted file mode 100644 index 4f8a533712a..00000000000 --- a/gdb/testsuite/gdb.rocm/multi-GPU.cpp +++ /dev/null @@ -1,112 +0,0 @@ -/* Copyright (C) 2019-2023 Free Software Foundation, Inc. - Copyright (C) 2019-2023 Advanced Micro Devices, Inc. All rights reserved. - - This file is part of GDB. - - This program is free software; you can redistribute it and/or modify - it under the terms of the GNU General Public License as published by - the Free Software Foundation; either version 3 of the License, or - (at your option) any later version. - - This program is distributed in the hope that it will be useful, - but WITHOUT ANY WARRANTY; without even the implied warranty of - MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the - GNU General Public License for more details. - - You should have received a copy of the GNU General Public License - along with this program. If not, see . */ - -#include -#include -#include - -// Defining number of elements in Array -#define N 5 - -#define MAX_GPU 8 - -#define HIPCHECK(cmd) \ -do { \ - hipError_t error = (cmd); \ - if (error != hipSuccess) { \ - std::cerr << "Encountered HIP error (" << error << ") at line " \ - << __LINE__ << " in file " << __FILE__ << "\n"; \ - exit(-1); \ - } \ -} while (0) - -// Defining Kernel function for vector addition -__global__ void gpu_kernel_add(int *d_a, int *d_b, int *d_c) { - // Getting block index of current kernel - //int tid = blockIdx.x; // handle the data at this index - int tid = blockIdx.x * blockDim.x + threadIdx.x; - printf("\n%d\n",tid); - if (tid < N) - d_c[tid] = d_a[tid] + d_b[tid]; -} - -int main(void) -{ - // Defining host arrays - int h_a[N], h_b[N], h_c[N]; - // Defining device pointers - int *d_a[N], *d_b[N], *d_c[N]; - - hipStream_t stream[MAX_GPU]; - int nGpu = 1; - HIPCHECK(hipGetDeviceCount(&nGpu)); - - for (int i = 0; i < nGpu; i ++) { - HIPCHECK(hipSetDevice(i)); - hipDeviceProp_t prop; - HIPCHECK(hipGetDeviceProperties(&prop, i)); - printf("# device %d [0x%02x] %s\n", - i, prop.pciBusID, prop.name); - //create stream - HIPCHECK(hipStreamCreate(&stream[i])); - - hipMalloc((void**)&d_a[i], N * sizeof(int)); - hipMalloc((void**)&d_b[i], N * sizeof(int)); - hipMalloc((void**)&d_c[i], N * sizeof(int)); - - // Initializing Arrays - for (int i = 0; i < N; i++) { - h_a[i] = 2*i; - h_b[i] = i ; - } - - // Copy input arrays from host to device memory - hipMemcpyAsync(d_a[i], h_a, N * sizeof(int), hipMemcpyHostToDevice, stream[i]); - hipMemcpyAsync(d_b[i], h_b, N * sizeof(int), hipMemcpyHostToDevice, stream[i]); - } - - for (int i = 0; i < nGpu; i ++) { - HIPCHECK(hipSetDevice(i)); - // Calling kernels with N blocks and one thread per block, passing - // device pointers as parameters - hipLaunchKernelGGL(gpu_kernel_add, dim3(N), dim3(1 ), 0, stream[i], d_a[i], d_b[i], d_c[i]); - } - - for (int i = 0; i < nGpu; i ++) { - HIPCHECK(hipSetDevice(i)); - - // Copy result back to host memory from device memory - hipMemcpyAsync(h_c, d_c[i], N * sizeof(int), hipMemcpyDeviceToHost, stream[i]); - HIPCHECK(hipStreamSynchronize(stream[i])); - printf("Vector addition on GPU \n"); - - // Printing result on console - for (int i = 0; i < N; i++) { - printf("Operation result of %d element is %d + %d = %d\n", - i, h_a[i], h_b[i],h_c[i]); - } - - // Free up memory - HIPCHECK(hipStreamDestroy(stream[i])); - hipFree(d_a[i]); - hipFree(d_b[i]); - hipFree(d_c[i]); - } - return 0; -} - diff --git a/gdb/testsuite/gdb.rocm/multi-GPU.exp b/gdb/testsuite/gdb.rocm/multi-GPU.exp deleted file mode 100644 index 1fefa3b51ec..00000000000 --- a/gdb/testsuite/gdb.rocm/multi-GPU.exp +++ /dev/null @@ -1,28 +0,0 @@ -# Copyright (C) 2019-2023 Free Software Foundation, Inc. -# Copyright (C) 2019-2023 Advanced Micro Devices, Inc. All rights reserved. - -# This file is part of GDB. - -# This program is free software; you can redistribute it and/or modify -# it under the terms of the GNU General Public License as published by -# the Free Software Foundation; either version 3 of the License, or -# (at your option) any later version. - -# This program is distributed in the hope that it will be useful, -# but WITHOUT ANY WARRANTY; without even the implied warranty of -# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the -# GNU General Public License for more details. - -# You should have received a copy of the GNU General Public License -# along with this program. If not, see . - -load_lib rocm.exp - -require allow_hipcc_tests - -standard_testfile .cpp - -# Compile the hip program -if {[prepare_for_testing "failed to prepare ${testfile}" $testfile $srcfile {debug hip}]} { - return -1 -} diff --git a/gdb/testsuite/gdb.rocm/multi-inferior-fork.cpp b/gdb/testsuite/gdb.rocm/multi-inferior-fork.cpp index 3ce5f14e59f..d3e2afa6ee7 100644 --- a/gdb/testsuite/gdb.rocm/multi-inferior-fork.cpp +++ b/gdb/testsuite/gdb.rocm/multi-inferior-fork.cpp @@ -19,6 +19,17 @@ #include #include +#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 kernel () {} @@ -39,7 +50,7 @@ main () /* Child. */ child_after_fork (); hipLaunchKernelGGL (kernel, dim3 (1), dim3 (1), 0, 0); - hipDeviceSynchronize (); + CHECK (hipDeviceSynchronize ()); } return 0; diff --git a/gdb/testsuite/gdb.rocm/multi-inferior-gpu.cpp b/gdb/testsuite/gdb.rocm/multi-inferior-gpu.cpp index d64afdd1994..4ba406f6398 100644 --- a/gdb/testsuite/gdb.rocm/multi-inferior-gpu.cpp +++ b/gdb/testsuite/gdb.rocm/multi-inferior-gpu.cpp @@ -95,7 +95,7 @@ child (int argc, char **argv) CHECK (hipSetDevice (dev_number)); kern<<<1, 1>>> (); - hipDeviceSynchronize (); + CHECK (hipDeviceSynchronize ()); return 0; } diff --git a/gdb/testsuite/gdb.rocm/nonstop-displaced.cpp b/gdb/testsuite/gdb.rocm/nonstop-displaced.cpp index 18e224d0ad3..ae3837ecf87 100644 --- a/gdb/testsuite/gdb.rocm/nonstop-displaced.cpp +++ b/gdb/testsuite/gdb.rocm/nonstop-displaced.cpp @@ -17,6 +17,17 @@ #include +#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); \ + } \ + } + __device__ void foo () { @@ -44,6 +55,6 @@ int main (int argc, char *argv[]) { kern<<<2, 1>>> (); - hipDeviceSynchronize (); + CHECK (hipDeviceSynchronize ()); return 0; } diff --git a/gdb/testsuite/gdb.rocm/nonstop-mode.cpp b/gdb/testsuite/gdb.rocm/nonstop-mode.cpp index 3309c186f95..a52bdfd3ef2 100644 --- a/gdb/testsuite/gdb.rocm/nonstop-mode.cpp +++ b/gdb/testsuite/gdb.rocm/nonstop-mode.cpp @@ -16,103 +16,89 @@ You should have received a copy of the GNU General Public License along with this program. If not, see . */ -#include -#include +#include #include -// Number of elements in Array. -#define N 64 - -#define HIPCHECK(cmd) \ -do { \ - hipError_t error = (cmd); \ - if (error != hipSuccess) \ - { \ - std::cerr << "Encountered HIP error (" << error << ") at line " \ - << __LINE__ << " in file " << __FILE__ << "\n"; \ - exit(-1); \ - } \ -} while (0) - -#define MAX_GPU 8 - +#if !defined(GRID_DIM) +# error "Missing definition of GRID_DIM" +#endif +#if !defined(BLOCK_DIM) +# error "Missing definition of BLOCK_DIM" +#endif + +#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); \ + } \ + } +/* Number of elements in Array. */ +constexpr size_t N = 64; -// Defining Kernel function for vector addition -__global__ void VectorAdd(int *d_a, int *d_b, int *d_c) +__global__ void +VectorAdd (int *d_a, int *d_b, int *d_c) { - // Getting block index of current kernel - int tid = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; + size_t tid = blockIdx.x * blockDim.x + threadIdx.x; if (tid < N) d_c[tid] = d_a[tid] + d_b[tid]; } - -int main(void) +int +main () { - // Defining host arrays - int h_a[N], h_b[N], h_c[N]; - // Defining device pointers - int *d_a[N], *d_b[N], *d_c[N]; - // allocate the memory - - hipStream_t stream[MAX_GPU]; - - int nGpu = 1; - // To do - // In multi gpu scenario on running kernel on every GPU causing multiple - // failures in nonstop mode test cases, which need to investigate. - // HIPCHECK(hipGetDeviceCount(&nGpu)); - for (int i = 0; i < nGpu; i ++) { - HIPCHECK(hipSetDevice(i)); - hipDeviceProp_t prop; - HIPCHECK(hipGetDeviceProperties(&prop, i)); - printf("# device %d [0x%02x] %s\n", - i, prop.pciBusID, prop.name); - //create stream - HIPCHECK(hipStreamCreate(&stream[i])); - - hipMalloc((void**)&d_a[i], N * sizeof(int)); - hipMalloc((void**)&d_b[i], N * sizeof(int)); - hipMalloc((void**)&d_c[i], N * sizeof(int)); - // Initializing Arrays - for (int i = 0; i < N; i++) { - h_a[i] = 2*i; - h_b[i] = i ; + int gpu_id; + CHECK (hipGetDevice (&gpu_id)); + + hipDeviceProp_t prop; + CHECK (hipGetDeviceProperties (&prop, gpu_id)); + printf ("# device %d [%04x:%02x.%02x] %s\n", gpu_id, prop.pciDomainID, + prop.pciBusID, prop.pciDeviceID, prop.name); + + /* Host allocations. */ + int h_a[N] = {}; + int h_b[N] = {}; + int h_c[N] = {}; + for (int i = 0; i < N; ++i) + { + h_a[i] = 2 * i; + h_b[i] = i; } - // Copy input arrays from host to device memory - hipMemcpyAsync(d_a[i], h_a, N * sizeof(int), hipMemcpyHostToDevice, stream[i]); - hipMemcpyAsync(d_b[i], h_b, N * sizeof(int), hipMemcpyHostToDevice, stream[i]); - } - - for (int i = 0; i < nGpu; i ++) { - HIPCHECK(hipSetDevice(i)); + /* Device allocations. */ + int *d_a = nullptr; + int *d_b = nullptr; + int *d_c = nullptr; - hipLaunchKernelGGL(VectorAdd, - dim3(GRID_DIM), dim3(BLOCK_DIM), - 0, stream[i], d_a[i], d_b[i], d_c[i]); - } - - for (int i = 0; i < nGpu; i ++) { - HIPCHECK(hipSetDevice(i)); - // Copy result back to host memory from device memory - hipMemcpyAsync(h_c, d_c[i], N * sizeof(int), hipMemcpyDeviceToHost, stream[i]); - HIPCHECK(hipStreamSynchronize(stream[i])); - //printf("Vector addition on GPU \n"); - // Printing result on console - for (int i = 0; i < N; i++) { - /*printf("Operation result of %d element is %d + %d = %d\n", - i, h_a[i], h_b[i],h_c[i]);*/ - if(h_a[i]+h_b[i] !=h_c[i]) { - HIPCHECK(hipErrorUnknown); - } + CHECK (hipMalloc (&d_a, sizeof (int) * N)); + CHECK (hipMalloc (&d_b, sizeof (int) * N)); + CHECK (hipMalloc (&d_c, sizeof (int) * N)); + + CHECK (hipMemcpy (d_a, h_a, sizeof (int) * N, hipMemcpyHostToDevice)); + CHECK (hipMemcpy (d_b, h_b, sizeof (int) * N, hipMemcpyHostToDevice)); + + VectorAdd<<>> (d_a, d_b, d_c); + + CHECK (hipMemcpy (h_c, d_c, sizeof (int) * N, hipMemcpyDeviceToHost)); + + CHECK (hipFree (d_a)); + CHECK (hipFree (d_b)); + CHECK (hipFree (d_c)); + + bool error_found = false; + for (int i = 0; i < N; ++i) + { + if (h_a[i] + h_b[i] != h_c[i]) + { + fprintf (stderr, "%d + %d != %d (at index %d)", h_a[i], h_b[i], + h_c[i], i); + error_found = true; + } } - // Free up memory - HIPCHECK(hipStreamDestroy(stream[i])); - hipFree(d_a[i]); - hipFree(d_b[i]); - hipFree(d_c[i]); - } - return 0; + + return (error_found ? EXIT_FAILURE : EXIT_SUCCESS); } diff --git a/gdb/testsuite/gdb.rocm/nonstop-mode.exp b/gdb/testsuite/gdb.rocm/nonstop-mode.exp index dbb93640805..0f3d92b8810 100644 --- a/gdb/testsuite/gdb.rocm/nonstop-mode.exp +++ b/gdb/testsuite/gdb.rocm/nonstop-mode.exp @@ -72,20 +72,12 @@ if ![runto_main] { return -1 } -# Fetch the thread id of the first wave. +# Fetch the thread id of the first wave and the number of available GPUs. # Set breakpoint in device code. gdb_breakpoint "VectorAdd" "allow-pending" -gdb_test "continue" {.+hit\sBreakpoint\s\d+.+\sVectorAdd\s\(.*\)\sat.*} - -set threadid "" - -gdb_test_multiple "info threads" "extract thread id" { - -re "\\s+(\\d+)\\s+AMDGPU Wave\\s+\\d+\:\\d+:\\d+:\\d+\\s+\\(\\d+,\\d+,\\d+\\)/\\d+.*$gdb_prompt $" { - set threadid "$expect_out(1,string)" - pass $gdb_test_name - } -} +gdb_continue_to_breakpoint "continue to VectorAdd" +set threadid [get_valueof "" "\$_thread" 0] if {$threadid == ""} { return diff --git a/gdb/testsuite/gdb.rocm/precise-memory-warning-sigsegv.cpp b/gdb/testsuite/gdb.rocm/precise-memory-warning-sigsegv.cpp index 783f42dbe75..6259fac9e38 100644 --- a/gdb/testsuite/gdb.rocm/precise-memory-warning-sigsegv.cpp +++ b/gdb/testsuite/gdb.rocm/precise-memory-warning-sigsegv.cpp @@ -17,6 +17,17 @@ #include +#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 kernel () { int *p = nullptr; @@ -26,7 +37,7 @@ kernel () { int main (int argc, char* argv[]) { - hipLaunchKernelGGL (kernel, dim3 (1), dim3 (1), 0, 0); - hipDeviceSynchronize (); - return 0; + kernel<<<1, 1>>> (); + CHECK (hipDeviceSynchronize ()); + return 0; } diff --git a/gdb/testsuite/gdb.rocm/precise-memory-warning-watchpoint.cpp b/gdb/testsuite/gdb.rocm/precise-memory-warning-watchpoint.cpp index 37762aeb699..ea5f1b2b866 100644 --- a/gdb/testsuite/gdb.rocm/precise-memory-warning-watchpoint.cpp +++ b/gdb/testsuite/gdb.rocm/precise-memory-warning-watchpoint.cpp @@ -18,6 +18,17 @@ #include #include +#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); \ + } \ + } + __device__ int global = 0; __global__ void @@ -33,6 +44,6 @@ main (int argc, char* argv[]) { printf("host global: %p\n", &global); hipLaunchKernelGGL (kernel, dim3 (1), dim3 (1), 0, 0); - hipDeviceSynchronize (); + CHECK (hipDeviceSynchronize ()); return 0; } diff --git a/gdb/testsuite/gdb.rocm/precise-memory.cpp b/gdb/testsuite/gdb.rocm/precise-memory.cpp index 13222f50b82..e94f4a0bce0 100644 --- a/gdb/testsuite/gdb.rocm/precise-memory.cpp +++ b/gdb/testsuite/gdb.rocm/precise-memory.cpp @@ -17,6 +17,17 @@ #include +#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 kernel () { @@ -26,7 +37,7 @@ kernel () int main (int argc, char* argv[]) { - hipLaunchKernelGGL (kernel, dim3 (1), dim3 (1), 0, 0); - hipDeviceSynchronize (); + kernel<<<1, 1>>> (); + CHECK (hipDeviceSynchronize ()); return 0; } diff --git a/gdb/testsuite/gdb.rocm/snapshot-objfile-on-load.cpp b/gdb/testsuite/gdb.rocm/snapshot-objfile-on-load.cpp index 01e74b91781..afed41c0666 100644 --- a/gdb/testsuite/gdb.rocm/snapshot-objfile-on-load.cpp +++ b/gdb/testsuite/gdb.rocm/snapshot-objfile-on-load.cpp @@ -76,10 +76,10 @@ main (int argc, char* argv[]) /* Now that the module is submitted to the device, try to be the worst possible citizen by unloading the module and scrambling the underlying buffer. */ - hipModuleUnload (m); + CHECK (hipModuleUnload (m)); std::fill (module_buffer.begin (), module_buffer.end (), 0); module_buffer.resize (0); module_buffer.shrink_to_fit (); - hipDeviceSynchronize (); + CHECK (hipDeviceSynchronize ()); } diff --git a/gdb/testsuite/gdb.rocm/static-global.cpp b/gdb/testsuite/gdb.rocm/static-global.cpp index 52dc815b10f..967b7ec9afa 100644 --- a/gdb/testsuite/gdb.rocm/static-global.cpp +++ b/gdb/testsuite/gdb.rocm/static-global.cpp @@ -19,6 +19,17 @@ #include #include +#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); \ + } \ + } + __device__ int extern_global; __device__ static int static_global; @@ -58,6 +69,6 @@ main (int argc, char* argv[]) printf ("static_global's address on host: %p\n", &static_global); hipLaunchKernelGGL (kernel, dim3 (1), dim3 (1), 0, 0); - hipDeviceSynchronize (); + CHECK (hipDeviceSynchronize ()); return 0; } diff --git a/gdb/testsuite/lib/hip/hip-driver.cc b/gdb/testsuite/lib/hip/hip-driver.cc index e9ae5988d7c..b11688e92e9 100644 --- a/gdb/testsuite/lib/hip/hip-driver.cc +++ b/gdb/testsuite/lib/hip/hip-driver.cc @@ -160,7 +160,7 @@ main (int argc, char **argv, char **envp) CHECK (hipMemcpy (&exitcode_h, exitcode_d, sizeof (int), hipMemcpyDeviceToHost)); /* Wait until kernel finishes. */ - hipDeviceSynchronize (); + CHECK (hipDeviceSynchronize ()); return exitcode_h; } diff --git a/gdb/testsuite/lib/rocm.exp b/gdb/testsuite/lib/rocm.exp index 249d0747952..9b9128fca00 100644 --- a/gdb/testsuite/lib/rocm.exp +++ b/gdb/testsuite/lib/rocm.exp @@ -90,7 +90,8 @@ gdb_caching_proc allow_hipcc_tests { main () { kern<<<1, 1>>> (); - hipDeviceSynchronize (); + if (hipDeviceSynchronize () != hipSuccess) + return -1; return 0; } } executable $flags]} {