AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

Disabling PC-sampling compilation

Просмотр исходного кода 
Change-Id: I56cc8a1c69ca32dc147cb945b18e708b3292beaf
Этот коммит содержится в:
gobhardw
2024-04-03 18:19:23 +05:30
коммит произвёл Ammar Elwazir
родитель 563b1b023e
Коммит f0daa910d8
16 изменённых файлов: 10 добавлений и 2188 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
CHANGELOG.md
+6
Просмотреть файл
@@ -328,3 +328,9 @@ Example for file plugin output:
- Fixed ROCprofiler to match versioning changes in HIP Runtime.
- Fixed plugins race condition.
- Updated metrics to MI300.
## ROCprofiler for rocm 6.2
### Removed
- pcsampler sample code has been removed due to deprecation from v2.
plugin/cli/cli.cpp
+1
Просмотреть файл
@@ -321,6 +321,7 @@ class file_plugin_t {
break;
}
case ROCPROFILER_PC_SAMPLING_RECORD: {
[[deprecated("PC Sampling is deprecated")]]
const rocprofiler_record_pc_sample_t* pc_sampling_record =
reinterpret_cast<const rocprofiler_record_pc_sample_t*>(begin);
FlushPCSamplingRecord(pc_sampling_record);
plugin/file/file.cpp
+1
Просмотреть файл
@@ -443,6 +443,7 @@ class file_plugin_t {
break;
}
case ROCPROFILER_PC_SAMPLING_RECORD: {
[[deprecated("PC Sampling is deprecated")]]
const rocprofiler_record_pc_sample_t* pc_sampling_record =
reinterpret_cast<const rocprofiler_record_pc_sample_t*>(begin);
FlushPCSamplingRecord(pc_sampling_record);
plugin/file/file_v1.cpp
+1
Просмотреть файл
@@ -441,6 +441,7 @@ class file_plugin_t {
break;
}
case ROCPROFILER_PC_SAMPLING_RECORD: {
[[deprecated("PC Sampling is deprecated")]]
const rocprofiler_record_pc_sample_t* pc_sampling_record =
reinterpret_cast<const rocprofiler_record_pc_sample_t*>(begin);
FlushPCSamplingRecord(pc_sampling_record);
samples/CMakeLists.txt
-44
Просмотреть файл
@@ -184,50 +184,6 @@ install(TARGETS tracer_hip_hsa_async
RUNTIME DESTINATION ${CMAKE_INSTALL_DATAROOTDIR}/${PROJECT_NAME}/samples
COMPONENT samples)
# ########################################################################################
# PC Sampling Samples
# ########################################################################################
set(CODE_PRINTING_SAMPLE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/pcsampler/code_printing_sample)
file(GLOB PC_SAMPLING_CODE_PRINTING_FILES ${CODE_PRINTING_SAMPLE_DIR}/*.cpp)
set_source_files_properties(${PC_SAMPLING_CODE_PRINTING_FILES}
PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1)
hip_add_executable(
pc_sampling_code_printing ${PC_SAMPLING_CODE_PRINTING_FILES} HIPCC_OPTIONS -std=c++17
# Include debugging symbols and source for the contextual disassembly
-gdwarf-4)
rocprofiler_sample_add_test(pc_sampling_code_printing "-d;0;-n;100000000;10;43532")
check_c_source_compiles(
"
#define _GNU_SOURCE
#include <sys/mman.h>
int main() { return memfd_create (\"cmake_test\", 0); }
"
HAVE_MEMFD_CREATE)
if(HAVE_MEMFD_CREATE)
target_compile_definitions(pc_sampling_code_printing PRIVATE HAVE_MEMFD_CREATE)
endif()
target_link_libraries(
pc_sampling_code_printing
PRIVATE rocprofiler-v2 rocm-dbgapi ${LIBELF_LIBRARIES} ${LIBDW_LIBRARIES}
hsa-runtime64::hsa-runtime64 Threads::Threads dl)
target_include_directories(
pc_sampling_code_printing PRIVATE ${TEST_DIR} ${ROOT_DIR} ${HSA_RUNTIME_INC_PATH}
${PROJECT_SOURCE_DIR})
target_link_options(pc_sampling_code_printing PRIVATE "-Wl,--build-id=md5")
add_dependencies(samples pc_sampling_code_printing)
install(TARGETS pc_sampling_code_printing
RUNTIME DESTINATION ${CMAKE_INSTALL_DATAROOTDIR}/${PROJECT_NAME}/samples
COMPONENT samples)
install(
DIRECTORY "${PROJECT_SOURCE_DIR}/samples/"
DESTINATION ${CMAKE_INSTALL_DATAROOTDIR}/${PROJECT_NAME}/samples-src
OPTIONAL
COMPONENT samples)
# ########################################################################################
# Scripts to run samples
# ########################################################################################
samples/common/common.h
+1
Просмотреть файл
@@ -249,6 +249,7 @@ int WriteBufferRecords(const rocprofiler_record_header_t* begin,
break;
}
case ROCPROFILER_PC_SAMPLING_RECORD: {
[[deprecated("PC Sampling is deprecated")]]
const rocprofiler_record_pc_sample_t* pc_sampling_record =
reinterpret_cast<const rocprofiler_record_pc_sample_t*>(begin);
FlushPCSamplingRecord(pc_sampling_record);
samples/pcsampler/code_printing_sample/.clangd
-10
Просмотреть файл
@@ -1,10 +0,0 @@
---
If:
PathMatch: main.cpp
CompileFlags:
Add: ['-x', 'hip']
# Local Variables:
# mode: yaml
# End:
samples/pcsampler/code_printing_sample/Makefile
-70
Просмотреть файл
@@ -1,70 +0,0 @@
# -*- makefile-gmake -*-
ROCM_PATH ?= /opt/rocm
HIP_PATH ?= $(ROCM_PATH)/hip
HIPCC := $(HIP_PATH)/bin/hipcc
ROCM_PATH ?=/opt/rocm
ROCPROFILER_LIBS_PATH ?=$(ROCM_PATH)/lib
ROCPROFILER_INCLUDES=$(ROCPROFILER_LIBS_PATH)/../include/rocprofiler/
ifndef ROCPROFILER_PATH
$(warning You may need to set ROCPROFILER_PATH to the path of the rocprofiler source)
endif
CXXFLAGS += -std=c++17 -Wall
ifdef DEBUG
CXXFLAGS += -gdwarf-4 -O0
else
ifdef DEBUGOPT
CXXFLAGS += -gdwarf-4 -Og
else
CXXFLAGS += -gdwarf-4 -O2
endif
endif
###
srcs := $(wildcard *.cpp)
prog := main
objs := $(srcs:%.cpp=%.o)
deps := $(srcs:%.cpp=%.d)
# Kernel program
CPPFLAGS += -DHAVE_MEMFD_CREATE
$(prog): CC = $(HIPCC)
$(prog): CPPFLAGS += -I$(ROCPROFILER_INCLUDES) -I$(ROCM_PATH)/include
$(prog): LDFLAGS := -L$(ROCPROFILER_LIBS_PATH) -L$(ROCM_PATH)/lib
$(prog): LDLIBS += -ldl -lpthread -lhsa-runtime64 -lrocprofiler64v2 -lrocm-dbgapi -ldw -lelf
$(objs): CXX = $(HIPCC)
# Targets
all: $(prog)
$(prog): $(objs)
-include $(deps)
OUTPUT_OPTION = -MMD -MP -o $@
%.so: %.o
$(LINK.o) $(OUTPUT_OPTION) $^ $(LDLIBS)
#COMPILE.hip = $(COMPILE.cpp)
#LINK.hip = $(LINK.cpp)
#%.o: %.hip
# $(COMPILE.hip) $(OUTPUT_OPTION) $<
clean:
$(RM) $(prog) $(objs) $(deps)
distclean: | clean
$(RM) compile_commands.json
.PHONY: all clean distclean
samples/pcsampler/code_printing_sample/README.md
-149
Просмотреть файл
@@ -1,149 +0,0 @@
# ROCProfiler PC sampling example code
The ROCProfiler library includes an API to enable periodic sampling of the GPU
program counter during kernel execution. This program demonstrates the PC
sampling API, with additional code to illustrate a typical non-trivial use case:
correlation of sampled PC addresses with their disassembled machine code, as
well as source code and symbolic debugging information if available.
## Building the demo program
If your ROCm installation already includes ROCProfiler, the only requirements to
build the demo program are:
* GNU `make`
* libdw (**not** libdwarf)
* libelf
If ROCm is installed in the standard location (`/opt/rocm`), running `make` in
the same directory as this README should work; otherwise, set `ROCM_PATH` to the
location of the ROCm installation in your environment and `ROCPROFILER_PATH` to
the location of the ROCProfiler source repo before running `make`.
If your ROCm installation does **not** include ROCProfiler, you will need to build
it yourself. This demo program will be built as part of that process. See the
main ROCProfiler README for additional requirements and directions.
## Running the demo program
The demo program simply fills a vector with random 64-bit unsigned integers and
tallies the count of those greater than the mandatory `MIN` argument:
```
usage: code_printing_sample [OPTION]... MIN [SEED]
-d DEV HIP device number
-n LEN Length of random integer array
-D Print kernel disassembly
-P Print source and disassembly of sampled PC locations
where
DEV : i32
MIN : u64
LEN : u64
SEED : u64
```
### Defaults and troubleshooting
* `-d`: use HIP device 0
* `-n`: 4194304 (1024 * 1024 * 4)
* `-D`: false
* `-P`: false
* `SEED`: random seed; taken from the system's monotonic clock
The program contains two trivial GPU kernels: an implementation of `memset`, and
the parallel counting procedure. Because the actual point is to demonstrate the
PC sampling functionality, it is recommended to use the `-n` option with an
argument such that the allocated vector fits in the smaller of available host as
well as device memory, but sufficiently large argument such that the kernels run
long enough for ROCProfiler to actually collect some samples.
In order for the `-P` option to display source, the demo program must have been
built with debug symbols (at least `-gdwarf-4`). Any optimization level is
fine, but if the kernels run too quickly for ROCProfiler to collect any samples
even when a very large vector is given with the `-n` option, try rebuilding the
demo program without optimizations by adding `-O0` to the `hipcc` compilation
flags.
## Files
* `main.cpp`: initializes ROCProfiler and PC sampling and runs the GPU kernels
* `code_printing.cpp`: inspects the ELF and DWARF info for the GPU programs
embedded in the host binary and uses amd-dbgapi to print disassembly and
source
* `disassembly.cpp`: wrapper for `code_printing.cpp`
## PC sampling API
Adding PC sampling to a program already using the ROCProfiler API requires only
two changes:
1. Call `rocprofiler_create_filter` to create a `ROCPROFILER_PC_SAMPLING_COLLECTION`
filter, then `rocprofiler_set_filter_buffer` to add the filter to the desired
buffer (see functions `main` and `run_kernel` in `main.cpp`)
2. Handle records of kind `ROCPROFILER_PC_SAMPLING_RECORD` in the buffer callback
function. These should be cast to `rocprofiler_record_pc_sample_t *` (see
function `callback_flush_fn` in `main.cpp`)
Like all ROCProfiler records, PC sample records contain a standard header followed
by one or more payloads:
```c
/**
* PC sample record: contains the program counter/instruction pointer observed
* during periodic sampling of a kernel
*/
typedef struct {
/**
* ROCProfiler General Record base header to identify the id and kind of every
* record
*/
rocprofiler_record_header_t header;
/**
* PC sample data
*/
rocprofiler_pc_sample_t pc_sample;
} rocprofiler_record_pc_sample_t;
```
PC samples are delivered via the normal ROCProfiler buffer callback mechanism,
along with some additional information allowing each sample to be associated
with a unique, individual kernel execution:
```c
/**
* An individual PC sample
*/
typedef struct {
/**
* Kernel dispatch ID. This is used by PC sampling to associate samples with
* individual dispatches and is unrelated to any user-supplied correlation ID
*/
rocprofiler_kernel_dispatch_id_t dispatch_id;
union {
/**
* Host timestamp
*/
rocprofiler_timestamp_t timestamp;
/**
* GPU clock counter (not currently used)
*/
uint64_t cycle;
};
/**
* Sampled program counter
*/
uint64_t pc;
/**
* Sampled shader element
*/
uint32_t se;
/**
* Sampled GPU agent
*/
rocprofiler_agent_id_t gpu_id;
} rocprofiler_pc_sample_t;
```
PC sampling is started and stopped with `rocprofiler_start_session` and
`rocprofiler_terminate_session`, just like other profiling activities.
samples/pcsampler/code_printing_sample/code_printing.cpp
-1122
Просмотреть файл
Разница между файлами не показана из-за своего большого размера Загрузить разницу
samples/pcsampler/code_printing_sample/code_printing.hpp
-104
Просмотреть файл
@@ -1,104 +0,0 @@
/* Copyright (c) 2022 Advanced Micro Devices, Inc.
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. */
#ifndef SAMPLES_PCSAMPLER_CODE_PRINTING_SAMPLE_CODE_PRINTING_HPP_
#define SAMPLES_PCSAMPLER_CODE_PRINTING_SAMPLE_CODE_PRINTING_HPP_
#include <map>
#include <optional>
#include <string>
#include <vector>
#include <amd-dbgapi/amd-dbgapi.h>
namespace amd::debug_agent {
class code_object_t {
struct symbol_info_t {
const std::string m_name;
amd_dbgapi_global_address_t m_value;
amd_dbgapi_size_t m_size;
};
using symbol_map_t = std::optional<
std::map<amd_dbgapi_global_address_t, std::pair<std::string, amd_dbgapi_size_t>>>;
public:
void load_symbol_map();
void load_debug_info();
std::optional<symbol_info_t> find_symbol(amd_dbgapi_global_address_t address);
code_object_t(amd_dbgapi_code_object_id_t code_object_id);
code_object_t(code_object_t&& rhs);
~code_object_t();
void open();
bool is_open() const { return m_fd.has_value(); }
amd_dbgapi_global_address_t load_address() const { return m_load_address; }
amd_dbgapi_size_t mem_size() const { return m_mem_size; }
// FIXME(?): extra function not in rocr-debug-agent
uint32_t elf_amdgpu_machine() const { return m_elf_amdgpu_machine; }
void disassemble_around(amd_dbgapi_architecture_id_t architecture_id,
amd_dbgapi_global_address_t pc);
void disassemble_kernel(amd_dbgapi_architecture_id_t architecture_id,
amd_dbgapi_global_address_t start_addr, bool const print_src = false);
bool save(const std::string& directory) const;
amd_dbgapi_global_address_t m_load_address{0};
amd_dbgapi_size_t m_mem_size{0};
std::optional<int> m_fd;
std::optional<std::map<amd_dbgapi_global_address_t, std::pair<std::string, size_t>>>
m_line_number_map;
std::optional<std::map<amd_dbgapi_global_address_t, amd_dbgapi_global_address_t>> m_pc_ranges_map;
symbol_map_t m_symbol_map;
std::string m_uri;
amd_dbgapi_code_object_id_t const m_code_object_id;
// FIXME(?): extra field not in rocr-debug-agent
uint32_t m_elf_amdgpu_machine{0};
};
} // namespace amd::debug_agent
enum struct disassembly_mode { AROUND, KERNEL };
std::tuple<amd_dbgapi_process_id_t,
std::map<amd_dbgapi_global_address_t, amd::debug_agent::code_object_t>>
init_disassembly();
void disassemble(
disassembly_mode const mode, amd_dbgapi_process_id_t const process_id,
std::map<amd_dbgapi_global_address_t, amd::debug_agent::code_object_t>& code_object_map,
uint64_t const addr);
void print_pc_context(
amd_dbgapi_process_id_t const process_id,
std::map<amd_dbgapi_global_address_t, amd::debug_agent::code_object_t>& code_object_map,
amd_dbgapi_global_address_t const pc);
#endif // SAMPLES_PCSAMPLER_CODE_PRINTING_SAMPLE_CODE_PRINTING_HPP_
samples/pcsampler/code_printing_sample/disassembly.cpp
-176
Просмотреть файл
@@ -1,176 +0,0 @@
/* Copyright (c) 2022 Advanced Micro Devices, Inc.
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 <algorithm>
#include <atomic>
#include <functional>
#include <map>
#include <mutex>
#include <optional>
#include <string>
#include <tuple>
#include <type_traits>
#include <unordered_map>
#include <vector>
#include <cassert>
#include <cinttypes>
#include <cstdint>
#include <cstdio>
#include <sys/mman.h>
#include <hsa/hsa.h>
#include <amd-dbgapi/amd-dbgapi.h>
#include <hsa/amd_hsa_kernel_code.h>
#include <hsa/hsa_ven_amd_loader.h>
#include <rocprofiler/v2/rocprofiler.h>
#include "code_printing.hpp"
#include "program.hpp"
struct libc_freer {
void operator()(char* p) { free(p); }
};
namespace util {
template <typename T, typename... Ts>
static void hash_combine(size_t& hsh, T const& v, Ts const&... rest) {
hsh ^= std::hash<T>{}(v) + 0x9e3779b9 + (hsh << 6) + (hsh >> 2);
(hash_combine(hsh, rest), ...);
}
} // namespace util
[[maybe_unused]] static inline bool operator==(hsa_executable_t const& l,
hsa_executable_t const& r) {
return l.handle == r.handle;
}
[[maybe_unused]] static inline bool operator==(rocprofiler_kernel_dispatch_id_t const& l,
rocprofiler_kernel_dispatch_id_t const& r) {
return l.value == r.value;
}
static inline bool operator==(amd_dbgapi_process_id_t const& l, amd_dbgapi_process_id_t const& r) {
return l.handle == r.handle;
}
static inline bool operator!=(amd_dbgapi_process_id_t const& l, amd_dbgapi_process_id_t const& r) {
return !(l == r);
}
namespace std {
template <> struct hash<hsa_executable_t> {
size_t operator()(hsa_executable_t const& v) const {
size_t ret = 0;
util::hash_combine(ret, v.handle);
return ret;
}
};
template <> struct hash<rocprofiler_kernel_dispatch_id_t> {
size_t operator()(rocprofiler_kernel_dispatch_id_t const& v) const {
size_t ret = 0;
util::hash_combine(ret, v.value);
return ret;
}
};
} // namespace std
struct disassembly_ctx_t {
disassembly_ctx_t();
~disassembly_ctx_t();
void disassemble_kernels(bool const reinitialize);
void init();
bool inited() const;
void reset();
amd_dbgapi_process_id_t process_id;
std::map<amd_dbgapi_global_address_t, amd::debug_agent::code_object_t> codeobjs;
};
disassembly_ctx_t::disassembly_ctx_t() : process_id(AMD_DBGAPI_PROCESS_NONE), codeobjs() {}
disassembly_ctx_t::~disassembly_ctx_t() { reset(); }
void disassembly_ctx_t::disassemble_kernels(bool const reinitialize) {
if (reinitialize) {
reset();
}
if (!inited()) {
init();
}
auto it = codeobjs.begin();
auto const end = codeobjs.end();
auto const pred = [](decltype(*it)& x) {
/*
* A lame filter for the kernels in the current file, because nothing
* else in this little demo will have the URL prefix of `file://`.
*/
return x.second.m_uri.find("file://", 0, 7) != std::string::npos;
};
while (end != (it = std::find_if(it, end, pred))) {
auto& codeobj = it->second;
codeobj.load_symbol_map();
if (!codeobj.m_symbol_map) {
fputs(PROGNAME ": error: failed to load symbol map\n", stderr);
break;
}
for (auto const& sym : *codeobj.m_symbol_map) {
auto const& addr = sym.first;
::disassemble(disassembly_mode::KERNEL, process_id, codeobjs, addr);
}
++it;
}
}
inline void disassembly_ctx_t::init() { std::tie(process_id, codeobjs) = init_disassembly(); }
inline bool disassembly_ctx_t::inited() const { return AMD_DBGAPI_PROCESS_NONE != process_id; }
void disassembly_ctx_t::reset() {
codeobjs.clear();
if (AMD_DBGAPI_PROCESS_NONE.handle != process_id.handle) {
amd_dbgapi_process_detach(process_id);
amd_dbgapi_finalize();
process_id = AMD_DBGAPI_PROCESS_NONE;
}
}
static disassembly_ctx_t g_dis;
void disassembly_disassemble_kernels(bool const reinitialize) {
g_dis.disassemble_kernels(reinitialize);
}
void disassembly_print_pc_sample_context(amd_dbgapi_global_address_t const pc) {
if (!g_dis.inited()) {
g_dis.init();
}
print_pc_context(g_dis.process_id, g_dis.codeobjs, pc);
}
samples/pcsampler/code_printing_sample/disassembly.hpp
-30
Просмотреть файл
@@ -1,30 +0,0 @@
/* Copyright (c) 2022 Advanced Micro Devices, Inc.
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. */
#ifndef SAMPLES_PCSAMPLER_CODE_PRINTING_SAMPLE_DISASSEMBLY_HPP_
#define SAMPLES_PCSAMPLER_CODE_PRINTING_SAMPLE_DISASSEMBLY_HPP_
#include <amd-dbgapi/amd-dbgapi.h>
void disassembly_disassemble_kernels(bool const);
void disassembly_print_pc_sample_context(amd_dbgapi_global_address_t const);
#endif // SAMPLES_PCSAMPLER_CODE_PRINTING_SAMPLE_DISASSEMBLY_HPP_
samples/pcsampler/code_printing_sample/main.cpp
-383
Просмотреть файл
@@ -1,383 +0,0 @@
/* Copyright (c) 2022 Advanced Micro Devices, Inc.
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 <algorithm>
#include <chrono>
#include <memory>
#include <numeric>
#include <vector>
#include <cfloat>
#include <cinttypes>
#include <cstdint>
#include <cstdlib>
#include <unistd.h>
#include <hip/hip_runtime.h>
#include <hsa/hsa.h>
#include <rocprofiler/v2/rocprofiler.h>
#include "program.hpp"
#include "program_options.hpp"
#include "disassembly.hpp"
#define XSTR(x) STR(x)
#define STR(x) #x
#define DBL_FMT "." XSTR(DBL_DECIMAL_DIG) "f"
namespace util {
struct hipMalloc_freer {
void operator()(void* const ptr) { (void)hipFree(ptr); }
};
} // namespace util
namespace prng {
static uint64_t splitmix64_next(uint64_t* const sm64_state) {
uint64_t z = (*sm64_state += 0x9e3779b97f4a7c15);
z = (z ^ (z >> 30)) * 0xbf58476d1ce4e5b9;
z = (z ^ (z >> 27)) * 0x94d049bb133111eb;
return z ^ (z >> 31);
}
static inline uint64_t rotl64(const uint64_t x, int k) { return (x << k) | (x >> (64 - k)); }
static uint64_t xrs_next(uint64_t* const xrs_state) {
const uint64_t result = rotl64(xrs_state[0] + xrs_state[3], 23) + xrs_state[0];
const uint64_t t = xrs_state[1] << 17;
xrs_state[2] ^= xrs_state[0];
xrs_state[3] ^= xrs_state[1];
xrs_state[1] ^= xrs_state[2];
xrs_state[0] ^= xrs_state[3];
xrs_state[2] ^= t;
xrs_state[3] = rotl64(xrs_state[3], 45);
return result;
}
} // namespace prng
namespace kernel {
template <typename T> __global__ static void memset_gpu(T* const s, T const c, size_t const n) {
size_t i_start = threadIdx.x + blockIdx.x * blockDim.x;
size_t i_shift = blockDim.x * gridDim.x;
for (size_t i = i_start; i < n; i += i_shift) {
s[i] = c;
}
}
template <typename T>
__global__ static void count_gpu(T const* const xs, T* const out, size_t const n,
size_t const nblocks, T const gt) {
size_t i_start = threadIdx.x + blockIdx.x * blockDim.x;
size_t i_shift = blockDim.x * gridDim.x;
for (size_t i = i_start; i < n; i += i_shift) {
if (xs[i] > gt) {
atomicAdd(&out[i % nblocks], 1);
}
}
}
} // namespace kernel
static char const GETOPT_ARGS[] = "cd:mn:DP";
static void usage() {
fputs("usage: " PROGNAME
" [OPTION]... MIN [SEED]\n"
" -d DEV\tHIP device number\n"
" -n LEN\tLength of random integer array\n"
" -D\t\tPrint kernel disassembly\n"
" -P\t\tPrint source and disassembly of sampled PC locations\n"
"where\n"
" DEV : i32\n"
" MIN : u64\n"
" LEN : u64\n"
" SEED : u64\n",
stderr);
}
static int get_options(int argc, char** argv, program_options* const opts) {
int opt;
while (-1 != (opt = getopt(argc, argv, GETOPT_ARGS))) {
switch (opt) {
case 'd':
// TODO error checking
opts->device = strtol(optarg, nullptr, 10);
break;
case 'n':
// TODO error checking
opts->rands_len = strtoul(optarg, nullptr, 10);
break;
case 'D':
opts->disassemble = true;
break;
case 'P':
opts->pc_sampling = true;
break;
default:
usage();
return EXIT_FAILURE;
}
}
auto const optcount = argc - optind;
if (!(1 == optcount || 2 == optcount)) {
usage();
return EXIT_FAILURE;
}
// TODO error checking
opts->gt = strtoul(argv[optind], nullptr, 10);
if (2 == argc - optind) {
opts->seed = strtoull(argv[optind + 1], nullptr, 10);
}
return EXIT_SUCCESS;
}
static program_options g_opts;
static void callback_flush_fn(rocprofiler_record_header_t const* record,
rocprofiler_record_header_t const* end_record,
rocprofiler_session_id_t session_id,
rocprofiler_buffer_id_t buffer_id) {
while (record < end_record) {
if (nullptr == record) {
break;
}
if (ROCPROFILER_PC_SAMPLING_RECORD == record->kind) {
auto const& pcr = (rocprofiler_record_pc_sample_t&)*record;
printf("dispatch[%" PRIu64 "] timestamp(%" PRIu64 ") gpu_id(%#" PRIx64 ") pc-sample(%#" PRIx64
") se(%" PRIu32 ")\n",
pcr.pc_sample.dispatch_id.value, pcr.pc_sample.timestamp.value,
pcr.pc_sample.gpu_id.handle, pcr.pc_sample.pc, pcr.pc_sample.se);
if (g_opts.pc_sampling) {
disassembly_print_pc_sample_context(pcr.pc_sample.pc);
}
}
rocprofiler_next_record(record, &record, session_id, buffer_id);
}
}
static int run_kernel(program_options const& opts) {
rocprofiler_session_id_t sid;
rocprofiler_filter_id_t fid, fid2;
rocprofiler_buffer_id_t bid;
auto rocprofiler_ok = ROCPROFILER_STATUS_SUCCESS;
if (opts.pc_sampling) {
ROCPROFILER_CHECK(rocprofiler_create_session(ROCPROFILER_NONE_REPLAY_MODE, &sid),
rocprofiler_ok);
if (ROCPROFILER_STATUS_SUCCESS != rocprofiler_ok) {
fputs("error: failed to create rocprofiler session\n", stderr);
return EXIT_FAILURE;
}
rocprofiler_filter_property_t property{};
ROCPROFILER_CHECK(
rocprofiler_create_buffer(sid, callback_flush_fn, static_cast<size_t>(0x1000), &bid),
rocprofiler_ok);
if (ROCPROFILER_STATUS_SUCCESS != rocprofiler_ok) {
fputs("error: failed to add PC sampling session mode\n", stderr);
goto out;
}
ROCPROFILER_CHECK(rocprofiler_create_filter(sid, ROCPROFILER_PC_SAMPLING_COLLECTION,
rocprofiler_filter_data_t{}, 0, &fid, property),
rocprofiler_ok);
if (ROCPROFILER_STATUS_SUCCESS != rocprofiler_ok) {
goto cleanup;
}
ROCPROFILER_CHECK(rocprofiler_create_filter(sid, ROCPROFILER_DISPATCH_TIMESTAMPS_COLLECTION,
rocprofiler_filter_data_t{}, 0, &fid2, property),
rocprofiler_ok);
if (ROCPROFILER_STATUS_SUCCESS != rocprofiler_ok) {
goto cleanup;
}
ROCPROFILER_CHECK(rocprofiler_set_filter_buffer(sid, fid, bid), rocprofiler_ok);
if (ROCPROFILER_STATUS_SUCCESS != rocprofiler_ok) {
goto cleanup;
}
ROCPROFILER_CHECK(rocprofiler_set_filter_buffer(sid, fid2, bid), rocprofiler_ok);
if (ROCPROFILER_STATUS_SUCCESS != rocprofiler_ok) {
goto cleanup;
}
ROCPROFILER_CHECK(rocprofiler_start_session(sid), rocprofiler_ok);
if (ROCPROFILER_STATUS_SUCCESS != rocprofiler_ok) {
goto cleanup;
}
}
{
printf("seed = %" PRIu64 "\n", opts.seed);
std::vector<uint64_t> rands(opts.rands_len);
using rands_elt_t = decltype(rands)::value_type;
uint64_t sm64_state = opts.seed, xrs_state[4];
{
using prng::splitmix64_next;
using prng::xrs_next;
// Initialize the Xoroshiro PRNG
xrs_state[0] = splitmix64_next(&sm64_state);
xrs_state[1] = splitmix64_next(&sm64_state);
xrs_state[2] = splitmix64_next(&sm64_state);
xrs_state[3] = splitmix64_next(&sm64_state);
// Fill rands with random integers
for (auto& i : rands) {
i = xrs_next(xrs_state);
}
}
struct tm {
using monoclk = std::chrono::steady_clock;
using dur = std::chrono::duration<double>;
};
using util::hipMalloc_freer;
auto const begin_time = tm::monoclk::now();
auto hip_ok = hipSuccess;
do {
HIP_CHECK_BREAK(hipSetDevice(opts.device), hip_ok);
auto const rands_nbytes = rands.size() * sizeof(rands_elt_t);
std::unique_ptr<rands_elt_t, hipMalloc_freer> rands_gpu;
{
rands_elt_t* rands_gpu_ptr;
HIP_CHECK_BREAK(hipMalloc(&rands_gpu_ptr, rands_nbytes), hip_ok);
rands_gpu.reset(rands_gpu_ptr);
}
HIP_CHECK_BREAK(hipMemcpy(rands_gpu.get(), rands.data(), rands_nbytes, hipMemcpyHostToDevice),
hip_ok);
(void)hipDeviceSynchronize();
uint32_t constexpr nthreads = 256U;
uint32_t const nblocks = (rands.size() + nthreads - 1) / nthreads;
using count_elt_t = size_t;
auto const count_subtotals_nbytes = nblocks * sizeof(count_elt_t);
std::unique_ptr<count_elt_t, hipMalloc_freer> count_subtotals_gpu;
{
count_elt_t* count_subtotals_gpu_ptr;
HIP_CHECK_BREAK(hipMalloc(&count_subtotals_gpu_ptr, count_subtotals_nbytes), hip_ok);
count_subtotals_gpu.reset(count_subtotals_gpu_ptr);
}
hipLaunchKernelGGL(kernel::memset_gpu, nblocks, nthreads, 0, 0, count_subtotals_gpu.get(),
0UL, static_cast<size_t>(nblocks));
HIP_CHECK_BREAK(hipGetLastError(), hip_ok);
(void)hipDeviceSynchronize();
auto const kernel_begin_time = tm::monoclk::now();
hipLaunchKernelGGL(kernel::count_gpu, nblocks, nthreads, 0, 0, rands_gpu.get(),
count_subtotals_gpu.get(), rands.size(), static_cast<size_t>(nblocks),
opts.gt);
HIP_CHECK_BREAK(hipGetLastError(), hip_ok);
(void)hipDeviceSynchronize();
auto const kernel_end_time = tm::monoclk::now();
std::vector<size_t> count_subtotals(nblocks);
HIP_CHECK_BREAK(hipMemcpy(count_subtotals.data(), count_subtotals_gpu.get(),
count_subtotals_nbytes, hipMemcpyDeviceToHost),
hip_ok);
(void)hipDeviceSynchronize();
// TODO parallel sum on GPU
auto const total =
std::accumulate(count_subtotals.cbegin(), count_subtotals.cend(), static_cast<size_t>(0));
auto const all_end_time = tm::monoclk::now();
tm::dur const kernel_time(kernel_end_time - kernel_begin_time);
auto total_time(all_end_time - begin_time);
tm::dur const total_time_without_tool_init(total_time);
printf(
"len(rands) = %zu; gt = %zu; count(rands, gt) = %zu\n"
"main kernel time elapsed: %" DBL_FMT
"\n"
"full time elapsed: %" DBL_FMT "\n",
rands.size(), opts.gt, total, kernel_time.count(), total_time_without_tool_init.count());
} while (false);
if (opts.disassemble) {
disassembly_disassemble_kernels(false);
}
}
cleanup:
if (opts.pc_sampling) {
rocprofiler_terminate_session(sid);
rocprofiler_flush_data(sid, bid);
rocprofiler_destroy_session(sid);
}
out:
return ROCPROFILER_STATUS_SUCCESS == rocprofiler_ok ? EXIT_SUCCESS : EXIT_FAILURE;
}
int main(int argc, char** argv) {
if (auto const ret = get_options(argc, argv, &g_opts); EXIT_SUCCESS != ret) {
return ret;
}
if (hsa_init() != HSA_STATUS_SUCCESS) {
return EXIT_FAILURE;
}
int ret = EXIT_FAILURE;
auto ok = ROCPROFILER_STATUS_SUCCESS;
ROCPROFILER_CHECK(rocprofiler_initialize(), ok);
if (ROCPROFILER_STATUS_SUCCESS == ok) {
ret = run_kernel(g_opts);
} else {
goto out;
}
rocprofiler_finalize();
out:
hsa_shut_down();
return ROCPROFILER_STATUS_SUCCESS == ok && EXIT_FAILURE != ret ? EXIT_SUCCESS : EXIT_FAILURE;
}
samples/pcsampler/code_printing_sample/program.hpp
-52
Просмотреть файл
@@ -1,52 +0,0 @@
/* Copyright (c) 2022 Advanced Micro Devices, Inc.
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. */
#ifndef SAMPLES_PCSAMPLER_CODE_PRINTING_SAMPLE_PROGRAM_HPP_
#define SAMPLES_PCSAMPLER_CODE_PRINTING_SAMPLE_PROGRAM_HPP_
#define PROGNAME "code_printing_sample"
#define HIP_ERROR(code) \
do { \
fprintf(stderr, PROGNAME ": Assertion failed at %s:%d, HIP error: %s\n", __FILE__, __LINE__, \
hipGetErrorString((code))); \
fflush(stderr); \
} while (false);
#define HIP_CHECK_BREAK(expr, var) \
if (auto const code = (expr); hipSuccess != code) { \
HIP_ERROR(code); \
(var) = code; \
break; \
}
#define ROCPROFILER_ERROR(code) \
do { \
fprintf(stderr, PROGNAME ": Assertion failed at %s:%d, ROCProfiler error: %s\n", __FILE__, \
__LINE__, rocprofiler_error_str(code)); \
fflush(stderr); \
} while (false);
#define ROCPROFILER_CHECK(expr, var) \
if ((var) = (expr); ROCPROFILER_STATUS_SUCCESS != (var)) { \
ROCPROFILER_ERROR((var)); \
}
#endif // SAMPLES_PCSAMPLER_CODE_PRINTING_SAMPLE_PROGRAM_HPP_
samples/pcsampler/code_printing_sample/program_options.hpp
-48
Просмотреть файл
@@ -1,48 +0,0 @@
/* Copyright (c) 2022 Advanced Micro Devices, Inc.
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. */
#ifndef SAMPLES_PCSAMPLER_CODE_PRINTING_SAMPLE_PROGRAM_OPTIONS_HPP_
#define SAMPLES_PCSAMPLER_CODE_PRINTING_SAMPLE_PROGRAM_OPTIONS_HPP_
#include <chrono>
#include <cstdint>
struct program_options {
program_options()
: device(0),
no_gpu(false),
hip_memset(false),
rands_len(1024 * 1024 * 4),
gt(0),
seed(std::chrono::steady_clock::now().time_since_epoch().count()),
disassemble(false),
pc_sampling(false) {}
int device;
bool no_gpu;
bool hip_memset;
size_t rands_len;
uint64_t gt;
uint64_t seed;
bool disassemble;
bool pc_sampling;
};
#endif // SAMPLES_PCSAMPLER_CODE_PRINTING_SAMPLE_PROGRAM_OPTIONS_HPP_