SWDEV-355608 - deprecate/cleanup hipcc link flags (#3128)

- deprecate -use-staticlib, -use-sharedlib which no longer provide any functional values - use --hip-link instead of specifying the HIP runtime by name when linking - fix linker option bug in HIT test's cmake - update build options for unit tests requiring pthread or rt Change-Id: Ib49978773c80fb40c71dc52b050ce921943ee3e4 [ROCm/hip commit: a388eb1e6f]
2023-01-24 06:42:52 +00:00
@@ -199,10 +199,6 @@ if ($HIP_PLATFORM eq "amd") {
        print ("HIP_CLANG_RT_LIB=$HIP_CLANG_RT_LIB\n");
    }

-    $HIPLDFLAGS .= " -L\"$HIP_LIB_PATH\"";
-    if ($isWindows) {
-      $HIPLDFLAGS .= " -lamdhip64";
-    }
    if ($HIP_CLANG_HCC_COMPAT_MODE) {
        ## Allow __fp16 as function parameter and return type.
        $HIPCXXFLAGS .= " -Xclang -fallow-half-arguments-and-returns -D__HIP_HCC_COMPAT_MODE__=1";
@@ -245,8 +241,6 @@ my $printCXXFlags = 0;      # print HIPCXXFLAGS
 my $printLDFlags = 0;       # print HIPLDFLAGS
 my $runCmd = 1;
 my $buildDeps = 0;
-my $linkType = 1;
-my $setLinkType = 0;
 my $hsacoVersion = 0;
 my $funcSupp = 0;      # enable function support
 my $rdc = 0;           # whether -fgpu-rdc is on
@@ -361,16 +355,11 @@ foreach $arg (@ARGV)
        $compileOnly = 1;
        $buildDeps = 1;
    }
-    if(($trimarg eq '-use-staticlib') and ($setLinkType eq 0))
-    {
-        $linkType = 0;
-        $setLinkType = 1;
-        $swallowArg = 1;
+    if($trimarg eq '-use-staticlib') {
+        print "Warning: The -use-staticlib option has been deprecated and is no longer needed.\n"
    }
-    if(($trimarg eq '-use-sharedlib') and ($setLinkType eq 0))
-    {
-        $linkType = 1;
-        $setLinkType = 1;
+    if($trimarg eq '-use-sharedlib') {
+        print "Warning: The -use-sharedlib option has been deprecated and is no longer needed.\n"
    }
    if($arg =~ m/^-O/)
    {
@@ -558,12 +547,6 @@ if ($buildDeps and $HIP_PLATFORM eq 'amd') {
    $HIPCXXFLAGS .= " --cuda-host-only";
 }

-# Add --hip-link only if it is compile only and -fgpu-rdc is on.
-if ($rdc and !$compileOnly and $HIP_PLATFORM eq 'amd') {
-    $HIPLDFLAGS .= " --hip-link";
-    $HIPLDFLAGS .= $HIPLDARCHFLAGS;
-}
-
 # hipcc currrently requires separate compilation of source files, ie it is not possible to pass
 # CPP files combined with .O files
 # Reason is that NVCC uses the file extension to determine whether to compile in CUDA mode or
@@ -588,18 +571,16 @@ if ($HIP_PLATFORM eq "amd") {
            $HIPCXXFLAGS .= " --hip-device-lib-path=\"$DEVICE_LIB_PATH\"";
        }
    }
-    if (not $isWindows) {
-        $HIPLDFLAGS .= " -lgcc_s -lgcc -lpthread -lm -lrt";
-    }

-    if (not $isWindows  and not $compileOnly) {
-      if ($linkType eq 0) {
-        $toolArgs = " -L$HIP_LIB_PATH -lamdhip64 -L$ROCM_PATH/lib -lhsa-runtime64 -ldl -lnuma " . ${toolArgs};
-      } else {
-        $toolArgs = ${toolArgs} . " -Wl,-rpath=$HIP_LIB_PATH:$ROCM_PATH/lib -lamdhip64 ";
-      }
+    if (!$compileOnly) {
+        $HIPLDFLAGS .= " --hip-link";
+        if ($rdc) {
+            $HIPLDFLAGS .= $HIPLDARCHFLAGS;
+        }
+        if (not $isWindows) {
+            $HIPLDFLAGS .= " --rtlib=compiler-rt -unwindlib=libgcc";

-      $toolArgs .= " -L$HIP_CLANG_RT_LIB -lclang_rt.builtins-x86_64 "
+        }
    }
 }

@@ -21,7 +21,7 @@ $ cd ~/hip/samples/0_Intro/square
 $ make
 /opt/rocm/hip/bin/hipify-perl square.cu > square.cpp
 /opt/rocm/hip/bin/hipcc  square.cpp -o square.out
-/opt/rocm/hip/bin/hipcc -use-staticlib  square.cpp -o square.out.static
+/opt/rocm/hip/bin/hipcc  square.cpp -o square.out
 ```
 - Execute file
 ```
@@ -19,6 +19,9 @@ set(LINUX_TEST_SRC
 )

 add_custom_target(dummy_kernel.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${CMAKE_CURRENT_SOURCE_DIR}/dummy_kernel.cpp -o ${CMAKE_CURRENT_BINARY_DIR}/../multiproc/dummy_kernel.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+if (UNIX)
+  set(__LINKER_LIBS__ pthread)
+endif()

 # the last argument linker libraries is required for this test but optional to the function
 if(HIP_PLATFORM MATCHES "nvidia")
@@ -30,7 +33,6 @@ elseif(HIP_PLATFORM MATCHES "amd")
 hip_add_exe_to_target(NAME MultiProc
                      TEST_SRC ${LINUX_TEST_SRC}
                      TEST_TARGET_NAME build_tests
-                      LINKER_LIBS ${CMAKE_DL_LIBS})
+                      LINKER_LIBS ${CMAKE_DL_LIBS} ${__LINKER_LIBS__})
 endif()
 add_dependencies(build_tests dummy_kernel.code)
-
@@ -43,10 +43,15 @@ set_source_files_properties(hipDeviceGetP2PAttribute.cc PROPERTIES COMPILE_FLAGS
 add_executable(getDeviceCount EXCLUDE_FROM_ALL getDeviceCount_exe.cc)
 add_executable(hipDeviceGetP2PAttribute EXCLUDE_FROM_ALL hipDeviceGetP2PAttribute_exe.cc)

+if (UNIX)
+  set(__LINKER_LIBS__ pthread)
+endif()
+
 hip_add_exe_to_target(NAME DeviceTest
                      TEST_SRC ${TEST_SRC}
                      TEST_TARGET_NAME build_tests
-                      COMPILE_OPTIONS -std=c++14)
+                      COMPILE_OPTIONS -std=c++14
+                      LINKER_LIBS ${__LINKER_LIBS__})

 add_dependencies(DeviceTest getDeviceCount)
 add_dependencies(DeviceTest hipDeviceGetP2PAttribute)
@@ -18,6 +18,7 @@ set(TEST_SRC
 if(UNIX)
    set(TEST_SRC ${TEST_SRC}
        deviceAllocation.cc)
+    set(__LINKER_LIBS__ pthread)
 endif()

 # AMD only tests
@@ -86,7 +87,7 @@ endif()
    hip_add_exe_to_target(NAME UnitDeviceTests
                      TEST_SRC ${TEST_SRC}
                      TEST_TARGET_NAME build_tests
-                      LINKER_LIBS hiprtc)
+                      LINKER_LIBS hiprtc ${__LINKER_LIBS__})
 elseif(HIP_PLATFORM MATCHES "nvidia")
    hip_add_exe_to_target(NAME UnitDeviceTests
                      TEST_SRC ${TEST_SRC}
@@ -6,7 +6,12 @@ set(TEST_SRC
    hipPeekAtLastError.cc
 )

+if (UNIX)
+  set(__LINKER_LIBS__ pthread)
+endif()
+
 hip_add_exe_to_target(NAME ErrorHandlingTest
                      TEST_SRC ${TEST_SRC}
                      TEST_TARGET_NAME build_tests
-                      COMPILE_OPTIONS -std=c++14)
+		      LINKER_LIBS ${__LINKER_LIBS__}
+		      COMPILE_OPTIONS -std=c++14)
@@ -20,6 +20,10 @@ if(HIP_PLATFORM MATCHES "amd")
  set(TEST_SRC ${TEST_SRC} ${AMD_SRC})
 endif()

+if (UNIX)
+  set(__LINKER_LIBS__ pthread)
+endif()
+
 hip_add_exe_to_target(NAME EventTest
                      TEST_SRC ${TEST_SRC}
                      TEST_TARGET_NAME build_tests)
@@ -83,6 +83,11 @@ set(TEST_SRC
  hipUserObjectCreate.cc
 )

+if (UNIX)
+  set(__LINKER_LIBS__ pthread)
+endif()
+
 hip_add_exe_to_target(NAME GraphsTest
                      TEST_SRC ${TEST_SRC}
-                      TEST_TARGET_NAME build_tests)
+                      TEST_TARGET_NAME build_tests
+                      LINKER_LIBS ${__LINKER_LIBS__})
@@ -182,8 +182,11 @@ if(UNIX)
    set(TEST_SRC ${TEST_SRC}
                 hipHmmOvrSubscriptionTst.cc
                 hipMemoryAllocateCoherent.cc)
+                 
+    set(__LINKER_LIBS__ pthread)
 endif()

 hip_add_exe_to_target(NAME MemoryTest
                      TEST_SRC ${TEST_SRC}
-                      TEST_TARGET_NAME build_tests)
+                      TEST_TARGET_NAME build_tests
+                      LINKER_LIBS ${__LINKER_LIBS__})
@@ -5,6 +5,11 @@ set(TEST_SRC
  hipMultiThreadStreams2.cc
 )

+if (UNIX)
+  set(__LINKER_LIBS__ pthread)
+endif()
+
 hip_add_exe_to_target(NAME MultiThreadTest
                      TEST_SRC ${TEST_SRC}
-                      TEST_TARGET_NAME build_tests)
+                      TEST_TARGET_NAME build_tests
+                      LINKER_LIBS ${__LINKER_LIBS__})
@@ -45,7 +45,12 @@ set(TEST_SRC
    # set_source_files_properties(hipStreamAttachMemAsync.cc PROPERTIES COMPILE_FLAGS -std=c++17)
 endif()

+if (UNIX)
+  set(__LINKER_LIBS__ pthread)
+endif()
+
 hip_add_exe_to_target(NAME StreamTest
                      TEST_SRC ${TEST_SRC}
                      TEST_TARGET_NAME build_tests
-                      COMPILE_OPTIONS -std=c++17)
+                      COMPILE_OPTIONS -std=c++17
+                      LINKER_LIBS ${__LINKER_LIBS__})
@@ -7,6 +7,11 @@ set(TEST_SRC
  hipStreamPerThrdTsts.cc
 )

