Add tests for amdgpu lldb-server plugin

lldb/packages/Python/lldbsuite/test/builders/builder.py

@@ -241,6 +241,11 @@ def getLibCxxArgs(self):
             return libcpp_args
         return []
 
+    def getHipccArgs(self):
+        if configuration.hipcc_path:
+            return ["HIPCC={}".format(configuration.hipcc_path)]
+        return []
+
     def getLLDBObjRoot(self):
         return ["LLDB_OBJ_ROOT={}".format(configuration.lldb_obj_root)]
 
@@ -293,6 +298,7 @@ def getBuildCommand(
             self.getSDKRootSpec(),
             self.getModuleCacheSpec(),
             self.getLibCxxArgs(),
+            self.getHipccArgs(),
             self.getLLDBObjRoot(),
             self.getCmdLine(dictionary),
         ]

lldb/packages/Python/lldbsuite/test/configuration.py

@@ -44,6 +44,7 @@
 dsymutil = None
 sdkroot = None
 make_path = None
+hipcc_path = None
 
 # The overriden dwarf verison.
 # Don't use this to test the current compiler's

lldb/packages/Python/lldbsuite/test/dotest.py

@@ -266,6 +266,9 @@ def parseOptionsAndInitTestdirs():
                     configuration.compiler = candidate
                     break
 
+    if args.hipcc_path:
+        configuration.hipcc_path = os.path.abspath(args.hipcc_path)
+
     if args.make:
         configuration.make_path = args.make
 

lldb/packages/Python/lldbsuite/test/dotest_args.py

@@ -58,6 +58,15 @@ def create_parser():
             """Specify the path to sysroot. This overrides apple_sdk sysroot."""
         ),
     )
