[CUDA] SYCL/USM/dep_events.cpp from llvm-test-suite fails with PI CUDA ERROR in plugins/cuda/pi_cuda.cpp:4674 on Windows

[yuxianch]

Describe the bug
SYCL/USM/dep_events.cpp from https://github.com/intel/llvm-test-suite fails with PI CUDA ERROR in plugins/cuda/pi_cuda.cpp:4674 with CUDA backend on Windows.

To Reproduce

1. Test source code can be found in https://github.com/intel/llvm-test-suite/blob/intel/SYCL/USM/dep_events.cpp
2. The command which should be used to compile the program and launch the program

clang-cl.exe /EHsc -fsycl -fsycl-targets=nvptx64-nvidia-cuda SYCL\USM\dep_events.cpp -o build\SYCL\USM\Output\dep_events.cpp.tmp1.out
env SYCL_DEVICE_FILTER=cuda:gpu,host build\SYCL\USM\Output\dep_events.cpp.tmp1.out

Output of the program:

PI CUDA ERROR:
        Value:           101
        Name:            CUDA_ERROR_INVALID_DEVICE
        Description:     invalid device ordinal
        Function:        cuda_piextUSMEnqueuePrefetch
        Source Location: ${compiler_build_path}\llvm.src\sycl\plugins\cuda\pi_cuda.cpp:4674

Environment

• OS: Windows
• Target device and vendor: Nvidia GPU
• DPC++ version: https://github.com/intel/llvm.git 3ef8c82
• CUDA version:

Cuda compilation tools, release 11.4, V11.4.120
Build cuda_11.4.r11.4/compiler.30300941_0

[yuxianch]

SYCL/USM/prefetch.cpp gets the same output when running with CUDA backend on Windows.

[AidanBeltonS]

I have been able to replicate this issue for both tests.
It comes from cuda_piextUSMEnqueuePrefetch in both situations.
This appears to be an issue with CUDA on windows machines. It can be reproduced without dpc++
I have filed a bug with nvidia.

To reproduce run code below with nvcc.

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
int main()
{
	int* data;
	size_t len = 10;
	cudaError_t err = cudaMallocManaged(reinterpret_cast<void **>(&data), len, cudaMemAttachGlobal);
	err = cudaMemPrefetchAsync(data, len, 0, 0);
	if (err != 0) {
        printf(cudaGetErrorName(err)); 
        }
}

Result: cudaErrorInvalidDevice

Note the code above is modified from https://forums.developer.nvidia.com/t/cudamemprefetchasync-returns-cudaerrorinvaliddevice/46705.

This was ran for a GeForce GTX 1660 GPU with Windows 10 and CUDA 11.4.
I propose setting these tests to unsupported for CUDA with Windows. intel/llvm-test-suite#551

[AerialMantis]

This issue is now awaiting a 3rd party update, a fix in CUDA for cudaMallocManaged on Windows.

[AidanBeltonS]

NVIDIA have responded to the bug report. It appears that this is not a bug and that it is a known limitation of CUDA.

Thanks for reporting this ticket .After checking with our engineers , we think this is expected behavior , please see https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY.html
Passing in cudaCpuDeviceId for dstDevice will prefetch the data to host memory. If dstDevice is a GPU, then the device attribute cudaDevAttrConcurrentManagedAccess must be non-zero. Additionally, stream must be associated with a device that has a non-zero value for the device attribute cudaDevAttrConcurrentManagedAccess.

Doing cudaMemPrefetchAsync on managed memory requires support for "cudaDevAttrConcurrentManagedAccess".

Currently, this is only supported on Linux. You are suggested to check for the above device attr before doing mem prefetch on managed mem. like
int* data;
size_t len = 10;
int featureSupported = 0;

CHECK_RT(cudaMallocManaged(reinterpret_cast(&data), len, cudaMemAttachGlobal));
CHECK_RT(cudaDeviceGetAttribute(&featureSupported, cudaDevAttrConcurrentManagedAccess, 0));

if (featureSupported) {
CHECK_RT(cudaMemPrefetchAsync(data, len, 0, 0));
}

[AerialMantis]

There's a pull request for introducing a check in cuda_piextUSMEnqueuePrefetch for CU_DEVICE_ATTRIBUTE_CONCURRENT_MANAGED_ACCESS which will address this - #5043