+if (UNIX)
+  set(__LINKER_LIBS__ pthread)
+endif()
+
 hip_add_exe_to_target(NAME StreamPerThreadTest
                      TEST_SRC ${TEST_SRC}
-                      TEST_TARGET_NAME build_tests)
+                      TEST_TARGET_NAME build_tests
+                      LINKER_LIBS ${__LINKER_LIBS__})
@@ -354,7 +354,7 @@ macro(HIT_ADD_FILES _config _dir _label _parent)
                set_source_files_properties(${_sources} PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1)
                hip_reset_flags()
                hip_add_executable(${target} ${_sources} HIPCC_OPTIONS ${_hipcc_options} CLANG_OPTIONS ${_clang_options} NVCC_OPTIONS ${_nvcc_options} EXCLUDE_FROM_ALL)
-                target_link_libraries(${target} PRIVATE ${_link_options})
+                target_link_options(${target} PRIVATE ${_link_options})
                set_target_properties(${target} PROPERTIES OUTPUT_NAME ${_target} RUNTIME_OUTPUT_DIRECTORY ${_label} LINK_DEPENDS "${HIP_LIB_FILES}")
                add_dependencies(${_parent} ${target})
                foreach(_dependency ${_depends})
@@ -21,7 +21,7 @@ THE SOFTWARE.
 */

 /* HIT_START
- * BUILD: %t %s ../test_common.cpp NVCC_OPTIONS -std=c++11
+ * BUILD: %t %s ../test_common.cpp NVCC_OPTIONS -std=c++11 LINK_OPTIONS -lpthread
 * TEST: %t
 * HIT_END
 */
