Skip to content

Commit 638ca35

Browse files
edgargabrieledgargabriel
authored
Merge pull request #10 from edgargabriel/topic/ipc-context
topic/ipc context [ROCm/rocshmem commit: d21d5aa]
2 parents 512c114 + 7e446cf commit 638ca35

7 files changed

Lines changed: 729 additions & 0 deletions

File tree

‎projects/rocshmem/src/ipc/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -27,4 +27,5 @@ target_sources(
2727
${PROJECT_NAME}
2828
PRIVATE
2929
ipc_policy.cpp
30+
context_ipc.cpp
3031
)

‎projects/rocshmem/src/ipc/context_ipc.cpp‎

Lines changed: 129 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,129 @@
1+
/******************************************************************************
2+
* Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
3+
*
4+
* Permission is hereby granted, free of charge, to any person obtaining a copy
5+
* of this software and associated documentation files (the "Software"), to
6+
* deal in the Software without restriction, including without limitation the
7+
* rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
8+
* sell copies of the Software, and to permit persons to whom the Software is
9+
* furnished to do so, subject to the following conditions:
10+
*
11+
* The above copyright notice and this permission notice shall be included in
12+
* all copies or substantial portions of the Software.
13+
*
14+
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
15+
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
16+
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
17+
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
18+
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
19+
* FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
20+
* IN THE SOFTWARE.
21+
*****************************************************************************/
22+
23+
#include "context_ipc.hpp"
24+
#include "context_ipc_tmpl_device.hpp"
25+
26+
#include <hip/hip_runtime.h>
27+
#include <hip/amd_detail/amd_device_functions.h>
28+
#include <unistd.h>
29+
30+
#include <cstdio>
31+
#include <cstdlib>
32+
33+
#include "config.h" // NOLINT(build/include_subdir)
34+
#include "roc_shmem/roc_shmem.hpp"
35+
36+
namespace rocshmem {
37+
38+
__host__ IPCContext::IPCContext(Backend *b)
39+
: Context(b, false) {
40+
}
41+
42+
__device__ void IPCContext::threadfence_system() {
43+
}
44+
45+
__device__ void IPCContext::ctx_create() {
46+
}
47+
48+
__device__ void IPCContext::ctx_destroy(){
49+
}
50+
51+
__device__ void IPCContext::putmem(void *dest, const void *source, size_t nelems,
52+
int pe) {
53+
}
54+
55+
__device__ void IPCContext::getmem(void *dest, const void *source, size_t nelems,
56+
int pe) {
57+
}
58+
59+
__device__ void IPCContext::putmem_nbi(void *dest, const void *source,
60+
size_t nelems, int pe) {
61+
}
62+
63+
__device__ void IPCContext::getmem_nbi(void *dest, const void *source,
64+
size_t nelems, int pe) {
65+
}
66+
67+
__device__ void IPCContext::fence() {
68+
}
69+
70+
__device__ void IPCContext::fence(int pe) {
71+
}
72+
73+
__device__ void IPCContext::quiet() {
74+
}
75+
76+
__device__ void *IPCContext::shmem_ptr(const void *dest, int pe) {
77+
void *ret = nullptr;
78+
return ret;
79+
}
80+
81+
__device__ void IPCContext::barrier_all() {
82+
__syncthreads();
83+
}
84+
85+
__device__ void IPCContext::sync_all() {
86+
__syncthreads();
87+
}
88+
89+
__device__ void IPCContext::sync(roc_shmem_team_t team) {
90+
__syncthreads();
91+
}
92+
93+
__device__ void IPCContext::putmem_wg(void *dest, const void *source,
94+
size_t nelems, int pe) {
95+
__syncthreads();
96+
}
97+
98+
__device__ void IPCContext::getmem_wg(void *dest, const void *source,
99+
size_t nelems, int pe) {
100+
__syncthreads();
101+
}
102+
103+
__device__ void IPCContext::putmem_nbi_wg(void *dest, const void *source,
104+
size_t nelems, int pe) {
105+
__syncthreads();
106+
}
107+
108+
__device__ void IPCContext::getmem_nbi_wg(void *dest, const void *source,
109+
size_t nelems, int pe) {
110+
__syncthreads();
111+
}
112+
113+
__device__ void IPCContext::putmem_wave(void *dest, const void *source,
114+
size_t nelems, int pe) {
115+
}
116+
117+
__device__ void IPCContext::getmem_wave(void *dest, const void *source,
118+
size_t nelems, int pe) {
119+
}
120+
121+
__device__ void IPCContext::putmem_nbi_wave(void *dest, const void *source,
122+
size_t nelems, int pe) {
123+
}
124+
125+
__device__ void IPCContext::getmem_nbi_wave(void *dest, const void *source,
126+
size_t nelems, int pe) {
127+
}
128+
129+
} // namespace rocshmem

