"invalid device ordinal" error on kernel launch

Question

I'm trying to use multiple CUDA devices from multiple OpenMP threads. The devices are initialized (i.e. memory is allocated on them) from the main thread, and then I use cudaSetDevice from different threads to then launch kernels on different devices. Threads are not sharing devices, each thread has exclusive access to its device.

From what I understand, this should work fine. However, as soon as I launch a kernel on a device from an OpenMP thread which is the not the main (i.e. omp_get_thread_num() != 0) I get an "invalid device ordinal error" from CUDA:

kernel<<<...>>>(...);
error = cudaDeviceSynchronize(); // returns cudaSuccess
error = cudaGetLastError(); // returns invalid device ordinal error

Am I missing something? Has anyone seen something like this before? I'm using CUDA 5.0.

From my perspective, there's not enough information in this question to tell what is wrong. Yes it should work. Something you are doing prior to the kernel launch is not quite right. You may want to check which device ordinals you are actually using with each of the cudaSetDevice calls that should precede your kernel invocation in each thread. If you are trying to get the code to work with more than 2 devices, try just 2 devices first. You may want to compare with the cuda openMP sample code. — Robert Crovella
Hi Robert, thanks for the response. Yes, I'm perfectly willing to admit this is programmer error :) (though I've checked the arguments to cudaSetDevice), I was just looking to see if someone had encountered something similar. The odd thing about it to me is that I do plenty of non-kernel operations (cudaMemcpy, cudaMemset) in a thread on a device before executing the kernel but it doesn't complain for any of those. From what I understand, an invalid device ordinal indicates a problem that should be triggered by just about anything using that device, not just a kernel launch. — agrippa
agreed. If you can post a small reproducer, I'd be happy to take a look. But as I stated, without having any of the code that precedes the kernel launch to look at, I'm just guessing myself. Maybe others will have better ideas. — Robert Crovella

agrippa agrippa · Accepted Answer · 2012-11-09T17:40:25

Just to close this issue, this problem was a result of me using cudaGetLastError to try and check for errors after a kernel launch, but not checking the error return value of one previous call. Therefore, it was returning the error code from a call to cudaGetDeviceInfo after the kernel launch which I mistakenly inferred to be coming from the launch itself. If you see this error, I would just advise making sure that you're checking the error values returned by all previous calls to the CUDA API.

"invalid device ordinal" error on kernel launch

1 Answers