@@ -16,7 +16,7 @@ DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE,
 OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */

 /* HIT_START
- * BUILD: %t %s test_common.cpp NVCC_OPTIONS -std=c++11
+ * BUILD: %t %s test_common.cpp LINK_OPTIONS -lpthread NVCC_OPTIONS -std=c++11
 * TEST: %t
 * HIT_END
 */
@@ -18,7 +18,7 @@ THE SOFTWARE.
 */

 /* HIT_START
- * BUILD: %t %s ../test_common.cpp
+ * BUILD: %t %s ../test_common.cpp LINK_OPTIONS -lrt
 * TEST: %t
 * HIT_END
 */
@@ -18,7 +18,7 @@ THE SOFTWARE.
 */

 /* HIT_START
- * BUILD: %t %s ../test_common.cpp
+ * BUILD: %t %s ../test_common.cpp LINK_OPTIONS -lrt
 * TEST: %t --N 4
 * HIT_END
 */
@@ -21,7 +21,7 @@ THE SOFTWARE.
 */

 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp
+ * BUILD: %t %s ../../test_common.cpp LINK_OPTIONS -lpthread
 * TEST: %t
 * HIT_END
 */
@@ -18,7 +18,7 @@ THE SOFTWARE.
 */

 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp
+ * BUILD: %t %s ../../test_common.cpp LINK_OPTIONS -lrt -lpthread
 * TEST: %t
 * HIT_END
 */
