[Openmp-commits] [openmp] 7e9351b - [libomptarget][amdgpu][nfc] Remove dead code from amdgpu plugin
Jon Chesterfield via Openmp-commits
openmp-commits at lists.llvm.org
Thu May 6 15:16:44 PDT 2021
Author: Jon Chesterfield
Date: 2021-05-06T23:16:32+01:00
New Revision: 7e9351b9dee225b9ab12ee9bbfc9f8b96ddd1a1d
URL: https://github.com/llvm/llvm-project/commit/7e9351b9dee225b9ab12ee9bbfc9f8b96ddd1a1d
DIFF: https://github.com/llvm/llvm-project/commit/7e9351b9dee225b9ab12ee9bbfc9f8b96ddd1a1d.diff
LOG: [libomptarget][amdgpu][nfc] Remove dead code from amdgpu plugin
[libomptarget][amdgpu][nfc] Remove dead code from amdgpu plugin
Drops an enum that was identical to a HSA one, localises some functions where
they were only called from one TU. Covers everything internalize + adce can
identify as dead, except for msgpack::dump which is useful when debugging.
Reviewed By: jdoerfert
Differential Revision: https://reviews.llvm.org/D102014
Added:
Modified:
openmp/libomptarget/plugins/amdgpu/impl/atmi.h
openmp/libomptarget/plugins/amdgpu/impl/data.cpp
openmp/libomptarget/plugins/amdgpu/impl/internal.h
openmp/libomptarget/plugins/amdgpu/impl/system.cpp
openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
Removed:
################################################################################
diff --git a/openmp/libomptarget/plugins/amdgpu/impl/atmi.h b/openmp/libomptarget/plugins/amdgpu/impl/atmi.h
index 35e61216ce780..b636d6b9aca13 100644
--- a/openmp/libomptarget/plugins/amdgpu/impl/atmi.h
+++ b/openmp/libomptarget/plugins/amdgpu/impl/atmi.h
@@ -51,25 +51,6 @@ typedef enum atmi_memtype_s {
ATMI_MEMTYPE_ANY
} atmi_memtype_t;
-/**
- * @brief ATMI Memory Fences for Tasks.
- */
-typedef enum atmi_task_fence_scope_s {
- /**
- * No memory fence applied; external fences have to be applied around the task
- * launch/completion.
- */
- ATMI_FENCE_SCOPE_NONE = 0,
- /**
- * The fence is applied to the device.
- */
- ATMI_FENCE_SCOPE_DEVICE = 1,
- /**
- * The fence is applied to the entire system.
- */
- ATMI_FENCE_SCOPE_SYSTEM = 2
-} atmi_task_fence_scope_t;
-
/** @} */
/** \defgroup common Common ATMI Structures
@@ -164,20 +145,9 @@ typedef struct atmi_machine_s {
// Below are some helper macros that can be used to setup
// some of the ATMI data structures.
-#define ATMI_PLACE_CPU(node, cpu_id) \
- { .node_id = node, .type = ATMI_DEVTYPE_CPU, .device_id = cpu_id }
#define ATMI_PLACE_GPU(node, gpu_id) \
{ .node_id = node, .type = ATMI_DEVTYPE_GPU, .device_id = gpu_id }
-#define ATMI_MEM_PLACE_CPU(node, cpu_id) \
- { \
- .node_id = node, .dev_type = ATMI_DEVTYPE_CPU, .dev_id = cpu_id, \
- .mem_id = -1 \
- }
-#define ATMI_MEM_PLACE_GPU(node, gpu_id) \
- { \
- .node_id = node, .dev_type = ATMI_DEVTYPE_GPU, .dev_id = gpu_id, \
- .mem_id = -1 \
- }
+
#define ATMI_MEM_PLACE_CPU_MEM(node, cpu_id, cpu_mem_id) \
{ \
.node_id = node, .dev_type = ATMI_DEVTYPE_CPU, .dev_id = cpu_id, \
diff --git a/openmp/libomptarget/plugins/amdgpu/impl/data.cpp b/openmp/libomptarget/plugins/amdgpu/impl/data.cpp
index 0d98d5c51ce1e..686ecec1efaab 100644
--- a/openmp/libomptarget/plugins/amdgpu/impl/data.cpp
+++ b/openmp/libomptarget/plugins/amdgpu/impl/data.cpp
@@ -22,17 +22,7 @@ extern ATLMachine g_atl_machine;
namespace core {
void allow_access_to_all_gpu_agents(void *ptr);
-const char *getPlaceStr(atmi_devtype_t type) {
- switch (type) {
- case ATMI_DEVTYPE_CPU:
- return "CPU";
- case ATMI_DEVTYPE_GPU:
- return "GPU";
- default:
- return NULL;
- }
-}
-
+namespace {
ATLProcessor &get_processor_by_mem_place(atmi_mem_place_t place) {
int dev_id = place.dev_id;
switch (place.dev_type) {
@@ -47,6 +37,7 @@ hsa_amd_memory_pool_t get_memory_pool_by_mem_place(atmi_mem_place_t place) {
ATLProcessor &proc = get_processor_by_mem_place(place);
return get_memory_pool(proc, place.mem_id);
}
+} // namespace
void register_allocation(void *ptr, size_t size, atmi_mem_place_t place) {
if (place.dev_type == ATMI_DEVTYPE_CPU)
diff --git a/openmp/libomptarget/plugins/amdgpu/impl/internal.h b/openmp/libomptarget/plugins/amdgpu/impl/internal.h
index 8ca66a9d478e0..d82216cbfc304 100644
--- a/openmp/libomptarget/plugins/amdgpu/impl/internal.h
+++ b/openmp/libomptarget/plugins/amdgpu/impl/internal.h
@@ -222,17 +222,11 @@ template <typename T> inline T *alignUp(T *value, size_t alignment) {
extern void register_allocation(void *addr, size_t size,
atmi_mem_place_t place);
-extern hsa_amd_memory_pool_t
-get_memory_pool_by_mem_place(atmi_mem_place_t place);
+
extern bool atl_is_atmi_initialized();
bool handle_group_signal(hsa_signal_value_t value, void *arg);
-void packet_store_release(uint32_t *packet, uint16_t header, uint16_t rest);
-uint16_t
-create_header(hsa_packet_type_t type, int barrier,
- atmi_task_fence_scope_t acq_fence = ATMI_FENCE_SCOPE_SYSTEM,
- atmi_task_fence_scope_t rel_fence = ATMI_FENCE_SCOPE_SYSTEM);
void allow_access_to_all_gpu_agents(void *ptr);
} // namespace core
diff --git a/openmp/libomptarget/plugins/amdgpu/impl/system.cpp b/openmp/libomptarget/plugins/amdgpu/impl/system.cpp
index d6cde1f699c23..f25020d7697b6 100644
--- a/openmp/libomptarget/plugins/amdgpu/impl/system.cpp
+++ b/openmp/libomptarget/plugins/amdgpu/impl/system.cpp
@@ -172,12 +172,12 @@ atmi_machine_t *Runtime::GetMachineInfo() {
return &g_atmi_machine;
}
-void atl_set_atmi_initialized() {
+static void atl_set_atmi_initialized() {
// FIXME: thread safe? locks?
g_atmi_initialized = true;
}
-void atl_reset_atmi_initialized() {
+static void atl_reset_atmi_initialized() {
// FIXME: thread safe? locks?
g_atmi_initialized = false;
}
@@ -233,7 +233,7 @@ atmi_status_t Runtime::Finalize() {
return ATMI_STATUS_SUCCESS;
}
-void atmi_init_context_structs() {
+static void atmi_init_context_structs() {
atlc_p = &atlc;
atlc.struct_initialized = true; /* This only gets called one time */
atlc.g_hsa_initialized = false;
@@ -609,7 +609,7 @@ atmi_status_t atl_init_gpu_context() {
return ATMI_STATUS_SUCCESS;
}
-bool isImplicit(KernelArgMD::ValueKind value_kind) {
+static bool isImplicit(KernelArgMD::ValueKind value_kind) {
switch (value_kind) {
case KernelArgMD::ValueKind::HiddenGlobalOffsetX:
case KernelArgMD::ValueKind::HiddenGlobalOffsetY:
diff --git a/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp b/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
index a6b426dc05579..073b309f5d57c 100644
--- a/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
+++ b/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
@@ -72,9 +72,6 @@ hostrpc_assign_buffer(hsa_agent_t, hsa_queue_t *, uint32_t device_id) {
int print_kernel_trace;
-// Size of the target call stack struture
-uint32_t TgtStackItemSize = 0;
-
#undef check // Drop definition from internal.h
#ifdef OMPTARGET_DEBUG
#define check(msg, status) \
@@ -275,21 +272,18 @@ static void callbackQueue(hsa_status_t status, hsa_queue_t *source,
}
namespace core {
+namespace {
void packet_store_release(uint32_t *packet, uint16_t header, uint16_t rest) {
__atomic_store_n(packet, header | (rest << 16), __ATOMIC_RELEASE);
}
-uint16_t create_header(hsa_packet_type_t type, int barrier,
- atmi_task_fence_scope_t acq_fence,
- atmi_task_fence_scope_t rel_fence) {
- uint16_t header = type << HSA_PACKET_HEADER_TYPE;
- header |= barrier << HSA_PACKET_HEADER_BARRIER;
- header |= (hsa_fence_scope_t) static_cast<int>(
- acq_fence << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE);
- header |= (hsa_fence_scope_t) static_cast<int>(
- rel_fence << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
+uint16_t create_header() {
+ uint16_t header = HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE;
+ header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE;
+ header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE;
return header;
}
+} // namespace
} // namespace core
/// Class containing all the device information
@@ -1904,11 +1898,8 @@ int32_t __tgt_rtl_run_target_team_region_locked(
hsa_signal_store_relaxed(packet->completion_signal, 1);
}
- core::packet_store_release(
- reinterpret_cast<uint32_t *>(packet),
- core::create_header(HSA_PACKET_TYPE_KERNEL_DISPATCH, 0,
- ATMI_FENCE_SCOPE_SYSTEM, ATMI_FENCE_SCOPE_SYSTEM),
- packet->setup);
+ core::packet_store_release(reinterpret_cast<uint32_t *>(packet),
+ core::create_header(), packet->setup);
hsa_signal_store_relaxed(queue->doorbell_signal, packet_id);
More information about the Openmp-commits
mailing list