SWDEV-355313 - Update latest code

Change-Id: I9c03cde09b42c8e3726153c2a177359efc8d6d29

.gitignore

@@ -0,0 +1,17 @@
+.*
+!.gitignore
+*.o
+*.exe
+*.swp
+lib
+packages
+build
+tags
+samples/0_Intro/module_api/runKernel.hip.out
+samples/0_Intro/module_api/vcpy_isa.code
+samples/0_Intro/module_api/vcpy_isa.hsaco
+samples/0_Intro/module_api/vcpy_kernel.co
+samples/0_Intro/module_api/vcpy_kernel.code
+samples/1_Utils/hipInfo/hipInfo
+samples/1_Utils/hipBusBandwidth/hipBusBandwidth
+samples/1_Utils/hipDispatchLatency/hipDispatchLatency

catch/hipTestMain/config/config_amd_linux_common.json

@@ -12,6 +12,8 @@
 	   "Unit_hipIpcOpenMemHandle_Negative_Open_In_Creating_Process",
 	   "Unit_hipDeviceGetPCIBusId_Negative_PartialFill",
 	   "Unit_hipInit_Negative",
-	   "Unit_hipMemset_Negative_OutOfBoundsPtr"
+	   "Unit_hipMemset_Negative_OutOfBoundsPtr",
+	   "Unit_hipDeviceReset_Positive_Basic",
+	   "Unit_hipDeviceReset_Positive_Threaded"
 	]
 }

catch/hipTestMain/config/config_amd_windows_MI2xx.json

@@ -90,6 +90,8 @@
 		"Unit_hipStreamValue_Wait64_Blocking_NoMask_And",
 		"Unit_hipStreamValue_Wait64_Blocking_NoMask_Nor",
 		"Unit_hipDeviceGetPCIBusId_Negative_PartialFill",
-		"Unit_hipInit_Negative"
+		"Unit_hipInit_Negative",
+		"Unit_hipGraphAddEventRecordNode_Functional_ElapsedTime",
+		"Unit_hipStreamBeginCapture_captureComplexGraph"
 	]
 }

catch/hipTestMain/config/config_amd_windows_common.json

@@ -80,15 +80,19 @@
 	   "Unit_hipDeviceGetPCIBusId_Negative_PartialFill",
 	   "Unit_hipDeviceGetSharedMemConfig_Positive_Basic",
 	   "Unit_hipDeviceGetSharedMemConfig_Positive_Threaded",
+	   "Unit_hipDeviceReset_Positive_Basic",
+	   "Unit_hipDeviceReset_Positive_Threaded",
 	   "Unit_hipInit_Negative",
 	   "Unit_hipGraphMemcpyNodeSetParams_Functional",
 	   "Unit_hipGraphNodeGetDependentNodes_Functional",
 	   "Unit_hipGraphNodeGetDependencies_Functional",
 	   "Unit_hipGraphExecChildGraphNodeSetParams_ChildTopology",
+	   "Unit_hipGraphAddEventRecordNode_MultipleRun",
+	   "Unit_hipGraphAddEventRecordNode_Functional_ElapsedTime",
+	   "Unit_hipStreamBeginCapture_captureComplexGraph",
 	   "Note: needs to be enabled when streamPerThread issues are fixed",
 	   "Unit_hipStreamSynchronize_NullStreamAndStreamPerThread",
 	   "Note: intermittent Seg fault failure ",
-	   "Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags",
-	   "Unit_hipGraphAddEventRecordNode_MultipleRun"
+	   "Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags"
 	]
 }

catch/hipTestMain/hip_test_context.cc

@@ -4,6 +4,7 @@
 #include <fstream>
 #include <sstream>
 #include <regex>
+#include "hip_test_context.hh"
 #include "hip_test_filesystem.hh"
 
 void TestContext::detectOS() {
@@ -98,11 +99,11 @@ void TestContext::getConfigFiles() {
     abort();
   }
 
-  const char* env_config = std::getenv("HIP_CATCH_EXCLUDE_FILE");
+  std::string env_config = TestContext::getEnvVar("HIP_CATCH_EXCLUDE_FILE");
   LogPrintf("Env Config file: %s",
-            (env_config != nullptr) ? env_config : "Not found, using common config");
+            (!env_config.empty()) ? env_config.c_str() : "Not found, using common config");
   // HIP_CATCH_EXCLUDE_FILE is set for custom file path
