Skip to content

Commit ddbdca2

Browse files
adeljo-amdadeljo-amd
committed
Add SOMA IPC examples
1 parent 415a48a commit ddbdca2

13 files changed

Lines changed: 831 additions & 0 deletions

File tree

‎HIP-Doc/Programming-Guide/Using-HIP-Runtime-API/Memory-Management/SOMA/CMakeLists.txt‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -29,6 +29,8 @@ set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/bin/${folder_bin})
2929
if(CMAKE_SYSTEM_NAME STREQUAL "Windows")
3030
message(STATUS "The memory pool examples are available only on Linux.")
3131
else()
32+
add_subdirectory(ipc_memory_pool_device_pointer)
33+
add_subdirectory(ipc_memory_pool_shareable_handle)
3234
add_subdirectory(memory_pool)
3335
add_subdirectory(memory_pool_resource_usage_statistics)
3436
add_subdirectory(memory_pool_threshold)

‎HIP-Doc/Programming-Guide/Using-HIP-Runtime-API/Memory-Management/SOMA/Makefile‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -21,6 +21,8 @@
2121
# SOFTWARE.
2222

2323
EXAMPLES := \
24+
ipc_memory_pool_device_pointer \
25+
ipc_memory_pool_shareable_handle \
2426
memory_pool \
2527
memory_pool_resource_usage_statistics \
2628
memory_pool_threshold \

‎HIP-Doc/Programming-Guide/Using-HIP-Runtime-API/Memory-Management/SOMA/ipc_memory_pool_device_pointer/.gitignore‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1 @@
1+
hip_ipc_memory_pool_device_pointer

‎HIP-Doc/Programming-Guide/Using-HIP-Runtime-API/Memory-Management/SOMA/ipc_memory_pool_device_pointer/CMakeLists.txt‎

Lines changed: 49 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,49 @@
1+
# MIT License
2+
#
3+
# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved.
4+
#
5+
# Permission is hereby granted, free of charge, to any person obtaining a copy
6+
# of this software and associated documentation files (the "Software"), to deal
7+
# in the Software without restriction, including without limitation the rights
8+
# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
9+
# copies of the Software, and to permit persons to whom the Software is
10+
# furnished to do so, subject to the following conditions:
11+
#
12+
# The above copyright notice and this permission notice shall be included in all
13+
# copies or substantial portions of the Software.
14+
#
15+
# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
16+
# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
17+
# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
18+
# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
19+
# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
20+
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
21+
# SOFTWARE.
22+
23+
set(example_name hip_ipc_memory_pool_device_pointer)
24+
25+
cmake_minimum_required(VERSION 3.21 FATAL_ERROR)
26+
project(${example_name} LANGUAGES CXX)
27+
28+
if(CMAKE_SYSTEM_NAME MATCHES "Windows")
29+
message(FATAL_ERROR "The IPC memory pool API is available only on Linux with the amdgpu-dkms driver.")
30+
endif()
31+
32+
include("${CMAKE_CURRENT_LIST_DIR}/../../../../../../Common/HipPlatform.cmake")
33+
select_gpu_language()
34+
enable_language(${ROCM_EXAMPLES_GPU_LANGUAGE})
35+
select_hip_platform()
36+
37+
include("${CMAKE_CURRENT_LIST_DIR}/../../../../../../Common/ROCmPath.cmake")
38+
39+
add_executable(${example_name} main.hip)
40+
set_source_files_properties(main.hip PROPERTIES LANGUAGE ${ROCM_EXAMPLES_GPU_LANGUAGE})
41+
set_target_properties(${example_name} PROPERTIES $<$<COMPILE_LANGUAGE:CUDA>:CUDA_EXTENSIONS OFF>
42+
$<$<COMPILE_LANGUAGE:HIP>:HIP_EXTENSIONS OFF>)
43+
target_compile_features(${example_name} PRIVATE cuda_std_17 hip_std_17)
44+
target_include_directories(${example_name} PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:${ROCM_PATH}/include>)
45+
46+
# Note: no ctest entry. This example requires two cooperating processes and
47+
# the amdgpu-dkms driver. See README.md for manual run instructions.
48+
49+
install(TARGETS ${example_name})

‎HIP-Doc/Programming-Guide/Using-HIP-Runtime-API/Memory-Management/SOMA/ipc_memory_pool_device_pointer/Makefile‎