‎projects/rocshmem/src/ipc/context_ipc.hpp‎

Lines changed: 241 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,241 @@
1+
/******************************************************************************
2+
* Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
3+
*
4+
* Permission is hereby granted, free of charge, to any person obtaining a copy
5+
* of this software and associated documentation files (the "Software"), to
6+
* deal in the Software without restriction, including without limitation the
7+
* rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
8+
* sell copies of the Software, and to permit persons to whom the Software is
9+
* furnished to do so, subject to the following conditions:
10+
*
11+
* The above copyright notice and this permission notice shall be included in
12+
* all copies or substantial portions of the Software.
13+
*
14+
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
15+
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
16+
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
17+
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
18+
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
19+
* FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
20+
* IN THE SOFTWARE.
21+
*****************************************************************************/
22+
23+
#ifndef LIBRARY_SRC_IPC_CONTEXT_DEVICE_HPP_
24+
#define LIBRARY_SRC_IPC_CONTEXT_DEVICE_HPP_
25+
26+
#include "../context.hpp"
27+
28+
namespace rocshmem {
29+
30+
class IPCContext : public Context {
31+
public:
32+
__host__ IPCContext(Backend *b);
33+
34+
__device__ IPCContext(Backend *b);
35+
36+
__device__ void threadfence_system();
37+
38+
__device__ void ctx_create();
39+
40+
__device__ void ctx_destroy();
41+
42+
__device__ void putmem(void *dest, const void *source, size_t nelems, int pe);
43+
44+
__device__ void getmem(void *dest, const void *source, size_t nelems, int pe);
45+
46+
__device__ void putmem_nbi(void *dest, const void *source, size_t nelems,
47+
int pe);
48+
49+
__device__ void getmem_nbi(void *dest, const void *source, size_t size,
50+
int pe);
51+
52+
__device__ void fence();
53+
54+
__device__ void fence(int pe);
55+
56+
__device__ void quiet();
57+
58+
__device__ void *shmem_ptr(const void *dest, int pe);
59+
60+
__device__ void barrier_all();
61+
62+
__device__ void sync_all();
63+
64+
__device__ void sync(roc_shmem_team_t team);
65+
66+
template <typename T>
67+
__device__ void p(T *dest, T value, int pe);
68+
69+
template <typename T>
70+
__device__ void put(T *dest, const T *source, size_t nelems, int pe);
71+
72+
template <typename T>
73+
__device__ void put_nbi(T *dest, const T *source, size_t nelems, int pe);
74+
75+
template <typename T>
76+
__device__ T g(const T *source, int pe);
77+
78+
template <typename T>
79+
__device__ void get(T *dest, const T *source, size_t nelems, int pe);
80+
81+
template <typename T>
82+
__device__ void get_nbi(T *dest, const T *source, size_t nelems, int pe);
83+
84+
// Atomic operations
85+
template <typename T>
86+
__device__ void amo_add(void *dst, T value, int pe);
87+
88+
template <typename T>
89+
__device__ void amo_set(void *dst, T value, int pe);
90+
91+
template <typename T>
92+
__device__ T amo_swap(void *dst, T value, int pe);
93+
94+
template <typename T>
95+
__device__ T amo_fetch_and(void *dst, T value, int pe);
96+
97+
template <typename T>
98+
__device__ void amo_and(void *dst, T value, int pe);
99+
100+
template <typename T>
101+
__device__ T amo_fetch_or(void *dst, T value, int pe);
102+
103+
template <typename T>
104+
__device__ void amo_or(void *dst, T value, int pe);
105+
106+
template <typename T>
107+
__device__ T amo_fetch_xor(void *dst, T value, int pe);
108+
109+
template <typename T>
110+
__device__ void amo_xor(void *dst, T value, int pe);
111+
112+
template <typename T>
113+
__device__ void amo_cas(void *dst, T value, T cond, int pe);
114+
115+
template <typename T>
116+
__device__ T amo_fetch_add(void *dst, T value, int pe);
117+
118+
template <typename T>
119+
__device__ T amo_fetch_cas(void *dst, T value, T cond, int pe);
120+
121+
// Collectives
122+
template <typename T, ROC_SHMEM_OP Op>
123+
__device__ void to_all(T *dest, const T *source, int nreduce, int PE_start,
124+
int logPE_stride, int PE_size, T *pWrk,
125+
long *pSync); // NOLINT(runtime/int)
126+
127+
template <typename T, ROC_SHMEM_OP Op>
128+
__device__ void to_all(roc_shmem_team_t team, T *dest, const T *source,
129+
int nreduce);
130+
131+
template <typename T>
132+
__device__ void broadcast(roc_shmem_team_t team, T *dest, const T *source,
133+
int nelems, int pe_root);
134+
135+
template <typename T>
136+
__device__ void broadcast(T *dest, const T *source, int nelems, int pe_root,
137+
int pe_start, int log_pe_stride, int pe_size,
138+
long *p_sync); // NOLINT(runtime/int)
139+
template <typename T>
140+
__device__ void alltoall(roc_shmem_team_t team, T *dest, const T *source,
141+
int nelems);
142+
template <typename T>
143+
__device__ void fcollect(roc_shmem_team_t team, T *dest, const T *source,
144+
int nelems);
145+
146+
147+
// Block/wave functions
148+
__device__ void putmem_wg(void *dest, const void *source, size_t nelems,
149+
int pe);
150+
151+
__device__ void getmem_wg(void *dest, const void *source, size_t nelems,
152+
int pe);
153+
154+
__device__ void putmem_nbi_wg(void *dest, const void *source, size_t nelems,
155+
int pe);
156+
157+
__device__ void getmem_nbi_wg(void *dest, const void *source, size_t size,
158+
int pe);
159+
160+
__device__ void putmem_wave(void *dest, const void *source, size_t nelems,
161+
int pe);
162+
163+
__device__ void getmem_wave(void *dest, const void *source, size_t nelems,
164+
int pe);
165+
166+
__device__ void putmem_nbi_wave(void *dest, const void *source, size_t nelems,
167+
int pe);
168+
169+
__device__ void getmem_nbi_wave(void *dest, const void *source, size_t size,
170+
int pe);
171+
172+
template <typename T>
173+
__device__ void put_wg(T *dest, const T *source, size_t nelems, int pe);
174+
175+
template <typename T>
176+
__device__ void put_nbi_wg(T *dest, const T *source, size_t nelems, int pe);
177+
178+
template <typename T>
179+
__device__ void put_wave(T *dest, const T *source, size_t nelems, int pe);
180+
181+
template <typename T>
182+
__device__ void put_nbi_wave(T *dest, const T *source, size_t nelems, int pe);
183+
184+
template <typename T>
185+
__device__ void get_wg(T *dest, const T *source, size_t nelems, int pe);
186+
187+
template <typename T>
188+
__device__ void get_nbi_wg(T *dest, const T *source, size_t nelems, int pe);
189+
190+
191+
template <typename T>
192+
__device__ void get_wave(T *dest, const T *source, size_t nelems, int pe);
193+
194+
template <typename T>
195+
__device__ void get_nbi_wave(T *dest, const T *source, size_t nelems, int pe);
196+
197+
// Wait / Test functions
198+
template <typename T>
199+
__device__ void wait_until(T* ptr, roc_shmem_cmps cmp, T val);
200+
201+
template <typename T>
202+
__device__ void wait_until_all(T* ptr, size_t nelems,
203+
const int *status,
204+
roc_shmem_cmps cmp, T val);
205+
206+
template <typename T>
207+
__device__ size_t wait_until_any(T* ptr, size_t nelems,
208+
const int *status,
209+
roc_shmem_cmps cmp, T val);
210+
211+
template <typename T>
212+
__device__ size_t wait_until_some(T* ptr, size_t nelems,
213+
size_t* indices,
214+
const int *status,
215+
roc_shmem_cmps cmp, T val);
216+
217+
template <typename T>
218+
__device__ void wait_until_all_vector(T* ptr, size_t nelems,
219+
const int *status,
220+
roc_shmem_cmps cmp, T* vals);
221+
222+
template <typename T>
223+
__device__ size_t wait_until_any_vector(T* ptr, size_t nelems,
224+
const int *status,
225+
roc_shmem_cmps cmp, T* vals);
226+
template <typename T>
227+
__device__ size_t wait_until_some_vector(T* ptr, size_t nelems,
228+
size_t* indices,
229+
const int *status,
230+
roc_shmem_cmps cmp, T* vals);
231+
232+
template <typename T>
233+
__device__ int test(T* ptr, roc_shmem_cmps cmp, T val);
234+
235+
private:
236+
237+
};
238+
239+
} // namespace rocshmem
240+
241+
#endif // LIBRARY_SRC_GPU_IB_CONTEXT_IB_DEVICE_HPP_

0 commit comments

Comments
 (0)