Actual source code: cupmdevice.cxx
1: #include <petsc/private/cpp/memory.hpp>
3: #include "cupmdevice.hpp"
5: #include <algorithm>
6: #include <csetjmp> // for cuda mpi awareness
7: #include <csignal> // SIGSEGV
8: #include <iterator>
9: #include <type_traits>
11: namespace Petsc
12: {
14: namespace device
15: {
17: namespace cupm
18: {
20: // internal "impls" class for CUPMDevice. Each instance represents a single cupm device
21: template <DeviceType T>
22: class Device<T>::DeviceInternal {
23: const int id_;
24: bool devInitialized_ = false;
25: cupmDeviceProp_t dprop_{}; // cudaDeviceProp appears to be an actual struct, i.e. you can't
26: // initialize it with nullptr or NULL (i've tried)
28: PETSC_CXX_COMPAT_DECL(PetscErrorCode CUPMAwareMPI_(bool *));
30: public:
31: // default constructor
32: explicit constexpr DeviceInternal(int dev) noexcept : id_(dev) { }
34: // gather all relevant information for a particular device, a cupmDeviceProp_t is
35: // usually sufficient here
36: PETSC_NODISCARD PetscErrorCode initialize() noexcept;
37: PETSC_NODISCARD PetscErrorCode configure() noexcept;
38: PETSC_NODISCARD PetscErrorCode view(PetscViewer) const noexcept;
39: PETSC_NODISCARD PetscErrorCode getattribute(PetscDeviceAttribute, void *) const noexcept;
41: PETSC_NODISCARD auto id() const -> decltype(id_) { return id_; }
42: PETSC_NODISCARD auto initialized() const -> decltype(devInitialized_) { return devInitialized_; }
43: PETSC_NODISCARD auto prop() const -> const decltype(dprop_) & { return dprop_; }
44: };
46: // the goal here is simply to get the cupm backend to create its context, not to do any type of
47: // modification of it, or create objects (since these may be affected by subsequent
48: // configuration changes)
49: template <DeviceType T>
50: PetscErrorCode Device<T>::DeviceInternal::initialize() noexcept
51: {
52: if (initialized()) return 0;
53: devInitialized_ = true;
54: // need to do this BEFORE device has been set, although if the user
55: // has already done this then we just ignore it
56: if (cupmSetDeviceFlags(cupmDeviceMapHost) == cupmErrorSetOnActiveProcess) {
57: // reset the error if it was cupmErrorSetOnActiveProcess
58: const auto PETSC_UNUSED unused = cupmGetLastError();
59: } else cupmGetLastError();
60: // cuda 5.0+ will create a context when cupmSetDevice is called
61: if (cupmSetDevice(id()) != cupmErrorDeviceAlreadyInUse) cupmGetLastError();
62: // and in case it doesn't, explicitly call init here
63: cupmInit(0);
64: // where is this variable defined and when is it set? who knows! but it is defined and set
65: // at this point. either way, each device must make this check since I guess MPI might not be
66: // aware of all of them?
67: if (use_gpu_aware_mpi) {
68: bool aware;
70: CUPMAwareMPI_(&aware);
71: // For OpenMPI, we could do a compile time check with
72: // "defined(PETSC_HAVE_OMPI_MAJOR_VERSION) && defined(MPIX_CUDA_AWARE_SUPPORT) &&
73: // MPIX_CUDA_AWARE_SUPPORT" to see if it is CUDA-aware. However, recent versions of IBM
74: // Spectrum MPI (e.g., 10.3.1) on Summit meet above conditions, but one has to use jsrun
75: // --smpiargs=-gpu to really enable GPU-aware MPI. So we do the check at runtime with a
76: // code that works only with GPU-aware MPI.
77: if (PetscUnlikely(!aware)) {
78: (*PetscErrorPrintf)("PETSc is configured with GPU support, but your MPI is not GPU-aware. For better performance, please use a GPU-aware MPI.\n");
79: (*PetscErrorPrintf)("If you do not care, add option -use_gpu_aware_mpi 0. To not see the message again, add the option to your .petscrc, OR add it to the env var PETSC_OPTIONS.\n");
80: (*PetscErrorPrintf)("If you do care, for IBM Spectrum MPI on OLCF Summit, you may need jsrun --smpiargs=-gpu.\n");
81: (*PetscErrorPrintf)("For OpenMPI, you need to configure it --with-cuda (https://www.open-mpi.org/faq/?category=buildcuda)\n");
82: (*PetscErrorPrintf)("For MVAPICH2-GDR, you need to set MV2_USE_CUDA=1 (http://mvapich.cse.ohio-state.edu/userguide/gdr/)\n");
83: (*PetscErrorPrintf)("For Cray-MPICH, you need to set MPICH_RDMA_ENABLED_CUDA=1 (https://www.olcf.ornl.gov/tutorials/gpudirect-mpich-enabled-cuda/)\n");
84: PETSCABORT(PETSC_COMM_SELF, PETSC_ERR_LIB);
85: }
86: }
87: return 0;
88: }
90: template <DeviceType T>
91: PetscErrorCode Device<T>::DeviceInternal::configure() noexcept
92: {
93: PetscAssert(initialized(), PETSC_COMM_SELF, PETSC_ERR_COR, "Device %d being configured before it was initialized", id());
94: // why on EARTH nvidia insists on making otherwise informational states into
95: // fully-fledged error codes is beyond me. Why couldn't a pointer to bool argument have
96: // sufficed?!?!?!
97: if (cupmSetDevice(id_) != cupmErrorDeviceAlreadyInUse) cupmGetLastError();
98: // need to update the device properties
99: cupmGetDeviceProperties(&dprop_, id_);
100: PetscInfo(nullptr, "Configured device %d\n", id_);
101: return 0;
102: }
104: template <DeviceType T>
105: PetscErrorCode Device<T>::DeviceInternal::view(PetscViewer viewer) const noexcept
106: {
107: PetscBool iascii;
109: PetscAssert(initialized(), PETSC_COMM_SELF, PETSC_ERR_COR, "Device %d being viewed before it was initialized or configured", id());
110: // we don't print device-specific info in CI-mode
111: if (PetscUnlikely(PetscCIEnabled)) return 0;
112: PetscObjectTypeCompare(PetscObjectCast(viewer), PETSCVIEWERASCII, &iascii);
113: if (iascii) {
114: MPI_Comm comm;
115: PetscMPIInt rank;
116: PetscViewer sviewer;
118: PetscObjectGetComm(PetscObjectCast(viewer), &comm);
119: MPI_Comm_rank(comm, &rank);
120: PetscViewerGetSubViewer(viewer, PETSC_COMM_SELF, &sviewer);
121: PetscViewerASCIIPrintf(sviewer, "[%d] name: %s\n", rank, dprop_.name);
122: PetscViewerASCIIPushTab(sviewer);
123: PetscViewerASCIIPrintf(sviewer, "Compute capability: %d.%d\n", dprop_.major, dprop_.minor);
124: PetscViewerASCIIPrintf(sviewer, "Multiprocessor Count: %d\n", dprop_.multiProcessorCount);
125: PetscViewerASCIIPrintf(sviewer, "Maximum Grid Dimensions: %d x %d x %d\n", dprop_.maxGridSize[0], dprop_.maxGridSize[1], dprop_.maxGridSize[2]);
126: PetscViewerASCIIPrintf(sviewer, "Maximum Block Dimensions: %d x %d x %d\n", dprop_.maxThreadsDim[0], dprop_.maxThreadsDim[1], dprop_.maxThreadsDim[2]);
127: PetscViewerASCIIPrintf(sviewer, "Maximum Threads Per Block: %d\n", dprop_.maxThreadsPerBlock);
128: PetscViewerASCIIPrintf(sviewer, "Warp Size: %d\n", dprop_.warpSize);
129: PetscViewerASCIIPrintf(sviewer, "Total Global Memory (bytes): %zu\n", dprop_.totalGlobalMem);
130: PetscViewerASCIIPrintf(sviewer, "Total Constant Memory (bytes): %zu\n", dprop_.totalConstMem);
131: PetscViewerASCIIPrintf(sviewer, "Shared Memory Per Block (bytes): %zu\n", dprop_.sharedMemPerBlock);
132: PetscViewerASCIIPrintf(sviewer, "Multiprocessor Clock Rate (KHz): %d\n", dprop_.clockRate);
133: PetscViewerASCIIPrintf(sviewer, "Memory Clock Rate (KHz): %d\n", dprop_.memoryClockRate);
134: PetscViewerASCIIPrintf(sviewer, "Memory Bus Width (bits): %d\n", dprop_.memoryBusWidth);
135: PetscViewerASCIIPrintf(sviewer, "Peak Memory Bandwidth (GB/s): %f\n", 2.0 * dprop_.memoryClockRate * (dprop_.memoryBusWidth / 8) / 1.0e6);
136: PetscViewerASCIIPrintf(sviewer, "Can map host memory: %s\n", dprop_.canMapHostMemory ? "PETSC_TRUE" : "PETSC_FALSE");
137: PetscViewerASCIIPrintf(sviewer, "Can execute multiple kernels concurrently: %s\n", dprop_.concurrentKernels ? "PETSC_TRUE" : "PETSC_FALSE");
138: PetscViewerASCIIPopTab(sviewer);
139: PetscViewerRestoreSubViewer(viewer, PETSC_COMM_SELF, &sviewer);
140: PetscViewerFlush(viewer);
141: }
142: return 0;
143: }
145: template <DeviceType T>
146: PetscErrorCode Device<T>::DeviceInternal::getattribute(PetscDeviceAttribute attr, void *value) const noexcept
147: {
148: PetscAssert(initialized(), PETSC_COMM_SELF, PETSC_ERR_COR, "Device %d was not initialized", id());
149: switch (attr) {
150: case PETSC_DEVICE_ATTR_SIZE_T_SHARED_MEM_PER_BLOCK:
151: *static_cast<std::size_t *>(value) = prop().sharedMemPerBlock;
152: case PETSC_DEVICE_ATTR_MAX:
153: break;
154: }
155: return 0;
156: }
158: static std::jmp_buf cupmMPIAwareJumpBuffer;
159: static bool cupmMPIAwareJumpBufferSet;
161: // godspeed to anyone that attempts to call this function
162: void SilenceVariableIsNotNeededAndWillNotBeEmittedWarning_ThisFunctionShouldNeverBeCalled()
163: {
164: PETSCABORT(MPI_COMM_NULL, INT_MAX);
165: if (cupmMPIAwareJumpBufferSet) (void)cupmMPIAwareJumpBuffer;
166: }
168: template <DeviceType T>
169: PETSC_CXX_COMPAT_DEFN(PetscErrorCode Device<T>::DeviceInternal::CUPMAwareMPI_(bool *awareness))
170: {
171: constexpr int bufSize = 2;
172: constexpr int hbuf[bufSize] = {1, 0};
173: int *dbuf = nullptr;
174: constexpr auto bytes = bufSize * sizeof(*dbuf);
175: const auto cupmSignalHandler = [](int signal, void *ptr) -> PetscErrorCode {
176: if ((signal == SIGSEGV) && cupmMPIAwareJumpBufferSet) std::longjmp(cupmMPIAwareJumpBuffer, 1);
177: return PetscSignalHandlerDefault(signal, ptr);
178: };
180: *awareness = false;
181: cupmMalloc(reinterpret_cast<void **>(&dbuf), bytes);
182: cupmMemcpy(dbuf, hbuf, bytes, cupmMemcpyHostToDevice);
183: cupmDeviceSynchronize();
184: PetscPushSignalHandler(cupmSignalHandler, nullptr);
185: cupmMPIAwareJumpBufferSet = true;
186: if (!setjmp(cupmMPIAwareJumpBuffer) && !MPI_Allreduce(dbuf, dbuf + 1, 1, MPI_INT, MPI_SUM, PETSC_COMM_SELF)) *awareness = true;
187: cupmMPIAwareJumpBufferSet = false;
188: PetscPopSignalHandler();
189: cupmFree(dbuf);
190: return 0;
191: }
193: template <DeviceType T>
194: PetscErrorCode Device<T>::finalize_() noexcept
195: {
196: if (PetscUnlikely(!initialized_)) return 0;
197: for (auto &&device : devices_) device.reset();
198: defaultDevice_ = PETSC_CUPM_DEVICE_NONE; // disabled by default
199: initialized_ = false;
200: return 0;
201: }
203: template <DeviceType T>
204: PETSC_CXX_COMPAT_DECL(PETSC_CONSTEXPR_14 const char *CUPM_VISIBLE_DEVICES())
205: {
206: switch (T) {
207: case DeviceType::CUDA:
208: return "CUDA_VISIBLE_DEVICES";
209: case DeviceType::HIP:
210: return "HIP_VISIBLE_DEVICES";
211: }
212: PetscUnreachable();
213: return "PETSC_ERROR_PLIB";
214: }
216: template <DeviceType T>
217: PetscErrorCode Device<T>::initialize(MPI_Comm comm, PetscInt *defaultDeviceId, PetscBool *defaultView, PetscDeviceInitType *defaultInitType) noexcept
218: {
219: auto initId = std::make_pair(*defaultDeviceId, PETSC_FALSE);
220: auto initView = std::make_pair(*defaultView, PETSC_FALSE);
221: auto initType = std::make_pair(*defaultInitType, PETSC_FALSE);
222: int ndev = 0;
224: if (initialized_) return 0;
225: initialized_ = true;
226: PetscRegisterFinalize(finalize_);
227: base_type::PetscOptionDeviceAll(comm, initType, initId, initView);
229: if (initType.first == PETSC_DEVICE_INIT_NONE) {
230: initId.first = PETSC_CUPM_DEVICE_NONE;
231: } else if (const auto cerr = cupmGetDeviceCount(&ndev)) {
232: auto PETSC_UNUSED ignored = cupmGetLastError();
233: // we won't be initializing anything anyways
234: initType.first = PETSC_DEVICE_INIT_NONE;
235: // save the error code for later
236: initId.first = -static_cast<decltype(initId.first)>(cerr);
239: }
241: // check again for init type, since the device count may have changed it
242: if (initType.first == PETSC_DEVICE_INIT_NONE) {
243: // id < 0 (excluding PETSC_DECIDE) indicates an error has occurred during setup
244: if ((initId.first > 0) || (initId.first == PETSC_DECIDE)) initId.first = PETSC_CUPM_DEVICE_NONE;
245: // initType overrides initView
246: initView.first = PETSC_FALSE;
247: } else {
248: PetscDeviceCheckDeviceCount_Internal(ndev);
249: if (initId.first == PETSC_DECIDE) {
250: if (ndev) {
251: PetscMPIInt rank;
253: MPI_Comm_rank(comm, &rank);
254: initId.first = rank % ndev;
255: } else initId.first = 0;
256: }
257: if (initView.first) initType.first = PETSC_DEVICE_INIT_EAGER;
258: }
260: static_assert(std::is_same<PetscMPIInt, decltype(defaultDevice_)>::value, "");
261: // initId.first is PetscInt, _defaultDevice is int
262: PetscMPIIntCast(initId.first, &defaultDevice_);
263: // record the results of the initialization
264: *defaultDeviceId = initId.first;
265: *defaultView = initView.first;
266: *defaultInitType = initType.first;
267: return 0;
268: }
270: template <DeviceType T>
271: PetscErrorCode Device<T>::init_device_id_(PetscInt *inid) const noexcept
272: {
273: const auto id = *inid == PETSC_DECIDE ? defaultDevice_ : *inid;
274: const auto cerr = static_cast<cupmError_t>(-defaultDevice_);
278: PetscAssert(static_cast<decltype(devices_.size())>(id) < devices_.size(), PETSC_COMM_SELF, PETSC_ERR_ARG_OUTOFRANGE, "Only supports %zu number of devices but trying to get device with id %" PetscInt_FMT, devices_.size(), id);
280: if (!devices_[id]) devices_[id] = util::make_unique<DeviceInternal>(id);
281: PetscAssert(id == devices_[id]->id(), PETSC_COMM_SELF, PETSC_ERR_PLIB, "Entry %" PetscInt_FMT " contains device with mismatching id %d", id, devices_[id]->id());
282: devices_[id]->initialize();
283: *inid = id;
284: return 0;
285: }
287: template <DeviceType T>
288: PetscErrorCode Device<T>::configure_device_(PetscDevice device) noexcept
289: {
290: devices_[device->deviceId]->configure();
291: return 0;
292: }
294: template <DeviceType T>
295: PetscErrorCode Device<T>::view_device_(PetscDevice device, PetscViewer viewer) noexcept
296: {
297: // now this __shouldn't__ reconfigure the device, but there is a petscinfo call to indicate
298: // it is being reconfigured
299: devices_[device->deviceId]->configure();
300: devices_[device->deviceId]->view(viewer);
301: return 0;
302: }
304: template <DeviceType T>
305: PetscErrorCode Device<T>::get_attribute_(PetscInt id, PetscDeviceAttribute attr, void *value) noexcept
306: {
307: devices_[id]->getattribute(attr, value);
308: return 0;
309: }
311: // explicitly instantiate the classes
312: #if PetscDefined(HAVE_CUDA)
313: template class Device<DeviceType::CUDA>;
314: #endif
315: #if PetscDefined(HAVE_HIP)
316: template class Device<DeviceType::HIP>;
317: #endif
319: } // namespace cupm
321: } // namespace device
323: } // namespace Petsc