Lines changed: 62 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,62 @@
1+
# MIT License
2+
#
3+
# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved.
4+
#
5+
# Permission is hereby granted, free of charge, to any person obtaining a copy
6+
# of this software and associated documentation files (the "Software"), to deal
7+
# in the Software without restriction, including without limitation the rights
8+
# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
9+
# copies of the Software, and to permit persons to whom the Software is
10+
# furnished to do so, subject to the following conditions:
11+
#
12+
# The above copyright notice and this permission notice shall be included in all
13+
# copies or substantial portions of the Software.
14+
#
15+
# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
16+
# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
17+
# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
18+
# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
19+
# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
20+
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
21+
# SOFTWARE.
22+
23+
EXAMPLE := hip_ipc_memory_pool_device_pointer
24+
GPU_RUNTIME ?= HIP
25+
26+
# HIP variables
27+
ROCM_PATH ?= /opt/rocm
28+
HIP_INCLUDE_DIR := $(ROCM_PATH)/include
29+
30+
HIPCXX ?= $(ROCM_PATH)/bin/hipcc
31+
32+
# Common variables and flags
33+
CXX_STD := c++17
34+
ICXXFLAGS := -std=$(CXX_STD)
35+
ICPPFLAGS :=
36+
ILDFLAGS :=
37+
ILDLIBS :=
38+
39+
ifeq ($(GPU_RUNTIME), CUDA)
40+
ICXXFLAGS += -x cu
41+
ICPPFLAGS += -isystem $(HIP_INCLUDE_DIR)
42+
else ifeq ($(GPU_RUNTIME), HIP)
43+
CXXFLAGS ?= -Wall -Wextra
44+
else
45+
$(error GPU_RUNTIME is set to "$(GPU_RUNTIME)". GPU_RUNTIME must be either CUDA or HIP)
46+
endif
47+
48+
ICXXFLAGS += $(CXXFLAGS)
49+
ICPPFLAGS += $(CPPFLAGS)
50+
ILDFLAGS += $(LDFLAGS)
51+
ILDLIBS += $(LDLIBS)
52+
53+
$(EXAMPLE): main.hip
54+
$(HIPCXX) $(ICXXFLAGS) $(ICPPFLAGS) $(ILDFLAGS) -o $@ $< $(ILDLIBS)
55+
56+
test:
57+
@echo "Note: $(EXAMPLE) requires two cooperating processes and the amdgpu-dkms driver. See README.md for instructions."
58+
59+
clean:
60+
$(RM) $(EXAMPLE)
61+
62+
.PHONY: clean test

‎HIP-Doc/Programming-Guide/Using-HIP-Runtime-API/Memory-Management/SOMA/ipc_memory_pool_device_pointer/README.md‎