@@ -58,7 +58,7 @@ Testcase Scenarios :
 */

 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 LINK_OPTIONS -lpthread
 * TEST_NAMED: %t hipMalloc_ArgValidation  --tests 1
 * TEST_NAMED: %t hipMalloc_LoopRegression_AllocFreeCycle --tests 2
 * TEST_NAMED: %t hipMalloc_LoopRegression_AllocPool --tests 3
@@ -21,7 +21,7 @@ THE SOFTWARE.
 */

 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 LINK_OPTIONS -lpthread
 * TEST_NAMED: %t hipMemcpy-modes --tests 0x1
 * TEST_NAMED: %t hipMemcpy-size --tests 0x6
 * TEST_NAMED: %t hipMemcpy-dev-offsets --tests 0x10
@@ -24,7 +24,7 @@ THE SOFTWARE.
 // of 8 hipmemcpy apis

 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 LINK_OPTIONS -lpthread
 * TEST_NAMED: %t hipMemcpyNegativeMThrdMSize_Negative_tests  --tests 1
 * TEST_NAMED: %t hipMemcpyNegativeMThrdMSize_MultiThread_tests --tests 2
 * TEST_NAMED: %t hipMemcpyNegativeMThrdMSize_MultiSize_singleType --tests 3 --memcpyPeersOnly 0 --testAllTypes 0
@@ -24,7 +24,7 @@ THE SOFTWARE.
 */

 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 LINK_OPTIONS -lpthread
 * TEST: %t
 * HIT_END
 */
@@ -24,7 +24,7 @@
 //

 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp
+ * BUILD: %t %s ../../test_common.cpp LINK_OPTIONS -lpthread
 * TEST: %t
 * HIT_END
 */
@@ -34,7 +34,7 @@ Testcase Scenarios :
 */

 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 LINK_OPTIONS -lpthread
 * TEST: %t --tests 1
 * HIT_END
 */
@@ -22,7 +22,7 @@
 */

 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp
+ * BUILD: %t %s ../../test_common.cpp LINK_OPTIONS -lpthread
 * TEST: %t
 * HIT_END
 */
@@ -21,7 +21,7 @@
 // and also launch hipMemcpyAsync() api on the same stream. This test case is simulate the scenario
 // reported in SWDEV-181598.
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 LINK_OPTIONS -lpthread
 * TEST: %t
 * HIT_END
 */
@@ -21,7 +21,7 @@
 // and also launch hipMemcpyAsync() api. This test case is simulate the scenario
 // reported in SWDEV-181598.
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 LINK_OPTIONS -lpthread
 * TEST: %t
 * HIT_END
 */