-  if (env_config != nullptr) {
+  if (!env_config.empty()) {
     if(fs::exists(env_config)) {
       config_.json_files.push_back(env_config);
     }

catch/include/hip_array_common.hh

@@ -0,0 +1,84 @@
+/*
+Copyright (c) 2022 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.
+*/
+
+#pragma once
+
+#include <hip_test_common.hh>
+
+template <class T, size_t N, hipArray_Format Format> struct type_and_size_and_format {
+  using type = T;
+  static constexpr size_t size = N;
+  static constexpr hipArray_Format format = Format;
+};
+
+// Create a map of type to scalar type, vector size and scalar type format enum.
+// This is useful for creating simpler function that depend on the vector size.
+template <typename T> struct vector_info;
+template <>
+struct vector_info<int> : type_and_size_and_format<int, 1, HIP_AD_FORMAT_SIGNED_INT32> {};
+template <> struct vector_info<float> : type_and_size_and_format<float, 1, HIP_AD_FORMAT_FLOAT> {};
+template <>
+struct vector_info<short> : type_and_size_and_format<short, 1, HIP_AD_FORMAT_SIGNED_INT16> {};
+template <>
+struct vector_info<char> : type_and_size_and_format<char, 1, HIP_AD_FORMAT_SIGNED_INT8> {};
+template <>
+struct vector_info<unsigned int>
+    : type_and_size_and_format<unsigned int, 1, HIP_AD_FORMAT_UNSIGNED_INT32> {};
+template <>
+struct vector_info<unsigned short>
+    : type_and_size_and_format<unsigned short, 1, HIP_AD_FORMAT_UNSIGNED_INT16> {};
+template <>
+struct vector_info<unsigned char>
+    : type_and_size_and_format<unsigned char, 1, HIP_AD_FORMAT_UNSIGNED_INT8> {};
+
+template <>
+struct vector_info<int2> : type_and_size_and_format<int, 2, HIP_AD_FORMAT_SIGNED_INT32> {};
+template <> struct vector_info<float2> : type_and_size_and_format<float, 2, HIP_AD_FORMAT_FLOAT> {};
+template <>
+struct vector_info<short2> : type_and_size_and_format<short, 2, HIP_AD_FORMAT_SIGNED_INT16> {};
+template <>
+struct vector_info<char2> : type_and_size_and_format<char, 2, HIP_AD_FORMAT_SIGNED_INT8> {};
+template <>
+struct vector_info<uint2>
+    : type_and_size_and_format<unsigned int, 2, HIP_AD_FORMAT_UNSIGNED_INT32> {};
+template <>
+struct vector_info<ushort2>
+    : type_and_size_and_format<unsigned short, 2, HIP_AD_FORMAT_UNSIGNED_INT16> {};
+template <>
+struct vector_info<uchar2>
+    : type_and_size_and_format<unsigned char, 2, HIP_AD_FORMAT_UNSIGNED_INT8> {};
+
+template <>
+struct vector_info<int4> : type_and_size_and_format<int, 4, HIP_AD_FORMAT_SIGNED_INT32> {};
+template <> struct vector_info<float4> : type_and_size_and_format<float, 4, HIP_AD_FORMAT_FLOAT> {};
+template <>
+struct vector_info<short4> : type_and_size_and_format<short, 4, HIP_AD_FORMAT_SIGNED_INT16> {};
+template <>
+struct vector_info<char4> : type_and_size_and_format<char, 4, HIP_AD_FORMAT_SIGNED_INT8> {};
+template <>
+struct vector_info<uint4>
+    : type_and_size_and_format<unsigned int, 4, HIP_AD_FORMAT_UNSIGNED_INT32> {};
+template <>
+struct vector_info<ushort4>
+    : type_and_size_and_format<unsigned short, 4, HIP_AD_FORMAT_UNSIGNED_INT16> {};
+template <>
+struct vector_info<uchar4>
+    : type_and_size_and_format<unsigned char, 4, HIP_AD_FORMAT_UNSIGNED_INT8> {};

catch/include/hip_test_context.hh

@@ -32,6 +32,7 @@ THE SOFTWARE.
 #include <set>
 #include <unordered_map>
 
+// OS Check
 #if defined(_WIN32)
 #define HT_WIN 1
 #define HT_LINUX 0
@@ -42,6 +43,7 @@ THE SOFTWARE.
 #error "OS not recognized"
 #endif
 
+// Platform check
 #if defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)
 #define HT_AMD 1
 #define HT_NVIDIA 0
@@ -52,16 +54,6 @@ THE SOFTWARE.
 #error "Platform not recognized"
 #endif
 
-static int _log_enable = (std::getenv("HT_LOG_ENABLE") ? 1 : 0);
-
-#define LogPrintf(format, ...)                                                                     \
-  {                                                                                                \
-    if (_log_enable) {                                                                             \
-      printf(format, __VA_ARGS__);                                                                 \
-      printf("%c", '\n');                                                                          \
-    }                                                                                              \
-  }
-
 typedef struct Config_ {
   std::vector<std::string> json_files;  // Json files
   std::string platform;   // amd/nvidia
@@ -124,6 +116,26 @@ class TestContext {
     return instance;
   }
 
+  static std::string getEnvVar(std::string var) {
+    #if defined(_WIN32)
+    rsize_t MAX_LEN = 4096;
+    char dstBuf[MAX_LEN];
+    size_t dstSize;
+    if (!::getenv_s(&dstSize, dstBuf, MAX_LEN, var.c_str())) {
+      return std::string(dstBuf);
+    }
+    #elif defined(__linux__)
+    char* val = std::getenv(var.c_str());
+    if (val != NULL) {
+      return std::string(val);
+    }
+    #else
+    #error "OS not recognized"
+    #endif
+    return std::string("");
+  }
+
+
   bool isWindows() const;
   bool isLinux() const;
   bool isNvidia() const;
@@ -171,3 +183,14 @@ class TestContext {
 
   ~TestContext();
 };
+
+static bool _log_enable = (!TestContext::getEnvVar("HT_LOG_ENABLE").empty() ? true : false);
+
+// printing logs
+#define LogPrintf(format, ...)                                                                   \
+{                                                                                                \
+  if(_log_enable) {                                                                              \
+    printf(format, __VA_ARGS__);                                                                 \
+    printf("%c", '\n');                                                                          \
+  }                                                                                              \
+}

catch/include/resource_guards.hh

@@ -19,6 +19,7 @@ THE SOFTWARE.
 
 #pragma once
 
+#include <hip_array_common.hh>
 #include <hip_test_common.hh>
 #include <hip/hip_runtime_api.h>
 
@@ -80,17 +81,121 @@ template <typename T> class LinearAllocGuard {
     }
   }
 
-  T* ptr() { return ptr_; };
-  T* const ptr() const { return ptr_; };
-  T* host_ptr() { return host_ptr_; }
-  T* const host_ptr() const { return host_ptr(); }
+  T* ptr() const { return ptr_; };
+  T* host_ptr() const { return host_ptr_; }
 
  private:
   const LinearAllocs allocation_type_;
   T* ptr_ = nullptr;
   T* host_ptr_ = nullptr;
 };
 
+template <typename T> class LinearAllocGuardMultiDim {
+ protected:
+  LinearAllocGuardMultiDim(hipExtent extent) : extent_{extent} {}
+
+  ~LinearAllocGuardMultiDim() { static_cast<void>(hipFree(pitched_ptr_.ptr)); }
+
+ public:
+  T* ptr() const { return reinterpret_cast<T*>(pitched_ptr_.ptr); };
+
+  size_t pitch() const { return pitched_ptr_.pitch; }
+
+  hipExtent extent() const { return extent_; }
+
+  hipPitchedPtr pitched_ptr() const { return pitched_ptr_; }
+
+  size_t width() const { return extent_.width; }
+
+  size_t width_logical() const { return extent_.width / sizeof(T); }
+
+  size_t height() const { return extent_.height; }
+
+ public:
+  hipPitchedPtr pitched_ptr_;
+  const hipExtent extent_;
+};
+
+template <typename T> class LinearAllocGuard2D : public LinearAllocGuardMultiDim<T> {
+ public:
+  LinearAllocGuard2D(const size_t width_logical, const size_t height)
+      : LinearAllocGuardMultiDim<T>{make_hipExtent(width_logical * sizeof(T), height, 1)} {
+    HIP_CHECK(hipMallocPitch(&this->pitched_ptr_.ptr, &this->pitched_ptr_.pitch,
+                             this->extent_.width, this->extent_.height));
+  }
+
+  LinearAllocGuard2D(const LinearAllocGuard2D&) = delete;
+  LinearAllocGuard2D(LinearAllocGuard2D&&) = delete;
+};
+
+template <typename T> class LinearAllocGuard3D : public LinearAllocGuardMultiDim<T> {
+ public:
+  LinearAllocGuard3D(const size_t width_logical, const size_t height, const size_t depth)
+      : LinearAllocGuardMultiDim<T>{make_hipExtent(width_logical * sizeof(T), height, depth)} {
+    HIP_CHECK(hipMalloc3D(&this->pitched_ptr_, this->extent_));
+  }
+
+  LinearAllocGuard3D(const hipExtent extent) : LinearAllocGuardMultiDim<T>(extent) {
+    HIP_CHECK(hipMalloc3D(&this->pitched_ptr_, this->extent_));
+  }
+
+  LinearAllocGuard3D(const LinearAllocGuard3D&) = delete;
+  LinearAllocGuard3D(LinearAllocGuard3D&&) = delete;
+
+  size_t depth() const { return this->extent_.depth; }
+};
+
+template <typename T> class ArrayAllocGuard {
+ public:
+  // extent should contain logical width
+  ArrayAllocGuard(const hipExtent extent, const unsigned int flags = 0u) : extent_{extent} {
+    hipChannelFormatDesc desc = hipCreateChannelDesc<T>();
+    HIP_CHECK(hipMalloc3DArray(&ptr_, &desc, extent_, flags));
+  }
+
+  ~ArrayAllocGuard() { static_cast<void>(hipFreeArray(ptr_)); }
+
+  ArrayAllocGuard(const ArrayAllocGuard&) = delete;
+  ArrayAllocGuard(ArrayAllocGuard&&) = delete;
+
+  hipArray_t ptr() const { return ptr_; }
+
+  hipExtent extent() const { return extent_; }
+
+ private:
+  hipArray_t ptr_ = nullptr;
+  const hipExtent extent_;
+};
+
+template <typename T> class DrvArrayAllocGuard {
+ public:
+  // extent should contain width in bytes
+  DrvArrayAllocGuard(const hipExtent extent, const unsigned int flags = 0u) : extent_{extent} {
+    HIP_ARRAY3D_DESCRIPTOR desc{};
+    using vec_info = vector_info<T>;
+    desc.Format = vec_info::format;
+    desc.NumChannels = vec_info::size;
+    desc.Width = extent_.width / sizeof(T);
+    desc.Height = extent_.height;
+    desc.Depth = extent_.depth;
+    desc.Flags = flags;
+    HIP_CHECK(hipArray3DCreate(&ptr_, &desc));
+  }
+
+  ~DrvArrayAllocGuard() { static_cast<void>(hipArrayDestroy(ptr_)); }
+
+  DrvArrayAllocGuard(const DrvArrayAllocGuard&) = delete;
+  DrvArrayAllocGuard(DrvArrayAllocGuard&&) = delete;
+
+  hiparray ptr() const { return ptr_; }
+
+  hipExtent extent() const { return extent_; }
+
+ private:
+  hiparray ptr_ = nullptr;
+  const hipExtent extent_;
+};
+
 enum class Streams { nullstream, perThread, created };
 
 class StreamGuard {