Lines changed: 80 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,80 @@
1+
# HIP-Doc IPC Memory Pool Device Pointer Example
2+
3+
## Description
4+
5+
This example demonstrates how to share a stream ordered memory allocation between
6+
two processes using the device pointer IPC mechanism. The exporter allocates memory
7+
from an IPC-capable pool, fills it using a GPU kernel, and exports an opaque handle
8+
(`hipMemPoolPtrExportData`) that the importer uses to access the same allocation.
9+
10+
For more information, see
11+
[HIP documentation](https://rocm.docs.amd.com/projects/HIP/en/latest/how-to/hip_runtime_api/memory_management/stream_ordered_allocator.html#device-pointer).
12+
13+
### Prerequisites
14+
15+
The IPC memory pool API requires the `amdgpu-dkms` driver and is available only
16+
on Linux. This example does not run on Windows.
17+
18+
The example consists of two cooperating processes. Run them in two terminals,
19+
starting the exporter first:
20+
21+
```bash
22+
# Terminal 1
23+
./hip_ipc_memory_pool_device_pointer export
24+
25+
# Terminal 2
26+
./hip_ipc_memory_pool_device_pointer import
27+
```
28+
29+
### Application flow
30+
31+
**Exporter:**
32+
33+
1. A memory pool is created with `hipMemHandleTypePosixFileDescriptor` to enable IPC.
34+
2. A HIP stream is created.
35+
3. Memory is allocated from the pool using `hipMallocFromPoolAsync`.
36+
4. A kernel fills the allocation with computed values.
37+
5. The stream is synchronized.
38+
6. The device pointer is exported with `hipMemPoolExportPointer`, producing a serializable `hipMemPoolPtrExportData` struct.
39+
7. The struct is written to a named pipe for the importer to read.
40+
8. The exporter waits for the importer to signal completion before freeing memory.
41+
9. Memory, pool, and stream are destroyed.
42+
43+
**Importer:**
44+
45+
1. The `hipMemPoolPtrExportData` struct is read from the named pipe.
46+
2. A memory pool is created with matching IPC-capable properties.
47+
3. The device pointer is imported with `hipMemPoolImportPointer`.
48+
4. The imported pointer is copied to host memory and verified.
49+
5. The importer signals the exporter that it has finished.
50+
6. The imported pointer and pool are destroyed.
51+
52+
## Key APIs and Concepts
53+
54+
* `hipMemPoolCreate` creates a memory pool. Setting `handleTypes` to `hipMemHandleTypePosixFileDescriptor` makes the pool IPC-capable.
55+
* `hipMallocFromPoolAsync` allocates memory from a specific pool with stream-ordered semantics.
56+
* `hipMemPoolExportPointer` exports an allocation as a serializable `hipMemPoolPtrExportData` struct that can be written to any IPC channel and read by another process.
57+
* `hipMemPoolImportPointer` imports a previously exported allocation into a local pool handle, giving the importing process access to the same GPU memory.
58+
* `hipFreeAsync` returns memory to the pool with stream-ordered semantics.
59+
* `hipMemPoolDestroy` destroys a memory pool.
60+
61+
## Demonstrated API Calls
62+
63+
### HIP Runtime
64+
65+
#### Device symbols
66+
67+
* `threadIdx`, `blockIdx`, `blockDim`
68+
69+
#### Host symbols
70+
71+
* `hipFreeAsync`
72+
* `hipMallocFromPoolAsync`
73+
* `hipMemPoolCreate`
74+
* `hipMemPoolDestroy`
75+
* `hipMemPoolExportPointer`
76+
* `hipMemPoolImportPointer`
77+
* `hipMemcpy`
78+
* `hipStreamCreate`
79+
* `hipStreamDestroy`
80+
* `hipStreamSynchronize`

‎HIP-Doc/Programming-Guide/Using-HIP-Runtime-API/Memory-Management/SOMA/ipc_memory_pool_device_pointer/main.hip‎

Lines changed: 183 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,183 @@
1+
// MIT License
2+
//
3+
// Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved.
4+
//
5+
// Permission is hereby granted, free of charge, to any person obtaining a copy
6+
// of this software and associated documentation files (the "Software"), to deal
7+
// in the Software without restriction, including without limitation the rights
8+
// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
9+
// copies of the Software, and to permit persons to whom the Software is
10+
// furnished to do so, subject to the following conditions:
11+
//
12+
// The above copyright notice and this permission notice shall be included in all
13+
// copies or substantial portions of the Software.
14+
//
15+
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
16+
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
17+
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
18+
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
19+
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
20+
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
21+
// SOFTWARE.
22+
23+
// [sphinx-exporter-start]
24+
#include <hip/hip_runtime.h>
25+
26+
#include <cstddef>
27+
#include <cstdlib>
28+
#include <fstream>
29+
#include <iostream>
30+
#include <sys/stat.h>
31+
#include <unistd.h>
32+
33+
#define HIP_CHECK(expression) \
34+
{ \
35+
const hipError_t status = expression; \
36+
if (status != hipSuccess) \
37+
{ \
38+
std::cerr << "HIP error " << status \
39+
<< ": " << hipGetErrorString(status) \
40+
<< " at " << __FILE__ << ":" \
41+
<< __LINE__ << std::endl; \
42+
std::exit(EXIT_FAILURE); \
43+
} \
44+
}
45+
46+
constexpr std::size_t numElements = 1024;
47+
constexpr const char* fifoPath = "/tmp/hip_ipc_devptr_fifo";
48+
constexpr const char* doneFifoPath = "/tmp/hip_ipc_devptr_done_fifo";
49+
50+
__global__ void fillKernel(int* data, std::size_t n)
51+
{
52+
int tid = threadIdx.x + blockIdx.x * blockDim.x;
53+
if (tid < static_cast<int>(n))
54+
data[tid] = tid * 2;
55+
}
56+
57+
void runExporter()
58+
{
59+
// Allocate memory from an IPC-capable pool.
60+
hipMemPoolProps poolProps = {};
61+
poolProps.allocType = hipMemAllocationTypePinned;
62+
poolProps.handleTypes = hipMemHandleTypePosixFileDescriptor;
63+
poolProps.location.type = hipMemLocationTypeDevice;
64+
poolProps.location.id = 0;
65+
66+
hipMemPool_t memPool;
67+
HIP_CHECK(hipMemPoolCreate(&memPool, &poolProps));
68+
69+
hipStream_t stream;
70+
HIP_CHECK(hipStreamCreate(&stream));
71+
72+
int* devPtr = nullptr;
73+
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&devPtr),
74+
numElements * sizeof(*devPtr),
75+
memPool,
76+
stream));
77+
78+
dim3 blockSize(256);
79+
dim3 gridSize((numElements + blockSize.x - 1) / blockSize.x);
80+
fillKernel<<<gridSize, blockSize, 0, stream>>>(devPtr, numElements);
81+
HIP_CHECK(hipStreamSynchronize(stream));
82+
83+
// Export the memory pool pointer.
84+
hipMemPoolPtrExportData exportData;
85+
HIP_CHECK(hipMemPoolExportPointer(&exportData, devPtr));
86+
87+
// Write the exported data to the named pipe.
88+
mkfifo(fifoPath, 0666);
89+
std::ofstream fifoStream(fifoPath, std::ios::out | std::ios::binary);
90+
fifoStream.write(reinterpret_cast<char*>(&exportData), sizeof(exportData));
91+
fifoStream.close();
92+
93+
// Wait for the importer to signal completion before freeing.
94+
mkfifo(doneFifoPath, 0666);
95+
char ack;
96+
std::ifstream doneStream(doneFifoPath, std::ios::in | std::ios::binary);
97+
doneStream.read(&ack, 1);
98+
doneStream.close();
99+
100+
HIP_CHECK(hipFreeAsync(devPtr, stream));
101+
HIP_CHECK(hipStreamSynchronize(stream));
102+
HIP_CHECK(hipMemPoolDestroy(memPool));
103+
HIP_CHECK(hipStreamDestroy(stream));
104+
unlink(fifoPath);
105+
unlink(doneFifoPath);
106+
std::cout << "Exporter: done." << std::endl;
107+
}
108+
// [sphinx-exporter-end]
109+
110+
// [sphinx-importer-start]
111+
void runImporter()
112+
{
113+
// Read the exported data from the named pipe.
114+
hipMemPoolPtrExportData importData;
115+
std::ifstream fifoStream(fifoPath, std::ios::in | std::ios::binary);
116+
fifoStream.read(reinterpret_cast<char*>(&importData), sizeof(importData));
117+
fifoStream.close();
118+
119+
// Create a memory pool with matching IPC-capable properties.
120+
hipMemPoolProps poolProps = {};
121+
poolProps.allocType = hipMemAllocationTypePinned;
122+
poolProps.handleTypes = hipMemHandleTypePosixFileDescriptor;
123+
poolProps.location.type = hipMemLocationTypeDevice;
124+
poolProps.location.id = 0;
125+
126+
hipMemPool_t memPool;
127+
HIP_CHECK(hipMemPoolCreate(&memPool, &poolProps));
128+
129+
// Import the device pointer from the exporter's pool.
130+
int* importedDevPtr = nullptr;
131+
HIP_CHECK(hipMemPoolImportPointer(reinterpret_cast<void**>(&importedDevPtr),
132+
memPool,
133+
&importData));
134+
135+
// Verify the values written by the exporter's kernel.
136+
int* hostData = new int[numElements];
137+
HIP_CHECK(hipMemcpy(hostData,
138+
importedDevPtr,
139+
numElements * sizeof(*hostData),
140+
hipMemcpyDeviceToHost));
141+
142+
bool passed = true;
143+
for (std::size_t i = 0; i < numElements; ++i)
144+
{
145+
if (hostData[i] != static_cast<int>(i * 2))
146+
{
147+
std::cerr << "Mismatch at [" << i << "]: expected " << i * 2
148+
<< ", got " << hostData[i] << std::endl;
149+
passed = false;
150+
break;
151+
}
152+
}
153+
delete[] hostData;
154+
155+
// Signal the exporter that this process has finished.
156+
char ack = 1;
157+
std::ofstream doneStream(doneFifoPath, std::ios::out | std::ios::binary);
158+
doneStream.write(&ack, 1);
159+
doneStream.close();
160+
161+
HIP_CHECK(hipFree(importedDevPtr));
162+
HIP_CHECK(hipMemPoolDestroy(memPool));
163+
164+
std::cout << "Importer: " << (passed ? "PASS" : "FAIL") << std::endl;
165+
std::exit(passed ? EXIT_SUCCESS : EXIT_FAILURE);
166+
}
167+
// [sphinx-importer-end]
168+
169+
int main(int argc, char* argv[])
170+
{
171+
if(argc < 2 || (std::string(argv[1]) != "export" && std::string(argv[1]) != "import"))
172+
{
173+
std::cerr << "Usage: " << argv[0] << " export|import" << std::endl;
174+
return EXIT_FAILURE;
175+
}
176+
177+
if(std::string(argv[1]) == "export")
178+
runExporter();
179+
else
180+
runImporter();
181+
182+
return EXIT_SUCCESS;
183+
}

‎HIP-Doc/Programming-Guide/Using-HIP-Runtime-API/Memory-Management/SOMA/ipc_memory_pool_shareable_handle/.gitignore‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1 @@
1+
hip_ipc_memory_pool_shareable_handle

0 commit comments

Comments
 (0)