-
Notifications
You must be signed in to change notification settings - Fork 74k
/
pjrt_client.h
1572 lines (1390 loc) · 71.3 KB
/
pjrt_client.h
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
375
376
377
378
379
380
381
382
383
384
385
386
387
388
389
390
391
392
393
394
395
396
397
398
399
400
401
402
403
404
405
406
407
408
409
410
411
412
413
414
415
416
417
418
419
420
421
422
423
424
425
426
427
428
429
430
431
432
433
434
435
436
437
438
439
440
441
442
443
444
445
446
447
448
449
450
451
452
453
454
455
456
457
458
459
460
461
462
463
464
465
466
467
468
469
470
471
472
473
474
475
476
477
478
479
480
481
482
483
484
485
486
487
488
489
490
491
492
493
494
495
496
497
498
499
500
501
502
503
504
505
506
507
508
509
510
511
512
513
514
515
516
517
518
519
520
521
522
523
524
525
526
527
528
529
530
531
532
533
534
535
536
537
538
539
540
541
542
543
544
545
546
547
548
549
550
551
552
553
554
555
556
557
558
559
560
561
562
563
564
565
566
567
568
569
570
571
572
573
574
575
576
577
578
579
580
581
582
583
584
585
586
587
588
589
590
591
592
593
594
595
596
597
598
599
600
601
602
603
604
605
606
607
608
609
610
611
612
613
614
615
616
617
618
619
620
621
622
623
624
625
626
627
628
629
630
631
632
633
634
635
636
637
638
639
640
641
642
643
644
645
646
647
648
649
650
651
652
653
654
655
656
657
658
659
660
661
662
663
664
665
666
667
668
669
670
671
672
673
674
675
676
677
678
679
680
681
682
683
684
685
686
687
688
689
690
691
692
693
694
695
696
697
698
699
700
701
702
703
704
705
706
707
708
709
710
711
712
713
714
715
716
717
718
719
720
721
722
723
724
725
726
727
728
729
730
731
732
733
734
735
736
737
738
739
740
741
742
743
744
745
746
747
748
749
750
751
752
753
754
755
756
757
758
759
760
761
762
763
764
765
766
767
768
769
770
771
772
773
774
775
776
777
778
779
780
781
782
783
784
785
786
787
788
789
790
791
792
793
794
795
796
797
798
799
800
801
802
803
804
805
806
807
808
809
810
811
812
813
814
815
816
817
818
819
820
821
822
823
824
825
826
827
828
829
830
831
832
833
834
835
836
837
838
839
840
841
842
843
844
845
846
847
848
849
850
851
852
853
854
855
856
857
858
859
860
861
862
863
864
865
866
867
868
869
870
871
872
873
874
875
876
877
878
879
880
881
882
883
884
885
886
887
888
889
890
891
892
893
894
895
896
897
898
899
900
901
902
903
904
905
906
907
908
909
910
911
912
913
914
915
916
917
918
919
920
921
922
923
924
925
926
927
928
929
930
931
932
933
934
935
936
937
938
939
940
941
942
943
944
945
946
947
948
949
950
951
952
953
954
955
956
957
958
959
960
961
962
963
964
965
966
967
968
969
970
971
972
973
974
975
976
977
978
979
980
981
982
983
984
985
986
987
988
989
990
991
992
993
994
995
996
997
998
999
1000
/* Copyright 2017 The OpenXLA Authors.
Licensed under the Apache License, Version 2.0 (the "License");
you may not use this file except in compliance with the License.
You may obtain a copy of the License at
http://www.apache.org/licenses/LICENSE-2.0
Unless required by applicable law or agreed to in writing, software
distributed under the License is distributed on an "AS IS" BASIS,
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and
limitations under the License.
==============================================================================*/
#ifndef XLA_PJRT_PJRT_CLIENT_H_
#define XLA_PJRT_PJRT_CLIENT_H_
#include <cstddef>
#include <cstdint>
#include <cstdlib>
#include <functional>
#include <memory>
#include <optional>
#include <string>
#include <string_view>
#include <utility>
#include <vector>
#include "absl/base/attributes.h"
#include "absl/base/thread_annotations.h"
#include "absl/container/flat_hash_map.h"
#include "absl/container/inlined_vector.h"
#include "absl/functional/any_invocable.h"
#include "absl/status/statusor.h"
#include "absl/strings/string_view.h"
#include "absl/synchronization/mutex.h"
#include "absl/synchronization/notification.h"
#include "absl/time/time.h"
#include "absl/types/span.h"
#include "mlir/IR/BuiltinOps.h" // from @llvm-project
#include "xla/client/xla_computation.h"
#include "xla/layout.h"
#include "xla/literal.h"
#include "xla/pjrt/pjrt_common.h"
#include "xla/pjrt/pjrt_compiler.h"
#include "xla/pjrt/pjrt_device_description.h"
#include "xla/pjrt/pjrt_executable.h"
#include "xla/pjrt/pjrt_future.h"
#include "xla/pjrt/pjrt_layout.h"
#include "xla/service/computation_placer.h"
#include "xla/service/hlo_cost_analysis.h"
#include "xla/shape.h"
#include "xla/shape_util.h"
#include "xla/status.h"
#include "xla/statusor.h"
#include "xla/util.h"
#include "xla/xla_data.pb.h"
#include "tsl/framework/allocator.h"
#include "tsl/platform/errors.h"
// API notes:
// PjRt stands for "Pretty much Just another RunTime".
namespace xla {
enum PjRtRuntimeType { kStreamExecutor, kTfrt };
inline constexpr absl::string_view PjRtRuntimeTypeString(PjRtRuntimeType type) {
switch (type) {
case kStreamExecutor:
return "stream_executor";
case kTfrt:
return "tfrt";
}
}
class PjRtClient;
class PjRtDevice;
class PjRtMemorySpace {
public:
virtual ~PjRtMemorySpace() = default;
// The owner of this memory space.
virtual PjRtClient* client() const = 0;
// The devices that this memory space is attached to.
virtual absl::Span<PjRtDevice* const> devices() const = 0;
// The ID of this memory space. IDs are globally unique across all hosts.
virtual int id() const = 0;
// A platform-dependent string that uniquely identifies the kind of the
// memory space.
virtual absl::string_view kind() const = 0;
// An ID uniquely identifies the kind of the memory space among those attached
// to the same `PjRtClient`. The IDs assigned to a kind is implementation
// specific.
virtual int kind_id() const = 0;
// Debug string suitable for logging when errors occur. Should be verbose
// enough to describe the current memory space unambiguously.
virtual absl::string_view DebugString() const = 0;
// Debug string suitable for reading by end users, should be reasonably terse.
virtual absl::string_view ToString() const = 0;
};
class PjRtDevice {
public:
virtual ~PjRtDevice() = default;
// Return the client that owns this device.
virtual PjRtClient* client() const = 0;
// Whether client can issue command to this device.
virtual bool IsAddressable() const = 0;
virtual const PjRtDeviceDescription& description() const {
LOG(FATAL) << "PjRtDeviceDescription not available (must override "
"PjRtDevice::description).";
}
// The ID of this device. IDs are unique among devices of this type
// (e.g. CPUs, GPUs). On multi-host platforms, this will be unique across all
// hosts' devices. This is the ID that should be used in a DeviceAssignment.
ABSL_DEPRECATED("Use global_device_id() instead")
virtual int id() const { return global_device_id().value(); }
// There are several different IDs for a PJRT device.
//
// - global_device_id: The logical global device ID. This is unique among
// devices of this type (e.g. CPUs, GPUs). On multi-host platforms, this will
// be unique across all hosts' devices. This is the ID that should be used in
// a DeviceAssignment.
//
// - local_device_id: The logical local device ID. This will be used to look
// up an addressable device local to a given client. It is -1 if undefined.
//
// - local_hardware_id: The physical local device ID, e.g., the CUDA device
// number. Multiple PJRT devices can have the same local_hardware_id if
// these PJRT devices share the same physical device. This is useful for
// identifying which physical device when interacting with non-JAX code. In
// general, not guaranteed to be dense, and -1 if undefined.
// TODO(b/314368788): Remove `id()` and replace it with this function.
virtual PjRtGlobalDeviceId global_device_id() const {
return PjRtGlobalDeviceId(description().id());
}
virtual PjRtLocalDeviceId local_device_id() const {
// By default, local_device_id is the same as local_hardware_id when there
// is only one PJRT device on a physical device.
return PjRtLocalDeviceId(local_hardware_id_typed().value());
}
// TODO(b/314368788): Remove `int local_hardware_id()` and rename this
// function to `local_hardware_id()`.
virtual PjRtLocalHardwareId local_hardware_id_typed() const = 0;
// The index of the process that this device belongs to, i.e. is addressable
// from. This is not always identical to PjRtClient::process_index() in a
// multi-process setting, where each client can see devices from all
// processes, but only a subset of them are addressable and have the same
// process_index as the client.
virtual int process_index() const { return description().process_index(); }
// Opaque hardware ID, e.g., the CUDA device number, useful for identifying
// which GPU when interacting with non-JAX code. In general, not guaranteed to
// be dense, and -1 if undefined.
ABSL_DEPRECATED("Use local_hardware_id_typed() instead")
virtual int local_hardware_id() const {
return local_hardware_id_typed().value();
}
// A vendor-dependent string that uniquely identifies the kind of device,
// e.g., "Tesla V100-SXM2-16GB". May be used to determine whether two GPUs are
// compatible compilation.
virtual absl::string_view device_kind() const {
return description().device_kind();
}
// Debug string suitable for logging when errors occur. Should be verbose
// enough to describe the current device unambiguously.
virtual absl::string_view DebugString() const {
return description().DebugString();
}
// Debug string suitable for reading by end users, should be reasonably terse,
// for example: "CpuDevice(id=0)".
virtual absl::string_view ToString() const {
return description().ToString();
}
// Returns vendor specific attributes about the device. For example the model
// number of a GPU, or the mesh coordinates of a TPU device. The returned
// reference will remain valid for the lifetime of the PjRtDevice.
virtual const absl::flat_hash_map<std::string, PjRtDeviceAttribute>&
Attributes() const {
return description().Attributes();
}
// Returns a scoped event that the caller uses to tell the PjRtClient that
// there is asynchronous work happening that depends on activity on the
// PjRtDevice. See comment on class definition in pjrt_future.h.
//
// Only some PjRtDevice implementations support ScopedAsyncTrackingEvent, and
// those that do not will return nullptr.
virtual std::unique_ptr<ScopedAsyncTrackingEvent> CreateAsyncTrackingEvent(
absl::string_view description) const = 0;
// Transfer the given literal to the infeed queue.
virtual Status TransferToInfeed(const LiteralSlice& literal) = 0;
// Transfer and return a value of the given shape from the outfeed queue.
virtual Status TransferFromOutfeed(MutableBorrowingLiteral literal) = 0;
// Returns allocator stats for the device. Only some PjRtDevice
// implementations support allocator_stats, and those that do not will return
// an Unimplemented error.
virtual StatusOr<tsl::AllocatorStats> GetAllocatorStats() const {
return Unimplemented("GetAllocatorStats is not supported");
}
// Returns all memory spaces attached to this device.
// The memory spaces are in no particular order.
virtual absl::Span<PjRtMemorySpace* const> memory_spaces() const = 0;
// Returns the default memory space attached to this device.
virtual StatusOr<PjRtMemorySpace*> default_memory_space() const = 0;
virtual absl::StatusOr<PjRtMemorySpace*> memory_space_by_kind(
absl::string_view memory_space_kind) const {
return Unimplemented("memory_space_by_kind not implemented");
}
// Returns a platform-specific stream handle that should be used to track when
// an externally-managed buffer is ready to use on this device. This is
// intended to support dlpack on GPU and is not expected to be implemented for
// all hardware platforms.
virtual StatusOr<std::intptr_t> GetStreamForExternalReadyEvents() const {
return Unimplemented(
"PjRtDevice::GetStreamForExternalReadyEvents only implemented for "
"GPU");
}
// Experimental: Poisons the earliest execution on this device with given
// launch_id if it's not finished yet, i.e. makes its output buffers error.
//
// Returns true if the output buffers have been successfully poisoned.
//
// Returns false if the output buffers were not successfully poisoned because
// launch_id is not in the list of executions that have not yet completed.
// This may happen either because the execution corresponding to launch_id has
// already completed, or because an incorrect launch_id was supplied.
//
// Returns error otherwise, including in the case that poisoning is not
// implemented by this client.
virtual StatusOr<bool> PoisonExecution(int32_t launch_id, Status error) {
return Unimplemented("PoisonExecution is not supported");
}
};
// Forward declaration.
class PjRtBuffer;
// Helper struct for cross host transfers, returned by the callback from a call
// to PjRtBuffer::MakeCrossHostReceiveBuffers or
// PjRtBuffer::MakeCrossHostReceiveBuffersForGather.
struct PjRtCrossHostRecvDescriptors {
// There is one serialized_descriptor per sub-buffer being gathered (i.e. a
// single descriptor if the buffer is returned from a call to
// MakeCrossHostReceiveBuffers). The descriptor should be transmitted to the
// sender(s) and passed to a call to src_buffer->CopyToRemoteDevice.
absl::InlinedVector<std::string, 1> serialized_descriptors;
};
// Function that the client should call at the receiver if it needs to cancel a
// cross-host send, for example because the buffer that the remote host wanted
// to send is not available. The serialized descriptor should match one of the
// descriptors returned in a PjRtCrossHostRecvDescriptors. on_canceled will be
// called once cancellation is complete and indicates whether cancellation was
// successful or not.
//
// For each serialized_descriptor provided in a PjRtCrossHostRecvDescriptors,
// *either* the sending host must successfully complete a CopyToRemoteDevice
// for that descriptor, *or* the receiving host must cancel. If there is a
// duplicate (e.g., both send and cancel) then the system will be left in an
// undefined state. If there is no send or cancellation then the system will
// hang indefinitely.
using PjRtCrossHostSendCancelNotifier =
std::function<void(absl::string_view serialized_descriptor, Status reason,
std::function<void(Status)> on_canceled)>;
// State asynchronously returned by MakeCrossHostReceiveBuffers. "descriptors"
// will match the returned PjRtBuffer objects 1:1. Specifically, each PjRtBuffer
// returned by MakeCrossHostReceiveBuffers will have one
// PjRtCrossHostRecvDescriptors object containing it descriptor(s).
struct PjRtCrossHostRecvState {
std::vector<PjRtCrossHostRecvDescriptors> descriptors;
PjRtCrossHostSendCancelNotifier cancel_notifier;
};
using PjRtCrossHostRecvNotifier =
std::function<void(StatusOr<PjRtCrossHostRecvState>)>;
// A sized chunk of host data. The host data can be either in host layout or in
// device layout, and it can be one part of the entire buffer. The PjRt
// implementations can customize how the memory is allocated and deallocated.
class PjRtChunk {
public:
// Allocate a PjRtChunk using malloc.
static PjRtChunk AllocateDefault(size_t size) {
return PjRtChunk(malloc(size), size, [](void* ptr) { free(ptr); });
}
PjRtChunk() = default;
PjRtChunk(void* data, size_t size, std::function<void(void*)> deleter)
: data_(static_cast<uint8_t*>(data)),
size_(size),
deleter_(std::move(deleter)) {}
~PjRtChunk() {
if (data_) {
deleter_(data_);
}
}
PjRtChunk(PjRtChunk&& other)
: data_(other.data_),
size_(other.size_),
deleter_(std::move(other.deleter_)) {
other.data_ = nullptr;
}
PjRtChunk& operator=(PjRtChunk&& other) {
if (data_) {
deleter_(data_);
}
data_ = other.data_;
size_ = other.size_;
deleter_ = std::move(other.deleter_);
other.data_ = nullptr;
return *this;
}
PjRtChunk(const PjRtChunk&) = delete;
PjRtChunk& operator=(const PjRtChunk&) = delete;
uint8_t* data() { return data_; }
const uint8_t* data() const { return data_; }
int64_t size() const { return size_; }
std::function<void(void*)> deleter() const { return deleter_; }
// Release the ownership of the data. Note that this does not free the data;
// the caller should copy `data()` and `deleter()` to manage the ownership
// before calling `release()`. This PjRtChunk is invalidated after calling.
void release() {
data_ = nullptr;
size_ = 0;
deleter_ = nullptr;
}
private:
// The ownership of the bytes pointed to by `data_` is controlled by the
// `deleter_`.
uint8_t* data_ = nullptr;
size_t size_ = 0;
std::function<void(void*)> deleter_;
};
// A stream of Chunks from the host to the device. Once the stream enters
// Complete state it never changes state again.
//
// This class is thread-safe.
class CopyToDeviceStream {
public:
CopyToDeviceStream(int64_t total_bytes, int64_t granule_bytes)
: total_bytes_(total_bytes), granule_bytes_(granule_bytes) {}
virtual ~CopyToDeviceStream();
// Emplaces a new Chunk of data to copy to the device. Returns an error future
// if the Chunk's size causes the amount of transferred data to exceed
// total_bytes(), if the stream is already complete, or if the chunk is not a
// multiple of granule_size_in_bytes().
//
// The transfer is started immediately, and the returned future is fulfilled
// when the transfer completes or fails.
virtual PjRtFuture<> AddChunk(PjRtChunk chunk) = 0;
// Returns the total amount of data the stream expects to be transferred.
int64_t total_bytes() const { return total_bytes_; }
// Returns the granule size in bytes. The size of the chunk added to this
// stream must be a multiple of this number.
int64_t granule_size_in_bytes() const { return granule_bytes_; }
// Returns the amount of data the stream currently has either transferred or
// has buffered to transfer.
int64_t current_bytes() const ABSL_LOCKS_EXCLUDED(mu_) {
absl::MutexLock lock(&mu_);
return current_bytes_;
}
// Returns true if the stream is complete; all expected bytes have been
// transferred or are buffered to transfer.
bool IsComplete() const ABSL_LOCKS_EXCLUDED(mu_) {
absl::MutexLock lock(&mu_);
return IsCompleteLocked();
}
// Returns true if the stream is empty; no data has been queued.
bool empty() const { return current_bytes() == 0; }
protected:
bool IsCompleteLocked() const ABSL_EXCLUSIVE_LOCKS_REQUIRED(mu_) {
return current_bytes_ == total_bytes_;
}
int64_t total_bytes_;
int64_t granule_bytes_;
int64_t current_bytes_ ABSL_GUARDED_BY(mu_) = 0;
mutable absl::Mutex mu_;
};
class PjRtHostMemoryForDeviceManager {
public:
virtual ~PjRtHostMemoryForDeviceManager();
// Transforms the host memory representations of a shape with the host layout
// to the host memory representation of the same shape with the device layout.
// `src_shape` and `dst_shape` may only differ in their layouts.
virtual StatusOr<PjRtChunk> ToDeviceLayout(const void* src_data,
size_t src_size,
const Shape& host_shape,
const Shape& device_shape) = 0;
// Transforms the host memory representations of a shape with the device
// layout to the host memory representation of the same shape with the host
// layout. `src_shape` and `dst_shape` may only differ in their layouts.
virtual Status ToHostLayout(const void* src_data, size_t src_size,
const Shape& src_shape, void* dst_data,
size_t dst_size, const Shape& dst_shape) = 0;
};
class PjRtLoadedExecutable;
struct PjRtPluginAttributes {
int64_t pjrt_c_api_major_version;
int64_t pjrt_c_api_minor_version;
absl::flat_hash_map<std::string, PjRtValueType> attributes;
};
// Encapsulates the state of Python session with XLA.
//
// It is the responsibility of the client of this API to keep the PjRtClient
// alive as long as any of the other runtime objects are alive.
//
// A note on the semantics of cross-device copies.
//
// There are two mechanisms to transfer a buffer from one device to another.
// When both devices are on the same host (more specifically, the user program
// ends up with pointers to both the source and destination buffers in the same
// address space), the caller can use:
// dst_buffer = src_buffer->CopyToDevice(dst_device)
//
// When the source and destination are on different hosts, but the transfer is
// made via native device networking (as opposed to the user program fetching
// the buffer and sending it using its own networking code), the caller can
// use:
// DstHost: dst_client->MakeCrossHostReceiveBuffers(...)
// DstHost: [...]
// DstHost: gets callback containing PjRtCrossHostRecvDescriptors
// DstHost: sends cross-host recv serialized descriptors to SrcHost
// SrcHost: src_buffer->CopyToRemoteDevice(serialized_descriptors)
//
// Note that in the cross-host case, the dst_client may call
// MakeCrossHostReceiveBuffers before the action that produces src_buffer has
// been enqueued at SrcHost.
//
// On some platforms, device-to-device transfers consume scarce hardware
// resources. If dst_client->MakeCrossHostReceiveBuffers immediately claimed
// those resources, then there would be a risk of system-wide deadlock, if the
// resources claimed by the recv prevented other transfers that are necessary
// to generate src_buffer from acquiring enough resources to proceed.
//
// In order to allow clients to avoid deadlocks such as those in the preceding
// paragraph, PjRtClient guarantees progress but not fairness with respect to
// the order that cross-device transfers are enqueued on a given host, as
// follows:
//
// The progress guarantee is that a cross-device transfer T on host A will not
// claim scarce hardware resources until it is guaranteed that all transfers
// enqueued on A before T have already either completed, or been assigned enough
// resources to ensure that they can eventually complete.
//
// The lack of a fairness guarantee means that, if cross-device transfer T1 is
// enqueued before transfer T2 at A, then T2 may complete before T1. T1 may be
// delayed for an unbounded time waiting for T2 if T2 is large, even though T1
// will eventually be able to make progress.
class PjRtClient {
public:
PjRtClient() = default;
explicit PjRtClient(std::unique_ptr<PjRtHostMemoryForDeviceManager>
host_memory_for_device_manager)
: host_memory_for_device_manager_(
std::move(host_memory_for_device_manager)) {}
virtual ~PjRtClient() = default;
// Return the process index of this client. Always 0 in single-process
// settings.
virtual int process_index() const = 0;
// Return the number of devices in the entire computation. In multi-headed
// client setting, some are addressable by this client, some are not. In a
// single-client setting, this is equal to the number of addressable devices.
virtual int device_count() const = 0;
// Return number of addressable devices. Addressable devices are those that
// the client can issue commands to.
virtual int addressable_device_count() const = 0;
// Return all devices known to the client, including addressable and
// non-addressable devices.
virtual absl::Span<PjRtDevice* const> devices() const = 0;
// Return only addressable devices. The devices are in no particular order.
virtual absl::Span<PjRtDevice* const> addressable_devices() const = 0;
// Lookup any PjRtDevice for a given PjRtDevice::id().
ABSL_DEPRECATED("Use LookupDevice(PjRtGlobalDeviceId) instead")
virtual StatusOr<PjRtDevice*> LookupDevice(int device_id) const {
return LookupDevice(PjRtGlobalDeviceId(device_id));
}
virtual StatusOr<PjRtDevice*> LookupDevice(
PjRtGlobalDeviceId global_device_id) const = 0;
// Return an addressable PjRtDevice for a given
// PjRtDevice::local_hardware_id().
ABSL_DEPRECATED("Use LookupAddressableDevice(PjRtLocalDeviceId) instead")
virtual StatusOr<PjRtDevice*> LookupAddressableDevice(
int local_hardware_id) const {
return LookupAddressableDevice(PjRtLocalDeviceId(local_hardware_id));
}
virtual StatusOr<PjRtDevice*> LookupAddressableDevice(
PjRtLocalDeviceId local_device_id) const = 0;
// Return all memory spaces owned by the client.
// The memory spaces are in no particular order.
virtual absl::Span<PjRtMemorySpace* const> memory_spaces() const = 0;
// Return an ID that identifies the platform (CPU/GPU/TPU).
virtual PjRtPlatformId platform_id() const = 0;
// Returns a string that identifies the platform (CPU/GPU/TPU).
virtual absl::string_view platform_name() const = 0;
// Returns a string containing human-readable, platform-specific version info
// (e.g. the CUDA version on GPU or libtpu version on Cloud TPU).
virtual absl::string_view platform_version() const = 0;
// Returns information about the underlying PJRT C API plugin if such a plugin
// is being used, otherwise returns nullopt.
virtual std::optional<PjRtPluginAttributes> plugin_attributes() const {
return std::nullopt;
}
// TODO(b/244756954): Rethink this function altogether
// Returns an enum that identifies the type of runtime being used under this
// client.
virtual PjRtRuntimeType runtime_type() const = 0;
// Return a device-specific default device assignment, e.g., GPU and TPU may
// be different.
virtual StatusOr<DeviceAssignment> GetDefaultDeviceAssignment(
int num_replicas, int num_partitions) const = 0;
// Returns a device-specific default device assignment for multi-slice system.
// If num_replicas_per_slice is not defined (nullopt) then we assume that
// all the partitions live entirely on a single slice and that all cross slice
// communication happens across replicas assuming then that
// num_replicas_per_slice is going to be "num_replicas / num_slices".
// TODO(zhangqiaorjc): Convert this to pure virtual and push down.
virtual StatusOr<DeviceAssignment> GetDefaultDeviceAssignment(
int num_replicas, std::optional<int> num_replicas_per_slice,
int num_partitions, const MultiSliceConfig* multi_slice_config) const {
return Unimplemented("Multi slice device assignment is not supported.");
}
// Returns the default device layout for a buffer with `element_type` and
// `dims`. The default layout is a platform-specific layout used when no other
// layout is specified, e.g. for host-to-device transfers. When compiling, the
// default layout is used for program arguments and outputs unless
// user-specified or compiler-chosen layouts are requested via the
// "mhlo.layout_mode" attribute.
virtual StatusOr<Layout> GetDefaultLayout(PrimitiveType element_type,
absl::Span<const int64_t> dims) = 0;
// Returns a backend-specific HLO cost analysis visitor.
virtual StatusOr<std::unique_ptr<HloCostAnalysis>> GetHloCostAnalysis()
const = 0;
// Compile `computation` with given `options`.
virtual StatusOr<std::unique_ptr<PjRtLoadedExecutable>> Compile(
const XlaComputation& computation, CompileOptions options) = 0;
// Variant of `Compile` that accepts an MLIR module.
virtual StatusOr<std::unique_ptr<PjRtLoadedExecutable>> Compile(
mlir::ModuleOp module, CompileOptions options) = 0;
// Deserializes a serialized executable as produced by
// PjRtExecutable::SerializeExecutable(). `serialized` must have been
// produced by a compiler of the same platform and version as this one.
//
// Pending completion of b/237720161, `options` is a mandatory argument in
// most implementations of this interface. They _are_ optional for
// implementations related to the PJRT C API.
virtual StatusOr<std::unique_ptr<PjRtLoadedExecutable>> DeserializeExecutable(
absl::string_view serialized, std::optional<CompileOptions> options) = 0;
// LoadSerializedExecutable takes the serialized output of PjRtExecutable. The
// returned executable is loaded by this client. The same checks are made as
// in Load that the serialized executable is compatible with the client.
virtual StatusOr<std::unique_ptr<PjRtLoadedExecutable>>
LoadSerializedExecutable(absl::string_view serialized,
std::optional<CompileOptions> options,
const LoadOptions& load_options) {
return Unimplemented("Loading serialized executable not supported.");
}
// Loads the executable returns aa PjRtLoadedExecutable runnable by this
// client. Returns an error if the PjRtExecutable was created with an
// incompatible topology or client.
// PjRtExecutable contains a copy of the CompileOptions that was used to
// generate the executable. Load will use the CompileOptions from within the
// executable.
virtual StatusOr<std::unique_ptr<PjRtLoadedExecutable>> Load(
std::unique_ptr<PjRtExecutable> executable,
const LoadOptions& load_options) {
return Unimplemented("Loading executable not supported.");
}
// Creates a buffer on the device without initializing or copying any data.
virtual StatusOr<std::unique_ptr<PjRtBuffer>> CreateUninitializedBuffer(
const Shape& shape, PjRtDevice* device) = 0;
// Creates buffer that carries an error future without allocating memory.
virtual StatusOr<std::unique_ptr<PjRtBuffer>> CreateErrorBuffer(
Status error, const Shape& shape, PjRtDevice* device) {
return Unimplemented("CreateErrorBuffer not supported.");
}
// Gets the pointer to the topology description held by the client.
virtual StatusOr<const PjRtTopologyDescription*> GetTopologyDescription()
const {
return Unimplemented("GetTopologyDescription not supported on platform %s",
platform_name());
}
// Returns topology object for compilation based on this client's topology.
virtual StatusOr<const PjRtTopologyDescription*>
GetFullTopologyForCompilation() const {
return GetTopologyDescription();
}
// A client may want to create a buffer, and hand the buffer to other PjRt
// methods, before the data to store in the buffer is available to the client.
// This is supported using CreateBuffersForAsyncHostToDevice, which returns an
// AsyncHostToDeviceTransferManager helper object.
//
// The PjRtBuffers can be retrieved from the AsyncHostToDeviceTransferManager
// and safely passed immediately to downstream PjRt method calls. Subsequently
// the client can call methods on the AsyncHostToDeviceTransferManager object
// to copy data into the buffers, and once the data copies are complete, the
// buffers' definition events will automatically become ready, unblocking
// downstream consumers of the buffers.
//
// A single call to CreateBuffersForAsyncHostToDevice creates a "batch" of
// buffers that share a single definition event, which may amortize some
// performance overheads, but means that none of the buffers are available to
// downstream consumers until all the transfers have completed. Multiple calls
// to CreateBuffersForAsyncHostToDevice should be made if it is desirable for
// buffers to become available as soon as transfers into them complete.
// Helper class to all clients to asynchronously transfer data into buffers
// that are created uninitialized, see comments immediately above.
class AsyncHostToDeviceTransferManager {
public:
virtual ~AsyncHostToDeviceTransferManager() = default;
// Returns the number of buffers managed by this object.
virtual size_t buffer_count() const = 0;
// Returns the destination device of the transfers.
virtual PjRtDevice* device() const = 0;
// Returns buffer_index, which can be passed to downstream consumers
// immediately and will become available once transfers complete. May not
// be called more than once for a given buffer_index.
//
// RetrieveBuffer can be called at any convenient time; transfer methods
// can safely be called for a buffer index after RetrieveBuffer has been
// called.
virtual std::unique_ptr<PjRtBuffer> RetrieveBuffer(int buffer_index) = 0;
// Transfers 'literal' into buffer_index. No transfer calls into
// buffer_index can be made after this call. on_done is called when the
// transfer is complete but before the buffers are made available to
// their consumers. 'literal' must remain in scope until on_done is
// called.
virtual Status TransferLiteralToBuffer(
int buffer_index, const LiteralSlice& literal,
absl::AnyInvocable<void() &&> on_done) = 0;
// Returns the on-device size in bytes of buffer buffer_index.
virtual size_t buffer_size(int buffer_index) const = 0;
// Transfers 'data' into buffer_index. 'data' must be already laid out in
// the correct on-device format, for example returned by a call to
// buffer->CopyRawToHost. No transfer calls (or SetBufferError calls) into
// buffer_index can be made after this call. on_done is called when the
// transfer is complete but before the buffers are made available to their
// consumers. 'data' must remain in scope until on_done is called.
virtual Status TransferRawDataToBuffer(
int buffer_index, absl::string_view data,
absl::AnyInvocable<void() &&> on_done) = 0;
// Transfers 'data' into a sub-buffer of buffer_index starting at offset, of
// length transfer_size. 'data' must be already laid out in the correct
// on-device format, for example returned by a call to
// buffer->CopyRawToHost. If is_last_transfer is false then the buffer
// remains unavailable to consumers after the transfer completes. If
// is_last_transfer is true then the buffer becomes available to consumers
// after the transfer completes, and no transfer calls (or SetBufferError
// calls) into buffer_index can be made after this call. on_done is called
// when the transfer is complete but before the buffers are made available
// to their consumers. 'data' must remain in scope until on_done is called.
virtual Status TransferRawDataToSubBuffer(
int buffer_index, const void* data, int64_t offset,
int64_t transfer_size, bool is_last_transfer,
absl::AnyInvocable<void() &&> on_done) = 0;
// Indicates that a specific buffer should result in an error status. No
// transfer calls (or further SetBufferError calls) into buffer_index can
// be made after this call.
virtual void SetBufferError(int buffer_index, Status error) = 0;
// Adds the specified key/value metadata for the transfer operation.
// This is typically used for debugging purposes, such as adding a handle
// that can be used to identify transfer operations.
using TransferMetadata = absl::flat_hash_map<std::string, std::string>;
virtual void AddTransferMetadata(const TransferMetadata& metadata) = 0;
};
// Returns a manager for async transfers into a set of buffers with on-host
// shapes 'shapes'.
virtual StatusOr<std::unique_ptr<AsyncHostToDeviceTransferManager>>
CreateBuffersForAsyncHostToDevice(absl::Span<const Shape> shapes,
PjRtDevice* device) = 0;
// Variant of CreateBuffersForAsyncHostToDevice with PjRtMemorySpace.
virtual StatusOr<std::unique_ptr<AsyncHostToDeviceTransferManager>>
CreateBuffersForAsyncHostToDevice(absl::Span<const Shape> shapes,
PjRtMemorySpace* memory_space) = 0;
// Creates a shapeless buffer on the device that can be partitioned into
// multiple PjRtBuffer. This class is an Arena version of
// `AsyncHostToDeviceTransferManager`.
// As a low-level interface, the user must make sure that invocations of
// `Slice` match properly with the writes from `TransferRawDataToSubBuffer`.
//
// For the intended application to Arena allocation / transfer, the user can
// use `GetOnDeviceSizeInBytes` to calculate the offsets for the host buffers
// that need to be transferred.
class PjRtRawDeviceBuffer {
public:
virtual ~PjRtRawDeviceBuffer() = default;
// Transfers data to the device buffer. Data should already be in the
// device layout.
virtual Status TransferRawDataToSubBuffer(
const void* data, int64_t offset, int64_t transfer_size,
bool is_last_transfer, absl::AnyInvocable<void() &&> on_done) = 0;
// The resulting buffer becomes ready when all transfers complete.
virtual StatusOr<std::unique_ptr<PjRtBuffer>> Slice(
int64_t offset, PrimitiveType type, absl::Span<int64_t const> dims,
const Layout& layout) = 0;
};
// Creates a raw device buffer of a given size in bytes.
virtual StatusOr<std::unique_ptr<PjRtRawDeviceBuffer>> CreateRawDeviceBuffer(
int64_t size, PjRtDevice* device) {
return Unimplemented("CreateRawDeviceBuffer is not implemented.");
}
// On-device bytes required for a PjRt buffer with these `Shape` attributes.
virtual StatusOr<int64_t> GetOnDeviceSizeInBytes(
PrimitiveType type, absl::Span<int64_t const> dims,
const Layout& layout) {
return Unimplemented("GetOnDeviceSizeInBytes is not implemented.");
};
// Describes the semantics the caller to BufferFromHostBuffer expects from the
// runtime, in a total order from most restrictive to least restrictive.
enum class HostBufferSemantics {
// The runtime may not hold references to `data` after the call to
// `BufferFromHostBuffer` completes. The caller promises that `data` is
// immutable and will not be freed only for the duration of the
// BufferFromHostBuffer call. `on_done_with_host_buffer` will be called
// before `BufferFromHostBuffer` returns.
kImmutableOnlyDuringCall,
// The runtime may hold onto `data` after the call to `BufferFromHostBuffer`
// returns while the runtime completes a transfer to the device. The caller
// promises not to mutate or free `data` until the transfer completes, at
// which point the runtime will call `on_done_with_host_buffer`. It is also
// correct to wait on the host (directly or indirectly) for the buffer's
// definition event to complete.
kImmutableUntilTransferCompletes,
// The PjRtBuffer may alias `data` internally and the runtime may use the
// `data` contents as long as the buffer is alive. The runtime promises not
// to mutate contents of the buffer (i.e. it will not use it for aliased
// output buffers). The caller promises to keep `data` alive and also not to
// mutate its contents as long as the buffer is alive; to notify the caller
// that the buffer may be freed, the runtime will call
// `on_done_with_host_buffer` when the PjRtBuffer is freed. On non-CPU
// platforms this acts identically to kImmutableUntilTransferCompletes.
kImmutableZeroCopy,
};
// on_done_with_host_buffer is optional and may be null.
// on_done_with_host_buffer will be called iff an OK status is returned.
//
// `data` points to the backing array of the host buffer. Caution:
// `byte_strides` are allowed to be negative, in which case `data` may need
// to point to the interior of the buffer, not necessarily its start.
//
// If byte_strides is omitted, the array is assumed to have a dense layout
// with dimensions in major-to-minor order.
virtual StatusOr<std::unique_ptr<PjRtBuffer>> BufferFromHostBuffer(
const void* data, PrimitiveType type, absl::Span<int64_t const> dims,
std::optional<absl::Span<int64_t const>> byte_strides,
HostBufferSemantics host_buffer_semantics,
absl::AnyInvocable<void() &&> on_done_with_host_buffer,
PjRtDevice* device) = 0;
// Variant of BufferFromHostBuffer that takes an optional device layout. It is
// used when non-compact layout is preferred.
// TODO(b/275645543): remove BufferFromHostBuffer without optional device
// layout after all the inherited classes and call sites are updated.
virtual StatusOr<std::unique_ptr<PjRtBuffer>> BufferFromHostBuffer(
const void* data, PrimitiveType type, absl::Span<int64_t const> dims,
std::optional<absl::Span<int64_t const>> byte_strides,
HostBufferSemantics host_buffer_semantics,
absl::AnyInvocable<void() &&> on_done_with_host_buffer,
PjRtDevice* device, const Layout* device_layout) {
return tsl::errors::Unimplemented(
"BufferFromHostBuffer with an optional device layout is not "
"implemented on platform: ",
platform_name());
}
// TODO(b/277820585): remove BufferFromHostBuffer with PjRtDevice after the
// migration is done.
virtual StatusOr<std::unique_ptr<PjRtBuffer>> BufferFromHostBuffer(
const void* data, PrimitiveType type, absl::Span<int64_t const> dims,
std::optional<absl::Span<int64_t const>> byte_strides,
HostBufferSemantics host_buffer_semantics,
absl::AnyInvocable<void() &&> on_done_with_host_buffer,
PjRtMemorySpace* memory_space, const Layout* device_layout) {
return tsl::errors::Unimplemented(
"BufferFromHostBuffer with PjRtMemorySpace is not implemented on "
"platform: ",
platform_name());
}
// Note that literal must remain in scope until the transfer has completed, so
// the caller should, for example, wait for GetReadyFuture().Await()
// completes on the return value before letting literal go out of scope.
virtual StatusOr<std::unique_ptr<PjRtBuffer>> BufferFromHostLiteral(
const LiteralSlice& literal, PjRtDevice* device) = 0;
// TODO(b/277820585): remove BufferFromHostLiteral with PjRtDevice after the
// migration is done.
virtual StatusOr<std::unique_ptr<PjRtBuffer>> BufferFromHostLiteral(
const LiteralSlice& literal, PjRtMemorySpace* memory_space) {
return tsl::errors::Unimplemented(
"BufferFromHostLiteral with PjRtMemorySpace is not implemented on "
"platform: ",
platform_name());
}
// Creates a PjRtBuffer that is a non-owned view of an on-device
// buffer (typically allocated by another library).
// on_delete_callback is called when the PjRtBuffer is done with the on-device
// buffer. The buffer may be mutated, for example, if the buffer is donated
// to an Execute operation.
//
// `stream`, if specified, is a platform-specific stream handle that should
// contain the work or events needed to materialize the on-device
// buffer. CreateViewOfDeviceBuffer will append an event to `stream` that
// indicates when the returned buffer is ready to use. This is intended to
// support dlpack on GPU and is not expected to be supported on all hardware
// platforms.
virtual StatusOr<std::unique_ptr<PjRtBuffer>> CreateViewOfDeviceBuffer(
void* device_ptr, const Shape& shape, PjRtDevice* device,
std::function<void()> on_delete_callback,
std::optional<std::intptr_t> stream = std::nullopt) = 0;
// Returns platform-dependent address for the given buffer that is often but
// not guaranteed to be the physical/device address.
virtual StatusOr<std::uintptr_t> UnsafeBufferPointer(PjRtBuffer* buffer);
// Returns a vector of PjRtBuffers that can be used to receive
// cross host transfers using `client` on `device'. Asynchronously calls
// `notifier` once receive descriptors are ready to be communicated to the
// sender. `shapes` must be the exact shapes, with identical layouts,
// corresponding to the buffers that will be sent. When resources for the
// transfer are available, notifier will be called with a vector of
// PjRtCrossHostRecvDescriptors structs, one for each shape in `shapes`. Each
// struct contains an opaque string that should be transmitted to the sending
// host and used in a call to CopyToRemoteDevice. None of the recv buffers
// will become ready until *all* of the sends have completed.
//
// If MakeCrossHostReceiveBuffers returns an error, then `notifier` will not
// be called. Otherwise `notifier` will be called exactly once. In the case
// where `notifier` is called with an error status, then the PjRtBuffers
// returned by MakeCrossHostReceiveBuffers will never yield data.
//
// See note on semantics of cross-device copies in the class definition
// comment for PjRtClient.
virtual StatusOr<std::vector<std::unique_ptr<PjRtBuffer>>>
MakeCrossHostReceiveBuffers(absl::Span<const Shape> shapes,
PjRtDevice* device,
PjRtCrossHostRecvNotifier notifier) = 0;
// Asynchronously makes a vector of PjRtBuffers that can be used to receive
// cross host transfers, as in MakeCrossHostReceiveBuffers above, however
// each buffer expects to be "gathered" using multiple sends, one for each of
// a set of sub-slices of the destination buffer.
//
// For each value in shapes there is a corresponding FullGatherDetails struct
// that describes the sub-slices.
struct GatherDetails {
// The dimensions of the corresponding buffer that the gather slices
// into. These dimensions must be the major dimensions in the on-device
// layout of the buffer, and must all be untiled. The scatter acts as if
// the buffer were transposed/reshaped so that all of these dimensions were
// combined into a single dimension whose size is the product of the
// dimensions, and the slice indices correspond to indices in that single
// combined dimension.
//
// For example, if the shape is [3, 4, 128, 128] with [3, 4] as the major
// dimensions in the layout, and dimensions = {0, 1}, then the buffer is
// treated as if it were shape [12, 128, 128] and the indices in
// slice_boundaries range in [0, 12].
absl::InlinedVector<int, 3> dimensions;
// The cumulative indices in dimension of the slices. For example, if
// shape.dimensions(dimension)==10, setting slice_boundaries to {2, 5, 10}
// would correspond to 3 slices of sizes {2, 3, 5} respectively. If the last
// entry in slice_boundaries is less than the size of the combined gather
// dimension, the trailing data in the buffer is undefined after the receive
// completes.
std::vector<int64_t> slice_boundaries;
};
virtual StatusOr<std::vector<std::unique_ptr<PjRtBuffer>>>
MakeCrossHostReceiveBuffersForGather(
absl::Span<const Shape> shapes, std::vector<GatherDetails> gather_details,
PjRtDevice* device, PjRtCrossHostRecvNotifier notifier) = 0;
// Create ChannelHandles for XLA send/recv.
virtual StatusOr<ChannelHandle> CreateChannelHandle() = 0;
virtual StatusOr<ChannelHandle> CreateDeviceToHostChannelHandle() = 0;
virtual StatusOr<ChannelHandle> CreateHostToDeviceChannelHandle() = 0;
// TODO(zhangqiaorjc): Experimental API to be removed.
// Defragment device memory.
virtual Status Defragment() = 0;
// If false, this client does not support send/recv host callbacks, and
// callers should not set the `send_callbacks` and `recv_callbacks` arguments
// in ExecuteOptions.
virtual bool SupportsSendRecvCallbacks() const { return false; }
// Return the PjRtHostMemoryForDeviceManager for this client. It can be
// nullptr if the implementation does not provide one.
virtual PjRtHostMemoryForDeviceManager* GetPjRtHostMemoryForDeviceManager()
const {
return host_memory_for_device_manager_.get();
}
private:
std::unique_ptr<PjRtHostMemoryForDeviceManager>
host_memory_for_device_manager_;
};
// Holds a reference from Python to a tuple of device buffers. A PjRtBuffer
// can be either valid or invalid. An invalid buffer is one that has never been
// initialized, or a buffer that has been deleted (e.g., by calling Delete, or
// by donating it to a computation that aliases an input parameter to an