This is an archive of the discontinued LLVM Phabricator instance.

Differential D95856

Set GPU context before {cu,hip}MemHostRegister.
ClosedPublic

Authored by csigg on Feb 2 2021, 1:43 AM.

Details

Reviewers
herhut
Commits
rG8a43ec7faa27: Set GPU context before {cu,hip}MemHostRegister.

Diff Detail

Event Timeline

csigg created this revision.Feb 2 2021, 1:43 AM
Herald added subscribers: teijeong, rdzhabarov, tatianashp and 16 others. · View Herald TranscriptFeb 2 2021, 1:43 AM
csigg requested review of this revision.Feb 2 2021, 1:43 AM
Herald added a project: Restricted Project. · View Herald TranscriptFeb 2 2021, 1:43 AM
Herald added subscribers: stephenneuendorffer, nicolasvasilache. · View Herald Transcript
Harbormaster completed remote builds in B87493: Diff 320715.Feb 2 2021, 2:18 AM
This revision was not accepted when it landed; it landed in state Needs Review.Feb 3 2021, 11:00 AM
Closed by commit rG8a43ec7faa27: Set GPU context before {cu,hip}MemHostRegister. (authored by csigg). · Explain Why
This revision was automatically updated to reflect the committed changes.
csigg added a commit: rG8a43ec7faa27: Set GPU context before {cu,hip}MemHostRegister..

Revision Contents

PathSize
mlir/
tools/
mlir-cuda-runner/
1 line
mlir-rocm-runner/
1 line

Diff 321166

mlir/tools/mlir-cuda-runner/cuda-runtime-wrappers.cpp

Show First 20 Lines • Show All 142 Lines • ▼ Show 20 Lines CUDA_REPORT_IF_ERROR(cuMemcpyAsync(reinterpret_cast<CUdeviceptr>(dst),
sizeBytes, stream)); sizeBytes, stream));
} }
/// Helper functions for writing mlir example code /// Helper functions for writing mlir example code
// Allows to register byte array with the CUDA runtime. Helpful until we have // Allows to register byte array with the CUDA runtime. Helpful until we have
// transfer functions implemented. // transfer functions implemented.
extern "C" void mgpuMemHostRegister(void *ptr, uint64_t sizeBytes) { extern "C" void mgpuMemHostRegister(void *ptr, uint64_t sizeBytes) {
ScopedContext scopedContext;
CUDA_REPORT_IF_ERROR(cuMemHostRegister(ptr, sizeBytes, /*flags=*/0)); CUDA_REPORT_IF_ERROR(cuMemHostRegister(ptr, sizeBytes, /*flags=*/0));
} }
// Allows to register a MemRef with the CUDA runtime. Helpful until we have // Allows to register a MemRef with the CUDA runtime. Helpful until we have
// transfer functions implemented. // transfer functions implemented.
extern "C" void extern "C" void
mgpuMemHostRegisterMemRef(int64_t rank, StridedMemRefType<char, 1> *descriptor, mgpuMemHostRegisterMemRef(int64_t rank, StridedMemRefType<char, 1> *descriptor,
int64_t elementSizeBytes) { int64_t elementSizeBytes) {
Show All 18 Lines

mlir/tools/mlir-rocm-runner/rocm-runtime-wrappers.cpp

Show First 20 Lines • Show All 142 Lines • ▼ Show 20 Lines HIP_REPORT_IF_ERROR(
hipMemcpyAsync(dst, src, sizeBytes, hipMemcpyDefault, stream)); hipMemcpyAsync(dst, src, sizeBytes, hipMemcpyDefault, stream));
} }
/// Helper functions for writing mlir example code /// Helper functions for writing mlir example code
// Allows to register byte array with the ROCM runtime. Helpful until we have // Allows to register byte array with the ROCM runtime. Helpful until we have
// transfer functions implemented. // transfer functions implemented.
extern "C" void mgpuMemHostRegister(void *ptr, uint64_t sizeBytes) { extern "C" void mgpuMemHostRegister(void *ptr, uint64_t sizeBytes) {
ScopedContext scopedContext;
HIP_REPORT_IF_ERROR(hipHostRegister(ptr, sizeBytes, /*flags=*/0)); HIP_REPORT_IF_ERROR(hipHostRegister(ptr, sizeBytes, /*flags=*/0));
} }
// Allows to register a MemRef with the ROCm runtime. Helpful until we have // Allows to register a MemRef with the ROCm runtime. Helpful until we have
// transfer functions implemented. // transfer functions implemented.
extern "C" void extern "C" void
mgpuMemHostRegisterMemRef(int64_t rank, StridedMemRefType<char, 1> *descriptor, mgpuMemHostRegisterMemRef(int64_t rank, StridedMemRefType<char, 1> *descriptor,
int64_t elementSizeBytes) { int64_t elementSizeBytes) {
▲ Show 20 Lines • Show All 41 Lines • Show Last 20 Lines