diff --git a/lldb/packages/Python/lldbsuite/test/builders/builder.py b/lldb/packages/Python/lldbsuite/test/builders/builder.py index ada6f9ff4a54f..04be523b2ef3a 100644 --- a/lldb/packages/Python/lldbsuite/test/builders/builder.py +++ b/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), ] diff --git a/lldb/packages/Python/lldbsuite/test/configuration.py b/lldb/packages/Python/lldbsuite/test/configuration.py index b2d91fd211477..7b4eeec19b5f4 100644 --- a/lldb/packages/Python/lldbsuite/test/configuration.py +++ b/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 diff --git a/lldb/packages/Python/lldbsuite/test/dotest.py b/lldb/packages/Python/lldbsuite/test/dotest.py index 90c8e32afa507..e6aa606177051 100644 --- a/lldb/packages/Python/lldbsuite/test/dotest.py +++ b/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 diff --git a/lldb/packages/Python/lldbsuite/test/dotest_args.py b/lldb/packages/Python/lldbsuite/test/dotest_args.py index e9c21388bc213..9e33224ff3a08 100644 --- a/lldb/packages/Python/lldbsuite/test/dotest_args.py +++ b/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", diff --git a/lldb/packages/Python/lldbsuite/test/make/Makefile.rules b/lldb/packages/Python/lldbsuite/test/make/Makefile.rules index 58833e1b0cc78..9cf2ee2885b77 100644 --- a/lldb/packages/Python/lldbsuite/test/make/Makefile.rules +++ b/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 diff --git a/lldb/test/API/gpu/amd/basic/Makefile b/lldb/test/API/gpu/amd/basic/Makefile new file mode 100644 index 0000000000000..c3dbfba929f3e --- /dev/null +++ b/lldb/test/API/gpu/amd/basic/Makefile @@ -0,0 +1,3 @@ +HIP_SOURCES := hello_world.hip + +include Makefile.rules diff --git a/lldb/test/API/gpu/amd/basic/TestBasicAmdGpuPlugin.py b/lldb/test/API/gpu/amd/basic/TestBasicAmdGpuPlugin.py new file mode 100644 index 0000000000000..b659cb2a0831e --- /dev/null +++ b/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") diff --git a/lldb/test/API/gpu/amd/basic/hello_world.hip b/lldb/test/API/gpu/amd/basic/hello_world.hip new file mode 100644 index 0000000000000..8da7899ec82b8 --- /dev/null +++ b/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 + +#include + +/// \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()); +} diff --git a/lldb/test/API/gpu/amd/lit.local.cfg b/lldb/test/API/gpu/amd/lit.local.cfg new file mode 100644 index 0000000000000..103bae4aa6476 --- /dev/null +++ b/lldb/test/API/gpu/amd/lit.local.cfg @@ -0,0 +1,2 @@ +if not "lldb-amdgpu" in config.enabled_plugins: + config.unsupported = True diff --git a/lldb/test/API/lit.cfg.py b/lldb/test/API/lit.cfg.py index 83713213ce1fe..935c699142077 100644 --- a/lldb/test/API/lit.cfg.py +++ b/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] diff --git a/lldb/test/API/lit.site.cfg.py.in b/lldb/test/API/lit.site.cfg.py.in index 86d58889cc4ad..89d7048d6bd45 100644 --- a/lldb/test/API/lit.site.cfg.py.in +++ b/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)