SWDEV-460144 - Adds test with inline static variable used in different translation unit

Change-Id: I9d5c8109f7264cdf469361b983b7a940706379c2

catch/unit/memory/CMakeLists.txt

@@ -215,6 +215,20 @@ hip_add_exe_to_target(NAME MemoryTest2
   TEST_SRC ${TEST_SRC}
   TEST_TARGET_NAME build_tests COMMON_SHARED_SRC ${COMMON_SHARED_SRC})
 
+if(HIP_PLATFORM MATCHES "amd")
+set(TEST_SRC
+  inlineVar.cc
+  memoryCommon.cc
+)
+
+hip_add_exe_to_target(NAME InlineVarTest
+  TEST_SRC ${TEST_SRC}
+  TEST_TARGET_NAME build_tests COMMON_SHARED_SRC ${COMMON_SHARED_SRC})
+
+  set_target_properties(InlineVarTest PROPERTIES COMPILE_FLAGS -fgpu-rdc)
+  set_target_properties(InlineVarTest PROPERTIES LINK_FLAGS -fgpu-rdc)
+endif()
+
 set(TEST_SRC
   hipSVMTestByteGranularity.cpp
   hipSVMTestFineGrainMemoryConsistency.cpp

catch/unit/memory/inlineVar.cc

@@ -0,0 +1,33 @@
+/*
+Copyright (c) 2024 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.
+*/
+
+// This test verifies the usage of inline constant variable in different translation units
+// The inline variable is declared in MemUtils.hh, it is used in memoryCommon.cc where
+// set_value and get_value are defined.
+
+#include "MemUtils.hh"
+#include "memoryCommon.hh"
+using namespace mem_utils;
+
+TEST_CASE("Unit_hipMemcpyToFromSymbol_GlobalConstVar") {
+  int const initialValue = 10;
+  set_value(initialValue);
+  int const finalValue = get_value();
+  REQUIRE(finalValue == initialValue);
+}

catch/unit/memory/memoryCommon.cc

@@ -0,0 +1,12 @@
+#include "memoryCommon.hh"
+#include "memoryGlobal.hh"
+
+void set_value(int const value) {
+  HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(globalVar), &value, sizeof(value)));
+}
+
+int get_value() {
+  int value;
+  HIP_CHECK(hipMemcpyFromSymbol(&value, HIP_SYMBOL(globalVar), sizeof(value)));
+  return value;
+}

catch/unit/memory/memoryCommon.hh

@@ -0,0 +1,24 @@
+/*
+Copyright (c) 2024 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
+
+void set_value(int const value);
+
+int get_value();

catch/unit/memory/memoryGlobal.hh

@@ -0,0 +1,5 @@
+#pragma once
+#include <memory>
+#include <hip_test_common.hh>
+
+inline __constant__ int globalVar;