diff --git a/clang/examples/DPCT/Driver/cuMemAddressFree.cu b/clang/examples/DPCT/Driver/cuMemAddressFree.cu new file mode 100644 index 000000000000..659322a9a6c6 --- /dev/null +++ b/clang/examples/DPCT/Driver/cuMemAddressFree.cu @@ -0,0 +1,7 @@ +// Option: --use-experimental-features=virtual_mem +#include +void test(CUdeviceptr ptr, size_t size) { + // Start + cuMemAddressFree(ptr /*CUdeviceptr*/, size /*size_t*/); + // End +} \ No newline at end of file diff --git a/clang/examples/DPCT/Driver/cuMemAddressReserve.cu b/clang/examples/DPCT/Driver/cuMemAddressReserve.cu new file mode 100644 index 000000000000..b87f5244a740 --- /dev/null +++ b/clang/examples/DPCT/Driver/cuMemAddressReserve.cu @@ -0,0 +1,10 @@ +// Option: --use-experimental-features=virtual_mem +#include +void test(CUdeviceptr *ptr, size_t size, size_t alignment, CUdeviceptr addr, + unsigned long long flags) { + // Start + cuMemAddressReserve(ptr /*CUdeviceptr **/, size /*size_t*/, + alignment /*size_t*/, addr /*CUdeviceptr*/, + flags /*unsigned long long*/); + // End +} \ No newline at end of file diff --git a/clang/examples/DPCT/Driver/cuMemCreate.cu b/clang/examples/DPCT/Driver/cuMemCreate.cu new file mode 100644 index 000000000000..a632ed09d23e --- /dev/null +++ b/clang/examples/DPCT/Driver/cuMemCreate.cu @@ -0,0 +1,9 @@ +// Option: --use-experimental-features=virtual_mem +#include +void test(size_t size, CUmemAllocationProp *prop, unsigned long long flags) { + // Start + CUmemGenericAllocationHandle *handle; + cuMemCreate(handle /*CUmemGenericAllocationHandle **/, size /*size_t*/, + prop /*CUmemAllocationProp **/, flags /*unsigned long long*/); + // End +} \ No newline at end of file diff --git a/clang/examples/DPCT/Driver/cuMemGetAllocationGranularity.cu b/clang/examples/DPCT/Driver/cuMemGetAllocationGranularity.cu new file mode 100644 index 000000000000..ee9b5e2a59c3 --- /dev/null +++ b/clang/examples/DPCT/Driver/cuMemGetAllocationGranularity.cu @@ -0,0 +1,12 @@ +// Option: --use-experimental-features=virtual_mem +#include +void test(size_t *granularity, CUmemAllocationProp *prop, + CUmemAllocationGranularity_flags option) { + // Start + cuMemGetAllocationGranularity(granularity /*size_t + **/ + , + prop /*CUmemAllocationProp **/, + option /*CUmemAllocationGranularity_flags*/); + // End +} \ No newline at end of file diff --git a/clang/examples/DPCT/Driver/cuMemMap.cu b/clang/examples/DPCT/Driver/cuMemMap.cu new file mode 100644 index 000000000000..2256b688be7b --- /dev/null +++ b/clang/examples/DPCT/Driver/cuMemMap.cu @@ -0,0 +1,11 @@ +// Option: --use-experimental-features=virtual_mem +#include +void test(CUdeviceptr ptr, size_t size, size_t offset, + unsigned long long flags) { + // Start + CUmemGenericAllocationHandle handle; + cuMemMap(ptr /*CUdeviceptr*/, size /*size_t*/, offset /*size_t*/, + handle /*CUmemGenericAllocationHandle*/, + flags /*unsigned long long */); + // End +} \ No newline at end of file diff --git a/clang/examples/DPCT/Driver/cuMemRelease.cu b/clang/examples/DPCT/Driver/cuMemRelease.cu new file mode 100644 index 000000000000..3f6e6093fdcb --- /dev/null +++ b/clang/examples/DPCT/Driver/cuMemRelease.cu @@ -0,0 +1,8 @@ +// Option: --use-experimental-features=virtual_mem +#include +void test() { + // Start + CUmemGenericAllocationHandle handle; + cuMemRelease(handle /*CUmemGenericAllocationHandle */); + // End +} \ No newline at end of file diff --git a/clang/examples/DPCT/Driver/cuMemSetAccess.cu b/clang/examples/DPCT/Driver/cuMemSetAccess.cu new file mode 100644 index 000000000000..2f50082136c4 --- /dev/null +++ b/clang/examples/DPCT/Driver/cuMemSetAccess.cu @@ -0,0 +1,8 @@ +// Option: --use-experimental-features=virtual_mem +#include +void test(CUdeviceptr ptr, size_t size, CUmemAccessDesc *desc, size_t count) { + // Start + cuMemSetAccess(ptr /*CUdeviceptr*/, size /*size_t*/, + desc /*CUmemAccessDesc **/, count /*size_t*/); + // End +} \ No newline at end of file diff --git a/clang/examples/DPCT/Driver/cuMemUnmap.cu b/clang/examples/DPCT/Driver/cuMemUnmap.cu new file mode 100644 index 000000000000..ec0cfad27ea4 --- /dev/null +++ b/clang/examples/DPCT/Driver/cuMemUnmap.cu @@ -0,0 +1,7 @@ +// Option: --use-experimental-features=virtual_mem +#include +void test(CUdeviceptr ptr, size_t size) { + // Start + cuMemUnmap(ptr /*CUdeviceptr*/, size /*size_t*/); + // End +} \ No newline at end of file diff --git a/clang/lib/DPCT/DPCT.cpp b/clang/lib/DPCT/DPCT.cpp index 6d0c0922ee84..807a3966c333 100644 --- a/clang/lib/DPCT/DPCT.cpp +++ b/clang/lib/DPCT/DPCT.cpp @@ -968,6 +968,8 @@ int runDPCT(int argc, const char **argv) { Experimentals.addValue(ExperimentalFeatures::Exp_LogicalGroup); else if (Option.ends_with("root-group")) Experimentals.addValue(ExperimentalFeatures::Exp_RootGroup); + else if (Option.ends_with("virtual_mem")) + Experimentals.addValue(ExperimentalFeatures::Exp_VirtualMemory); else if (Option.ends_with("masked-sub-group-operation")) Experimentals.addValue( ExperimentalFeatures::Exp_MaskedSubGroupFunction); diff --git a/clang/test/dpct/query_api_mapping/Driver/test_virtual_memory.cu b/clang/test/dpct/query_api_mapping/Driver/test_virtual_memory.cu new file mode 100644 index 000000000000..fa484ecf6e51 --- /dev/null +++ b/clang/test/dpct/query_api_mapping/Driver/test_virtual_memory.cu @@ -0,0 +1,59 @@ +// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1, cuda-10.2 +// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1, v10.2 +/// Virtual Memory Management + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuMemMap | FileCheck %s -check-prefix=cuMemMap +// cuMemMap: CUDA API: +// cuMemMap-NEXT: CUmemGenericAllocationHandle handle; +// cuMemMap-NEXT: cuMemMap(ptr /*CUdeviceptr*/, size /*size_t*/, offset /*size_t*/, +// cuMemMap-NEXT: handle /*CUmemGenericAllocationHandle*/, +// cuMemMap-NEXT: flags /*unsigned long long */); +// cuMemMap-NEXT: Is migrated to (with the option --use-experimental-features=virtual_mem): +// cuMemMap-NEXT: dpct::experimental::physical_mem_ptr handle; +// cuMemMap-NEXT: handle->map((uintptr_t)ptr, size, sycl::ext::oneapi::experimental::address_access_mode::read_write, offset); + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuMemRelease | FileCheck %s -check-prefix=cuMemRelease +// cuMemRelease: CUDA API: +// cuMemRelease-NEXT: CUmemGenericAllocationHandle handle; +// cuMemRelease-NEXT: cuMemRelease(handle /*CUmemGenericAllocationHandle */); +// cuMemRelease-NEXT: Is migrated to (with the option --use-experimental-features=virtual_mem): +// cuMemRelease-NEXT: dpct::experimental::physical_mem_ptr handle; +// cuMemRelease-NEXT: delete (handle); + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuMemUnmap | FileCheck %s -check-prefix=cuMemUnmap +// cuMemUnmap: CUDA API: +// cuMemUnmap-NEXT: cuMemUnmap(ptr /*CUdeviceptr*/, size /*size_t*/); +// cuMemUnmap-NEXT: Is migrated to (with the option --use-experimental-features=virtual_mem): +// cuMemUnmap-NEXT: sycl::ext::oneapi::experimental::unmap(ptr, size, dpct::get_current_device().get_context()); + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuMemCreate | FileCheck %s -check-prefix=cuMemCreate +// cuMemCreate: CUDA API: +// cuMemCreate-NEXT: CUmemGenericAllocationHandle *handle; +// cuMemCreate-NEXT: cuMemCreate(handle /*CUmemGenericAllocationHandle **/, size /*size_t*/, +// cuMemCreate-NEXT: prop /*CUmemAllocationProp **/, flags /*unsigned long long*/); +// cuMemCreate-NEXT: Is migrated to (with the option --use-experimental-features=virtual_mem): +// cuMemCreate-NEXT: dpct::experimental::physical_mem_ptr *handle; +// cuMemCreate-NEXT: *handle = new sycl::ext::oneapi::experimental::physical_mem(dpct::get_device(prop->location.id), dpct::get_device(prop->location.id).get_context(), size); + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuMemAddressFree | FileCheck %s -check-prefix=cuMemAddressFree +// cuMemAddressFree: CUDA API: +// cuMemAddressFree-NEXT: cuMemAddressFree(ptr /*CUdeviceptr*/, size /*size_t*/); +// cuMemAddressFree-NEXT: Is migrated to (with the option --use-experimental-features=virtual_mem): +// cuMemAddressFree-NEXT: sycl::ext::oneapi::experimental::free_virtual_mem((uintptr_t)ptr, size, dpct::get_current_device().get_context()); + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuMemSetAccess | FileCheck %s -check-prefix=cuMemSetAccess +// cuMemSetAccess: CUDA API: +// cuMemSetAccess-NEXT: cuMemSetAccess(ptr /*CUdeviceptr*/, size /*size_t*/, +// cuMemSetAccess-NEXT: desc /*CUmemAccessDesc **/, count /*size_t*/); +// cuMemSetAccess-NEXT: Is migrated to (with the option --use-experimental-features=virtual_mem): +// cuMemSetAccess-NEXT: sycl::ext::oneapi::experimental::set_access_mode(ptr, size, desc->flags, dpct::get_device(desc->location.id).get_context()); + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuMemGetAllocationGranularity | FileCheck %s -check-prefix=cuMemGetAllocationGranularity +// cuMemGetAllocationGranularity: CUDA API: +// cuMemGetAllocationGranularity-NEXT: cuMemGetAllocationGranularity(granularity /*size_t +// cuMemGetAllocationGranularity-NEXT: **/ +// cuMemGetAllocationGranularity-NEXT: , +// cuMemGetAllocationGranularity-NEXT: prop /*CUmemAllocationProp **/, +// cuMemGetAllocationGranularity-NEXT: option /*CUmemAllocationGranularity_flags*/); +// cuMemGetAllocationGranularity-NEXT: Is migrated to (with the option --use-experimental-features=virtual_mem): +// cuMemGetAllocationGranularity-NEXT: *granularity = sycl::ext::oneapi::experimental::get_mem_granularity(dpct::get_device(prop->location.id), dpct::get_device(prop->location.id).get_context(), option); diff --git a/clang/test/dpct/query_api_mapping/test_all.cu b/clang/test/dpct/query_api_mapping/test_all.cu index f39767d7c319..9297c864c0e8 100644 --- a/clang/test/dpct/query_api_mapping/test_all.cu +++ b/clang/test/dpct/query_api_mapping/test_all.cu @@ -645,20 +645,28 @@ // CHECK-NEXT: cuGetErrorString // CHECK-NEXT: cuInit // CHECK-NEXT: cuLaunchKernel +// CHECK-NEXT: cuMemAddressFree +// CHECK-NEXT: cuMemAddressReserve // CHECK-NEXT: cuMemAdvise // CHECK-NEXT: cuMemAlloc // CHECK-NEXT: cuMemAllocHost // CHECK-NEXT: cuMemAllocManaged // CHECK-NEXT: cuMemAllocPitch +// CHECK-NEXT: cuMemCreate // CHECK-NEXT: cuMemFree // CHECK-NEXT: cuMemFreeHost +// CHECK-NEXT: cuMemGetAllocationGranularity // CHECK-NEXT: cuMemGetInfo // CHECK-NEXT: cuMemHostAlloc // CHECK-NEXT: cuMemHostGetDevicePointer // CHECK-NEXT: cuMemHostGetFlags // CHECK-NEXT: cuMemHostRegister // CHECK-NEXT: cuMemHostUnregister +// CHECK-NEXT: cuMemMap // CHECK-NEXT: cuMemPrefetchAsync +// CHECK-NEXT: cuMemRelease +// CHECK-NEXT: cuMemSetAccess +// CHECK-NEXT: cuMemUnmap // CHECK-NEXT: cuMemcpy // CHECK-NEXT: cuMemcpy2D // CHECK-NEXT: cuMemcpy2DAsync