SWDEV-379358 - [catch2][dtest] module tests migrated from direct to c…

…atch2

Change-Id: I1bac203c9f04642227616752393c255682d09f1b

catch/unit/CMakeLists.txt

@@ -35,6 +35,7 @@ add_subdirectory(compiler)
 add_subdirectory(errorHandling)
 add_subdirectory(cooperativeGrps)
 add_subdirectory(context)
+add_subdirectory(module)
 if(HIP_PLATFORM STREQUAL "amd")
 add_subdirectory(callback)
 #add_subdirectory(clock)

catch/unit/module/CMakeLists.txt

@@ -0,0 +1,43 @@
+# Common Tests - Test independent of all platforms
+set(TEST_SRC
+    hipModuleNegative.cc
+    hipModuleGetGlobal.cc
+    hipModuleOccupancyMaxPotentialActiveBlockSize.cc
+    hipModuleLoadMultiThreaded.cc
+    hipModuleLoadDataMultThreaded.cc
+    hipModuleLoadDataMultThreadOnMultGPU.cc
+    hipModuleLoadUnloadStress.cc
+    hipModuleTexture2dDrv.cc
+)
+
+add_custom_target(vcpy_kernel.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/vcpy_kernel.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../module/vcpy_kernel.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+
+add_custom_target(global_kernel.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/global_kernel.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../module/global_kernel.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+
+add_custom_target(empty_kernel.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/empty_kernel.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../module/empty_kernel.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+
+add_custom_target(kernel_composite_test.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/kernel_composite_test.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../module/kernel_composite_test.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+
+add_custom_target(tex2d_kernel.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/tex2d_kernel.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../module/tex2d_kernel.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+
+if (UNIX)
+add_custom_target(opencl_add.cc COMMAND cp ${CMAKE_CURRENT_SOURCE_DIR}/opencl_add.cc ${CMAKE_CURRENT_BINARY_DIR})
+endif()
+
+# only for AMD
+if(HIP_PLATFORM MATCHES "amd")
+  set(AMD_SRC
+      hipOpenCLCOTest.cc
+  )
+  set(TEST_SRC ${TEST_SRC} ${AMD_SRC})
+endif()
+
+hip_add_exe_to_target(NAME ModuleTest
+                      TEST_SRC ${TEST_SRC}
+                      TEST_TARGET_NAME build_tests)
+
+if (UNIX)
+add_dependencies(build_tests vcpy_kernel.code global_kernel.code empty_kernel.code kernel_composite_test.code tex2d_kernel.code opencl_add.cc)
+else()
+add_dependencies(build_tests vcpy_kernel.code global_kernel.code empty_kernel.code kernel_composite_test.code tex2d_kernel.code)
+endif()

catch/unit/module/empty_kernel.cc

@@ -0,0 +1,24 @@
+/*
+Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "hip/hip_runtime.h"
+
+extern "C" __global__ void EmptyKernel() {
+}
+

catch/unit/module/global_kernel.cc

@@ -0,0 +1,36 @@
+/*
+Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "hip/hip_runtime.h"
+
+#define ARRAY_SIZE (16)
+
+__device__ float myDeviceGlobal;
+__device__ float myDeviceGlobalArray[16];
+
+
+extern "C" __global__ void hello_world(const float* a, float* b) {
+    int tx = threadIdx.x;
+    b[tx] = a[tx];
+}
+
+extern "C" __global__ void test_globals(const float* a, float* b) {
+    int tx = threadIdx.x;
+    b[tx] = a[tx] + myDeviceGlobal + myDeviceGlobalArray[tx % ARRAY_SIZE];
+}

catch/unit/module/hipModuleGetGlobal.cc

@@ -0,0 +1,136 @@
+/*
+Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip_test_kernels.hh>
+#include <hip_test_common.hh>
+#include <hip_test_checkers.hh>
+#include <fstream>
+#include <vector>
+
+#define fileName "global_kernel.code"
+#define LEN 64
+#define SIZE LEN * sizeof(float)
+#define ARRAY_SIZE 16
+
+struct {
+  void* _Ad;
+  void* _Bd;
+} args;
+
+TEST_CASE("Unit_hipModuleGetGlobal") {
+  float *A, *B;
+  float *Ad, *Bd;
+  A = new float[LEN];
+  B = new float[LEN];
+
+  for (uint32_t i = 0; i < LEN; i++) {
+      A[i] = i * 1.0f;
+      B[i] = 0.0f;
+  }
+
+  HIP_CHECK(hipInit(0));
+  hipDevice_t device;
+  hipCtx_t context;
+  HIP_CHECK(hipDeviceGet(&device, 0));
+  #if HT_NVIDIA
+  HIP_CHECK(hipCtxCreate(&context, 0, device));
+  #endif
+  HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
+  HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), SIZE));
+
+  HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(Ad), A, SIZE));
+  HIP_CHECK(hipMemcpyHtoD((hipDeviceptr_t)(Bd), B, SIZE));
+  hipModule_t Module;
+  HIP_CHECK(hipModuleLoad(&Module, fileName));
+
+  float myDeviceGlobal_h = 42.0;
+  hipDeviceptr_t deviceGlobal;
+  size_t deviceGlobalSize;
+  HIP_CHECK(hipModuleGetGlobal(&deviceGlobal, &deviceGlobalSize,
+                               Module, "myDeviceGlobal"));
+  HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(deviceGlobal), &myDeviceGlobal_h,
+                                                        deviceGlobalSize));
+
+  float myDeviceGlobalArray_h[ARRAY_SIZE];
+  hipDeviceptr_t myDeviceGlobalArray;
+  size_t myDeviceGlobalArraySize;
+
+  HIP_CHECK(hipModuleGetGlobal(reinterpret_cast<hipDeviceptr_t*>(&myDeviceGlobalArray), //NOLINT
+           &myDeviceGlobalArraySize, Module, "myDeviceGlobalArray"));
+
+  for (int i = 0; i < ARRAY_SIZE; i++) {
+    myDeviceGlobalArray_h[i] = i * 1000.0f;
+    HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(myDeviceGlobalArray),
+               &myDeviceGlobalArray_h, myDeviceGlobalArraySize));
+  }
+
+  args._Ad = reinterpret_cast<void*>(Ad);
+  args._Bd = reinterpret_cast<void*>(Bd);
+
+  size_t size = sizeof(args);
+  void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args,
+                    HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, HIP_LAUNCH_PARAM_END};
+
+  SECTION("Running test for hello world kernel") {
+    hipFunction_t Function;
+    HIP_CHECK(hipModuleGetFunction(&Function, Module, "hello_world"));
+    HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL,
+                                    reinterpret_cast<void**>(&config)));
+
+    HIP_CHECK(hipMemcpyDtoH(B, hipDeviceptr_t(Bd), SIZE));
+
+    int mismatchCount = 0;
+    for (uint32_t i = 0; i < LEN; i++) {
+      if (A[i] != B[i]) {
+        mismatchCount++;
+      if (mismatchCount >= 10) {
+        break;
+      }
+    }
+    }
+  REQUIRE(mismatchCount == 0);
+  }
+
+  SECTION("running test for tests_globals kernel") {
+    hipFunction_t Function;
+    HIP_CHECK(hipModuleGetFunction(&Function, Module,
+                                   "test_globals"));
+    HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL,
+                                    reinterpret_cast<void**>(&config)));
+    HIP_CHECK(hipMemcpyDtoH(B, hipDeviceptr_t(Bd), SIZE));
+
+    int mismatchCount = 0;
+    for (uint32_t i = 0; i < LEN; i++) {
+      float expected;
+      expected = A[i] + myDeviceGlobal_h + + myDeviceGlobalArray_h[i % 16];
+      if (expected != B[i]) {
+        mismatchCount++;
+      if (mismatchCount >= 10) {
+        break;
+      }
+    }
+    }
+  REQUIRE(mismatchCount == 0);
+  }
+
+  HIP_CHECK(hipModuleUnload(Module));
+  #if HT_NVIDIA
+  HIP_CHECK(hipCtxDestroy(context));
+  #endif
+}

catch/unit/module/hipModuleLoadDataMultThreadOnMultGPU.cc

@@ -0,0 +1,142 @@
+/*
+Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip_test_kernels.hh>
+#include <hip_test_common.hh>
+#include <hip_test_checkers.hh>
+#include <hip_test_helper.hh>
+#include <fstream>
+#include <vector>
+
+#define LEN 64
+#define SIZE LEN << 2
+#define THREADS 8
+
+#define FILENAME "vcpy_kernel.code"
+#define kernel_name "hello_world"
+
+static std::vector<char> load_file() {
+  std::ifstream file(FILENAME, std::ios::binary | std::ios::ate);
+  std::streamsize fsize = file.tellg();
+  file.seekg(0, std::ios::beg);
+
+  std::vector<char> buffer(fsize);
+  if (!file.read(buffer.data(), fsize)) {
+    printf("Info:could not open code object '%s'\n", FILENAME);
+  }
+  return buffer;
+}
+
+static void run(const std::vector<char>& buffer, int device) {
+  HIP_CHECK(hipSetDevice(device));
+  hipModule_t Module;
+  hipFunction_t Function;
+
+  float *A, *B, *Ad, *Bd;
+  A = new float[LEN];
+  B = new float[LEN];
+
+  for (uint32_t i = 0; i < LEN; i++) {
+    A[i] = i * 1.0f;
+    B[i] = 0.0f;
+  }
+
+  HIP_CHECK(hipMalloc(&Ad, SIZE));
+  HIP_CHECK(hipMalloc(&Bd, SIZE));
+
+  HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
+  HIP_CHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
+
+  HIP_CHECK(hipModuleLoadData(&Module, &buffer[0]));
+  HIP_CHECK(hipModuleGetFunction(&Function, Module, kernel_name));
+
+  hipStream_t stream;
+  HIP_CHECK(hipStreamCreate(&stream));
+
+  struct {
+    void* _Ad;
+    void* _Bd;
+  } args;
+  args._Ad = static_cast<void*>(Ad);
+  args._Bd = static_cast<void*>(Bd);
+  size_t size = sizeof(args);
+
+  void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args,
+                    HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, HIP_LAUNCH_PARAM_END};
+  HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, stream,
+            NULL, reinterpret_cast<void**>(&config)));
+
+  HIP_CHECK(hipStreamSynchronize(stream));
+  HIP_CHECK(hipStreamDestroy(stream));
+  HIP_CHECK(hipModuleUnload(Module));
+  HIP_CHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost));
+
+  for (uint32_t i = 0; i < LEN; i++) {
+    REQUIRE(A[i] == B[i]);
+  }
+  HIP_CHECK(hipFree(Ad));
+  HIP_CHECK(hipFree(Bd));
+  delete[] A;
+  delete[] B;
+}
+
+struct joinable_thread : std::thread {
+  template <class... Xs>
+  joinable_thread(Xs&&... xs) : std::thread(std::forward<Xs>(xs)...) {} // NOLINT
+
+  joinable_thread& operator=(joinable_thread&& other) = default;
+  joinable_thread(joinable_thread&& other)      = default;
+
+  ~joinable_thread() {
+    if (this->joinable())
+      this->join();
+  }
+};
+
+static void run_multi_threads(uint32_t n, const std::vector<char>& buffer) {
+  int numDevices = 0;
+  HIP_CHECK(hipGetDeviceCount(&numDevices));
+  if (numDevices < 2) {
+    HipTest::HIP_SKIP_TEST("Skipping because devices < 2");
+    return;
+  }
+  std::vector<joinable_thread> threads;
+
+  for (int device =0; device < numDevices; ++device) {
+    for (uint32_t i = 0; i < n; i++) {
+      threads.emplace_back(std::thread{[&, device] {
+        run(buffer, device);
+      }
+    });
+    }
+  }
+}
+
+TEST_CASE("Unit_hipModuleLoadDataMultGPUOnMultThread") {
+  HIP_CHECK(hipInit(0));
+  auto buffer = load_file();
+  auto file_size = buffer.size() / (1024 * 1024);
+  auto thread_count = HipTest::getHostThreadCount(file_size + 10, THREADS);
+  if (thread_count == 0) {
+    HipTest::HIP_SKIP_TEST("Skipping because thread_count is 0");
+    return;
+  }
+  // run multi thread on multi devices
+  run_multi_threads(thread_count, buffer);
+}