| /* Copyright 2017 The TensorFlow Authors. All Rights Reserved. |
| |
| 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. |
| ==============================================================================*/ |
| |
| #include "tensorflow/compiler/xla/service/gpu/gpu_transfer_manager.h" |
| |
| #include <string> |
| #include <utility> |
| #include <vector> |
| |
| #include "absl/memory/memory.h" |
| #include "llvm/IR/DataLayout.h" |
| #include "tensorflow/compiler/xla/literal.h" |
| #include "tensorflow/compiler/xla/literal_util.h" |
| #include "tensorflow/compiler/xla/service/gpu/outfeed_manager.h" |
| #include "tensorflow/compiler/xla/service/gpu/target_constants.h" |
| #include "tensorflow/compiler/xla/shape_util.h" |
| #include "tensorflow/compiler/xla/status_macros.h" |
| #include "tensorflow/compiler/xla/statusor.h" |
| #include "tensorflow/compiler/xla/types.h" |
| #include "tensorflow/compiler/xla/util.h" |
| #include "tensorflow/core/lib/core/errors.h" |
| #include "tensorflow/core/lib/gtl/cleanup.h" |
| #include "tensorflow/core/platform/logging.h" |
| #include "tensorflow/core/platform/stream_executor_no_cuda.h" |
| |
| namespace xla { |
| namespace gpu { |
| |
| // TODO(b/30467474) Once GPU infeed implementation settles, consider |
| // folding back the cpu and gpu infeed implementations into a generic |
| // one if possible. |
| GpuTransferManager::GpuTransferManager(se::Platform::Id id, |
| unsigned pointer_size) |
| : GenericTransferManager(id, pointer_size) {} |
| |
| Status GpuTransferManager::TransferLiteralToInfeed( |
| se::StreamExecutor* executor, const LiteralSlice& literal) { |
| return gpu::GetOrCreateInfeedManager(executor)->TransferLiteralToInfeed( |
| executor, literal); |
| } |
| |
| Status GpuTransferManager::TransferLiteralFromOutfeed( |
| se::StreamExecutor* executor, MutableBorrowingLiteral literal) { |
| return gpu::GetOrCreateOutfeedManager(executor)->TransferLiteralFromOutfeed( |
| executor, literal); |
| } |
| |
| } // namespace gpu |
| } // namespace xla |
| |
| static std::unique_ptr<xla::TransferManager> CreateNVPTXTransferManager() { |
| return absl::make_unique<xla::gpu::GpuTransferManager>( |
| /*id=*/stream_executor::cuda::kCudaPlatformId, |
| /*pointer_size=*/llvm::DataLayout(xla::gpu::nvptx::kDataLayout) |
| .getPointerSize(0 /* default address space */)); |
| } |
| |
| static std::unique_ptr<xla::TransferManager> CreateAMDGPUTransferManager() { |
| return absl::make_unique<xla::gpu::GpuTransferManager>( |
| /*id=*/stream_executor::rocm::kROCmPlatformId, |
| /*pointer_size=*/llvm::DataLayout(xla::gpu::amdgpu::kDataLayout) |
| .getPointerSize(0 /* default address space */)); |
| } |
| |
| static bool InitModule() { |
| xla::TransferManager::RegisterTransferManager( |
| stream_executor::cuda::kCudaPlatformId, &CreateNVPTXTransferManager); |
| xla::TransferManager::RegisterTransferManager( |
| stream_executor::rocm::kROCmPlatformId, &CreateAMDGPUTransferManager); |
| return true; |
| } |
| static bool module_initialized = InitModule(); |