Catch2 Test Fixes (#2961)

* test fixes

* address PR comment

* PR comment fixing Nvidia pass

* add additional kernel launch checks

* pr comments

Author: pvelesko

tests/catch/ABM/AddKernels/add.cc

@@ -29,6 +29,7 @@ TEMPLATE_TEST_CASE("ABM_AddKernel_MultiTypeMultiSize", "", int, long, float, lon
   REQUIRE(res == hipSuccess);
 
   hipLaunchKernelGGL(add<TestType>, 1, size, 0, 0, d_a, d_b, d_c, size);
+  HIP_CHECK(hipGetLastError());
 
   res = hipMemcpy(a.data(), d_c, sizeof(TestType) * size, hipMemcpyDeviceToHost);
   REQUIRE(res == hipSuccess);

tests/catch/TypeQualifiers/hipManagedKeyword.cc

@@ -53,6 +53,7 @@ TEST_CASE("Unit_hipManagedKeyword_SingleGpu") {
   hipLaunchKernelGGL(add, dimGrid, dimBlock, 0, 0, static_cast<const float*>(A),
                      static_cast<float*>(B));
 
+  HIP_CHECK(hipGetLastError());
   HIP_CHECK(hipDeviceSynchronize());
 
   float maxError = 0.0f;

tests/catch/include/hip_test_checkers.hh

@@ -26,7 +26,7 @@ THE SOFTWARE.
 #include <fstream>
 #include <regex>
 #include <type_traits>
-
+#define TOL 0.001
 #define guarantee(cond, str)                                                                       \
   {                                                                                                \
     if (!(cond)) {                                                                                 \
@@ -45,7 +45,7 @@ size_t checkVectors(T* A, T* B, T* Out, size_t N, T (*F)(T a, T b), bool expectM
   size_t mismatchesToPrint = 10;
   for (size_t i = 0; i < N; i++) {
     T expected = F(A[i], B[i]);
-    if (Out[i] != expected) {
+    if (std::fabs(Out[i] - expected) > TOL) {
       if (mismatchCount == 0) {
         firstMismatch = i;
       }

tests/catch/include/hip_test_common.hh

@@ -270,6 +270,7 @@ void launchKernel(K kernel, Dim numBlocks, Dim numThreads, std::uint32_t memPerB
   launchRTCKernel<Typenames...>(kernel, numBlocks, numThreads, memPerBlock, stream,
                                 std::forward<Args>(packedArgs)...);
 #endif
+HIP_CHECK(hipGetLastError());
 }
 
 //---

tests/catch/multiproc/hipIpcEventHandle.cc

@@ -221,6 +221,7 @@ void runMultiProcKernel(ipcEventInfo_t *shmEventInfo, int index) {
     const dim3 blocks(BUF_SIZE / threads.x, 1);
     hipLaunchKernelGGL(computeKernel, dim3(blocks), dim3(threads), 0, 0,
                                     d_ptr + index *BUF_SIZE, d_ptr, index + 1);
+    HIP_CHECK(hipGetLastError());
     HIP_CHECK(hipEventRecord(event));
 
     // Barrier 2 : Signals that event is recorded

tests/catch/multiproc/hipMallocConcurrencyMproc.cc

@@ -120,7 +120,7 @@ static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) {
   hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
                      0, 0, static_cast<const int*>(A_d),
                      static_cast<const int*>(B_d), C_d, N);
-
+  HIP_CHECK(hipGetLastError());
   HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
 
   if (!HipTest::checkVectorADD(A_h, B_h, C_h, N)) {

tests/catch/stress/printf/Stress_printf_ComplexKernels.cc

@@ -276,12 +276,15 @@ bool test_printf_multistream(uint32_t num_blocks,
       hipLaunchKernelGGL(kernel_complex_opX, dim3(num_blocks, 1, 1),
                          dim3(threads_per_block, 1, 1),
                          0, stream[i], Ad, Bd, iterCount);
+      HIP_CHECK(hipGetLastError());
       hipLaunchKernelGGL(kernel_complex_opY, dim3(1, num_blocks, 1),
                          dim3(1, threads_per_block, 1),
                          0, stream[i], Ad, Bd, iterCount);
+      HIP_CHECK(hipGetLastError());
       hipLaunchKernelGGL(kernel_complex_opZ, dim3(1, 1, num_blocks),
                          dim3(1, 1, threads_per_block),
                          0, stream[i], Ad, Bd, iterCount);
+      HIP_CHECK(hipGetLastError());
     }
     HIP_CHECK(hipDeviceSynchronize());
     for (int i = 0; i < NUM_STREAM; i++) {
@@ -368,12 +371,15 @@ bool test_printf_multigpu(int gpu,
     hipLaunchKernelGGL(kernel_complex_opX, dim3(num_blocks, 1, 1),
                        dim3(threads_per_block, 1, 1),
                        0, 0, Ad, Bd, iterCount);
+    HIP_CHECK(hipGetLastError());
     hipLaunchKernelGGL(kernel_complex_opY, dim3(1, num_blocks, 1),
                        dim3(1, threads_per_block, 1),
                        0, 0, Ad, Bd, iterCount);
+    HIP_CHECK(hipGetLastError());
     hipLaunchKernelGGL(kernel_complex_opZ, dim3(1, 1, num_blocks),
                        dim3(1, 1, threads_per_block),
                        0, 0, Ad, Bd, iterCount);
+    HIP_CHECK(hipGetLastError());
     HIP_CHECK(hipDeviceSynchronize());
     std::ifstream CapturedData = captured.getCapturedData();
     char *buffer = new char[CHUNK_SIZE];

tests/catch/stress/printf/Stress_printf_SimpleKernels.cc

@@ -246,6 +246,7 @@ bool test_printf_conststr(uint32_t num_blocks, uint32_t threads_per_block,
     hipLaunchKernelGGL(kernel_printf_conststr, dim3(num_blocks, 1, 1),
                        dim3(threads_per_block, 1, 1),
                        0, 0, iterCount);
+    HIP_CHECK(hipGetLastError());
     HIP_CHECK(hipStreamSynchronize(0));
     std::ifstream CapturedData = captured.getCapturedData();
     char *buffer = new char[CHUNK_SIZE];
@@ -308,6 +309,7 @@ bool test_printf_two_conditionalstr(uint32_t num_blocks,
                        dim3(num_blocks, 1, 1),
                        dim3(threads_per_block, 1, 1),
                        0, 0, iterCount);
+    HIP_CHECK(hipGetLastError());
     HIP_CHECK(hipStreamSynchronize(0));
     std::ifstream CapturedData = captured.getCapturedData();
     char *buffer = new char[CHUNK_SIZE];
@@ -370,6 +372,7 @@ bool test_printf_single_conditionalstr(uint32_t num_blocks,
                        dim3(num_blocks, 1, 1),
                        dim3(threads_per_block, 1, 1),
                        0, 0, iterCount);
+    HIP_CHECK(hipGetLastError());
     HIP_CHECK(hipStreamSynchronize(0));
     std::ifstream CapturedData = captured.getCapturedData();
     char *buffer = new char[CHUNK_SIZE];
@@ -427,6 +430,7 @@ bool test_variable_str(uint32_t print_limit,
     hipLaunchKernelGGL(func, dim3(num_blocks, 1, 1),
                        dim3(threads_per_block, 1, 1),
                        0, 0, iterCount, Ad);
+    HIP_CHECK(hipGetLastError());
     HIP_CHECK(hipStreamSynchronize(0));
     HIP_CHECK(hipMemcpy(Ah, Ad, buffsize*sizeof(int32_t),
              hipMemcpyDeviceToHost));
@@ -483,6 +487,7 @@ bool test_decimal_str(uint32_t num_blocks, uint32_t threads_per_block,
     hipLaunchKernelGGL(kernel_decimal_calculation, dim3(num_blocks, 1, 1),
                        dim3(threads_per_block, 1, 1),
                        0, 0, iterCount, maxPrecision);
+    HIP_CHECK(hipGetLastError());
     HIP_CHECK(hipStreamSynchronize(0));
     std::ifstream CapturedData = captured.getCapturedData();
     char *buffer = new char[CHUNK_SIZE];

tests/catch/unit/cooperativeGrps/coalesced_groups_shfl_down.cc

@@ -157,6 +157,7 @@ static void test_group_partition(unsigned int tileSz) {
     // Launch Kernel
     hipLaunchKernelGGL(kernel_cg_group_partition, blockSize, threadsPerBlock,
                        threadsPerBlock * sizeof(int), 0, dResult, tileSz, i);
+    HIP_CHECK(hipGetLastError()); 
     err = hipDeviceSynchronize();
     if (err != hipSuccess) {
       fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
@@ -217,6 +218,7 @@ static void test_shfl_down() {
     // Launch Kernel
     hipLaunchKernelGGL(kernel_shfl_down, blockSize, threadsPerBlock,
                        threadsPerBlock * sizeof(int), 0, dPtr, dResults, lane_delta, i);
+    HIP_CHECK(hipGetLastError()); 
     hipMemcpy(hPtr, dResults, group_size_in_bytes, hipMemcpyDeviceToHost);
     err = hipDeviceSynchronize();
     if (err != hipSuccess) {

tests/catch/unit/cooperativeGrps/coalesced_groups_shfl_up.cc

@@ -138,6 +138,7 @@ static void test_group_partition(unsigned tileSz) {
     // Launch Kernel
     hipLaunchKernelGGL(kernel_cg_group_partition, blockSize, threadsPerBlock,
                      threadsPerBlock * sizeof(int), 0, dPtr, tileSz, i);
+    HIP_CHECK(hipGetLastError()); 
     hipMemcpy(hPtr, dPtr, arrSize, hipMemcpyDeviceToHost);
     err = hipDeviceSynchronize();
     if (err != hipSuccess) {
@@ -207,6 +208,7 @@ static void test_shfl_up() {
     hipLaunchKernelGGL(kernel_shfl_up, blockSize, threadsPerBlock,
                        threadsPerBlock * sizeof(int), 0, dPtr, dResults, lane_delta, i);
     hipMemcpy(hPtr, dResults, group_size_in_bytes, hipMemcpyDeviceToHost);
+    HIP_CHECK(hipGetLastError()); 
     err = hipDeviceSynchronize();
     if (err != hipSuccess) {
       fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));