+    group.add_argument(
+        "--hipcc-path",
+        metavar="path",
+        dest="hipcc_path",
+        default="",
+        help=textwrap.dedent(
+            """Specify the path to a ROCM hipcc compiler."""
+        ),
+    )
     if sys.platform == "darwin":
         group.add_argument(
             "--apple-sdk",

lldb/packages/Python/lldbsuite/test/make/Makefile.rules

@@ -103,6 +103,17 @@ else
 endif
 endif
 
+#----------------------------------------------------------------------
+# Override CC and CXX when building a hip application.
+#
+# Use the same hipcc compiler for compiling and linking everything.
+#----------------------------------------------------------------------
+ifneq "$(strip $(HIP_SOURCES))" ""
+	override CC := $(HIPCC)
+	override CXX := $(HIPCC)
+	override CC_TYPE := hip
+endif
+
 #----------------------------------------------------------------------
 # CC defaults to clang.
 #
@@ -303,6 +314,7 @@ endif
 
 CFLAGS += $(CFLAGS_EXTRAS)
 CXXFLAGS += -std=c++11 $(CFLAGS) $(ARCH_CXXFLAGS)
+HIPFLAGS ?= $(CFLAGS)
 # Copy common options to the linker flags (dwarf, arch. & etc).
 # Note: we get some 'garbage' options for linker here (such as -I, --isystem & etc).
 LDFLAGS += $(CFLAGS)
@@ -495,6 +507,13 @@ ifneq "$(strip $(OBJCXX_SOURCES))" ""
 	endif
 endif
 
+#----------------------------------------------------------------------
+# Check if we have any hip source files
+#----------------------------------------------------------------------
+ifneq "$(strip $(HIP_SOURCES))" ""
+	OBJECTS +=$(strip $(HIP_SOURCES:.hip=.o))
+endif
+
 ifeq ($(CC_TYPE), clang)
 	CXXFLAGS += --driver-mode=g++
 endif
@@ -632,6 +651,9 @@ endif
 %.o: %.mm %.d
 	$(CXX) $(CXXFLAGS) -MT $@ -MD -MP -MF $*.d -c -o $@ $<
 
+%.o: %.hip %.d
+	$(HIPCC) $(HIPFLAGS) -MT $@ -MD -MP -MF $*.d -c -o $@ $<
+
 #----------------------------------------------------------------------
 # Automatic variables based on items already entered. Below we create
 # an object's lists from the list of sources by replacing all entries

lldb/test/API/gpu/amd/basic/Makefile

@@ -0,0 +1,3 @@
+HIP_SOURCES := hello_world.hip
+
+include Makefile.rules

lldb/test/API/gpu/amd/basic/TestBasicAmdGpuPlugin.py

@@ -0,0 +1,81 @@
+"""
+Basic tests for the AMDGPU plugin.
+"""
+
+
+import lldb
+import lldbsuite.test.lldbutil as lldbutil
+from lldbsuite.test.lldbtest import *
+
+
+class BasicAmdGpuTestCase(TestBase):
+    NO_DEBUG_INFO_TESTCASE = True
+
+    def test_gpu_target_created_on_demand(self):
+        """Test that we create the gpu target automatically."""
+        self.build()
+
+        # There should be no targets before we run the program.
+        self.assertEqual(self.dbg.GetNumTargets(), 0, "There are no targets")
+
+        # Set a breakpoint in the CPU source and run to it.
+        source_spec = lldb.SBFileSpec("hello_world.hip", False)
+        (cpu_target, cpu_process, cpu_thread, cpu_bkpt) = lldbutil.run_to_source_breakpoint(
+            self, "// CPU BREAKPOINT - BEFORE LAUNCH", source_spec
+        )
+        self.assertEqual(self.dbg.GetTargetAtIndex(0), cpu_target)
+
+        # Get the GPU target.
+        # This target is created at the first CPU stop.
+        self.assertEqual(self.dbg.GetNumTargets(), 2, "There are two targets")
+        gpu_target = self.dbg.GetTargetAtIndex(1)
+        gpu_thread = gpu_target.GetProcess().GetThreadAtIndex(0)
+        self.assertEqual(gpu_thread.GetName(), "AMD Native Shadow Thread", "GPU thread has the right name")
+
+    def test_gpu_breakpoint_hit(self):
+        """Test that we create the gpu target automatically."""
+        self.build()
+
+        # Set a breakpoint in the CPU source and run to it.
+        source = "hello_world.hip"
+        source_spec = lldb.SBFileSpec(source, False)
+        (cpu_target, cpu_process, cpu_thread, cpu_bkpt) = lldbutil.run_to_source_breakpoint(
+            self, "// CPU BREAKPOINT - BEFORE LAUNCH", source_spec
+        )
+
+        # Switch to the GPU target so we can set a breakpoint.
+        gpu_target = self.dbg.GetTargetAtIndex(1)
+        gpu_process = gpu_target.GetProcess()
+        self.dbg.SetSelectedTarget(gpu_target)
+
+
+        # Set a breakpoint in the GPU source.
+        # This will not yet resolve to a location.
+        line = line_number(source, "// GPU BREAKPOINT")
+        gpu_bkpt = lldbutil.run_break_set_by_file_and_line(
+            self, source, line, num_expected_locations=0, loc_exact=False
+        )
+
+        # Need to run these commands asynchronously to be able to switch targets.
+        self.setAsync(True)
+        listener = self.dbg.GetListener()
+
+        # Continue the GPU process.
+        self.runCmd("c")
+        lldbutil.expect_state_changes(self, listener, gpu_process, [lldb.eStateRunning])
+
+        # Continue the CPU process.
+        self.dbg.SetSelectedTarget(cpu_target)
+        self.runCmd("c")
+        lldbutil.expect_state_changes(self, listener, cpu_process, [lldb.eStateRunning])
+
+        # TODO: Looks like the CPU is hitting an extra SIGSTOP for some reason so continue again after it stops.
+        lldbutil.expect_state_changes(self, listener, cpu_process, [lldb.eStateStopped])
+        self.dbg.SetSelectedTarget(cpu_target)
+        self.runCmd("c")
+        lldbutil.expect_state_changes(self, listener, cpu_process, [lldb.eStateRunning])
+
+        # GPU breakpoint should get hit.
+        lldbutil.expect_state_changes(self, listener, gpu_process, [lldb.eStateStopped])
+        threads = lldbutil.get_threads_stopped_at_breakpoint_id(gpu_process, gpu_bkpt)
+        self.assertNotEqual(None, threads, "GPU thread should be stopped at breakpoint")

lldb/test/API/gpu/amd/basic/hello_world.hip

@@ -0,0 +1,110 @@
+// MIT License
+//
+// Copyright (c) 2015-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>
+
+#include <iostream>
+
+/// \brief Checks if the provided error code is \p hipSuccess and if not,
+/// prints an error message to the standard error output and terminates the
+/// program with an error code.
+constexpr int error_exit_code = -1;
+#define HIP_CHECK(condition)                                              \
+  {                                                                       \
+    const hipError_t error = condition;                                   \
+    if (error != hipSuccess) {                                            \
+      std::cerr << "An error encountered: \"" << hipGetErrorString(error) \
+                << "\" at " << __FILE__ << ':' << __LINE__ << std::endl;  \
+      std::exit(error_exit_code);                                         \
+    }                                                                     \
+  }
+
+// Functions marked with __device__ are executed on the device and called from
+// the device only.
+__device__ unsigned int get_thread_idx() {
+  // Built-in threadIdx returns the 3D coordinate of the active work item in the
+  // block of threads.
+  return threadIdx.x;
+}
+
+// Functions marked with __global__ are executed on the device and called from
+// the host only.
+__global__ void helloworld_kernel(char *str, size_t size) {
+  unsigned int thread_idx = get_thread_idx();
+  unsigned int block_idx = blockIdx.x;
+  unsigned int idx = thread_idx + block_idx * blockDim.x;
+
+  char c = 'x';
+  switch (idx) {
+    case  0: c = 'H'; break;
+    case  1: c = 'e'; break;
+    case  2: c = 'l'; break;
+    case  3: c = 'l'; break;
+    case  4: c = 'o'; break;
+    case  5: c = ','; break;
+    case  6: c = ' '; break;
+    case  7: c = 'w'; break;
+    case  8: c = 'o'; break;
+    case  9: c = 'r'; break;
+    case 10: c = 'l'; break;
+    case 11: c = 'd'; break;
+    case 12: c = '!'; break;
+  }
+
+  if (idx < size) { // GPU BREAKPOINT
+    str[idx] = c;
+  }
+}
+
+int main() {
+  // Allocate host vectors
+  const size_t size = strlen("Hello, world!");
+  std::string h_str(size, '?');
+  printf("%s\n", h_str.c_str()); // CPU BREAKPOINT - BEFORE LAUNCH
+
+  // Allocate device memory for the output data
+  char *d_str;
+  HIP_CHECK(hipMalloc(&d_str, size));
+
+  // Copy data from host to device
+  printf("Copying data to device...\n");
+  HIP_CHECK(hipMemcpy(d_str, h_str.data(), size, hipMemcpyHostToDevice));
+
+  // Launch the kernel.
+  printf("Launching kernel...\n");
+  helloworld_kernel<<<
+      dim3(1), // 3D grid specifying number of blocks to launch: (1, 1, 1)
+      dim3(h_str.size()), // 3D grid specifying number of threads to launch: (strlen, 1, 1)
+      0, // number of bytes of additional shared memory to allocate
+      hipStreamDefault // stream where the kernel should execute: default stream
+      >>>(d_str, size);
+
+  // Copy data from device to host
+  printf("Copying data to host...\n");   // CPU BREAKPOINT - AFTER LAUNCH
+  HIP_CHECK(hipMemcpy(h_str.data(), d_str, h_str.size(), hipMemcpyDeviceToHost));
+
+  // Free device memory
+  HIP_CHECK(hipFree(d_str)); // CPU BREAKPOINT - AFTER FINISH
+
+  // Print the output
+  printf("%s\n", h_str.c_str());
+}

lldb/test/API/gpu/amd/lit.local.cfg

@@ -0,0 +1,2 @@
+if not "lldb-amdgpu" in config.enabled_plugins:
+    config.unsupported = True

lldb/test/API/lit.cfg.py

@@ -250,6 +250,9 @@ def delete_module_cache(path):
 if is_configured("test_compiler"):
     dotest_cmd += ["--compiler", config.test_compiler]
 
+if is_configured("hipcc_path"):
+    dotest_cmd += ["--hipcc-path", config.hipcc_path]
+
 if is_configured("dsymutil"):
     dotest_cmd += ["--dsymutil", config.dsymutil]
 

lldb/test/API/lit.site.cfg.py.in

@@ -51,6 +51,12 @@ lldb_build_intel_pt = '@LLDB_BUILD_INTEL_PT@'
 if lldb_build_intel_pt == '1':
     config.enabled_plugins.append('intel-pt')
 
+# GPU Plugins
+lldb_enable_amdgpu_plugin = '@LLDB_ENABLE_AMDGPU_PLUGIN@'
+if lldb_enable_amdgpu_plugin == 'ON' or lldb_enable_amdgpu_plugin == '1':
+    config.enabled_plugins.append('lldb-amdgpu')
+    config.hipcc_path = os.path.join('@ROCM_PATH@', 'bin', 'hipcc')
+
 # Additional dotest arguments can be passed to lit by providing a
 # semicolon-separates list: --param dotest-args="arg;arg".
 dotest_lit_args_str = lit_config.params.get('dotest-args', None)