@@ -21,7 +21,7 @@ THE SOFTWARE.
 */

 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp
+ * BUILD: %t %s ../../test_common.cpp LINK_OPTIONS -lpthread
 * TEST: %t
 * HIT_END
 */
@@ -18,7 +18,7 @@ THE SOFTWARE.
 */

 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 LINK_OPTIONS -lpthread
 * TEST: %t
 * HIT_END
 */
@@ -18,7 +18,7 @@ THE SOFTWARE.
 */

 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 EXCLUDE_HIP_PLATFORM nvidia
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 EXCLUDE_HIP_PLATFORM nvidia LINK_OPTIONS -lpthread
 * TEST: %t
 * HIT_END
 */
@@ -19,7 +19,7 @@ THE SOFTWARE.

 /* HIT_START
 * BUILD_CMD: empty_kernel.code %hc --genco %S/empty_kernel.cpp -o empty_kernel.code
- * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 LINK_OPTIONS -lpthread
 * TEST: %t
 * HIT_END
 */
@@ -23,7 +23,7 @@ THE SOFTWARE.
 /* Tests 6 and 7 are skipped for CUDA 11.2 due to cuda runtime issues */
 /* HIT_START
 * BUILD_CMD: tex2d_kernel.code %hc --genco %S/tex2d_kernel.cpp -o tex2d_kernel.code
- * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 LINK_OPTIONS -lpthread
 * TEST: %t --tests 0x01
 * TEST: %t --tests 0x02
 * TEST: %t --tests 0x03
@@ -1,5 +1,5 @@
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 LINK_OPTIONS -lpthread
 * TEST_NAMED: %t hipMultiThreadDevice-serial --tests 0x1
 * TEST_NAMED: %t hipMultiThreadDevice-pyramid --tests 0x4
 * TEST_NAMED: %t hipMultiThreadDevice-nearzero --tests 0x10
@@ -21,7 +21,7 @@ THE SOFTWARE.
 */

 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 LINK_OPTIONS -lpthread
 * TEST: %t
 * HIT_END
 */
@@ -21,7 +21,7 @@ THE SOFTWARE.
 */

 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 LINK_OPTIONS -lpthread
 * TEST: %t
 * HIT_END
 */
@@ -21,7 +21,7 @@ THE SOFTWARE.
 */

 /* HIT_START
- * BUILD_CMD: %t %hc %S/%s -o %T/%t %S/../../test_common.cpp -I %S/../../ -L%rocm-path/rocm_smi/lib -lrocm_smi64 -ldl EXCLUDE_HIP_PLATFORM nvidia
+ * BUILD_CMD: %t %hc %S/%s -o %T/%t %S/../../test_common.cpp -I %S/../../ -L%rocm-path/rocm_smi/lib -lrocm_smi64 -lpthread -ldl EXCLUDE_HIP_PLATFORM nvidia
 * TEST: %t --tests 0x1
 * TEST: %t --tests 0x2
 * TEST: %t --tests 0x3
@@ -21,7 +21,7 @@
 // kernel. Verify that all the kernels queued are executed before the callback.

 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 LINK_OPTIONS -lpthread
 * TEST: %t
 * HIT_END
 */
@@ -22,7 +22,7 @@ THE SOFTWARE.
 // This test case is disabled currently.

 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 LINK_OPTIONS -lpthread
 * TEST: %t
 * HIT_END
 */
@@ -11,7 +11,7 @@
 #include "test_common.h"

 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 LINK_OPTIONS -lpthread
 * TEST: %t
 * HIT_END
 */
@@ -39,7 +39,7 @@ THE SOFTWARE.
 */

 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 EXCLUDE_HIP_PLATFORM nvidia
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 EXCLUDE_HIP_PLATFORM nvidia LINK_OPTIONS -lpthread
 * TEST: %t
 * HIT_END
 */
@@ -19,7 +19,7 @@ THE SOFTWARE.
 // Simple test for Fine Grained CPU-GPU coherency.

 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp HIPCC_OPTIONS -std=c++11 -lpthread EXCLUDE_HIP_PLATFORM nvidia
+ * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvidia LINK_OPTIONS -lpthread
 * TEST: %t
 * HIT_END
 */