diff options
author | Robert Maynard <robert.maynard@kitware.com> | 2017-04-25 20:09:56 (GMT) |
---|---|---|
committer | Robert Maynard <robert.maynard@kitware.com> | 2017-04-26 20:18:25 (GMT) |
commit | 493671a5212c6548b2d7376c7065f5f76692a792 (patch) | |
tree | b6e861120d50f15946603919b3637806a4e93596 /Tests/CudaOnly/ResolveDeviceSymbols/main.cu | |
parent | 8fb85c68bb0090b44df22c27dabc03da63602e5e (diff) | |
download | CMake-493671a5212c6548b2d7376c7065f5f76692a792.zip CMake-493671a5212c6548b2d7376c7065f5f76692a792.tar.gz CMake-493671a5212c6548b2d7376c7065f5f76692a792.tar.bz2 |
CUDA: Static libraries can now explicitly resolve device symbols
If a static library has the property CUDA_RESOLVE_DEVICE_SYMBOLS enabled
it will now perform the device link step. The normal behavior is
to delay calling device link until the static library is consumed by
a shared library or an executable.
Diffstat (limited to 'Tests/CudaOnly/ResolveDeviceSymbols/main.cu')
-rw-r--r-- | Tests/CudaOnly/ResolveDeviceSymbols/main.cu | 85 |
1 files changed, 85 insertions, 0 deletions
diff --git a/Tests/CudaOnly/ResolveDeviceSymbols/main.cu b/Tests/CudaOnly/ResolveDeviceSymbols/main.cu new file mode 100644 index 0000000..b4b5b9e --- /dev/null +++ b/Tests/CudaOnly/ResolveDeviceSymbols/main.cu @@ -0,0 +1,85 @@ + +#include <iostream> + +#include "file1.h" +#include "file2.h" + +int file2_launch_kernel(int x); + +result_type_dynamic __device__ file2_func(int x); +static __global__ void main_kernel(result_type_dynamic& r, int x) +{ + // call function that was not device linked to us, this will cause + // a runtime failure of "invalid device function" + r = file2_func(x); +} + +int main_launch_kernel(int x) +{ + result_type_dynamic r; + main_kernel<<<1, 1>>>(r, x); + return r.sum; +} + +int choose_cuda_device() +{ + int nDevices = 0; + cudaError_t err = cudaGetDeviceCount(&nDevices); + if (err != cudaSuccess) { + std::cerr << "Failed to retrieve the number of CUDA enabled devices" + << std::endl; + return 1; + } + for (int i = 0; i < nDevices; ++i) { + cudaDeviceProp prop; + cudaError_t err = cudaGetDeviceProperties(&prop, i); + if (err != cudaSuccess) { + std::cerr << "Could not retrieve properties from CUDA device " << i + << std::endl; + return 1; + } + std::cout << "prop.major: " << prop.major << std::endl; + if (prop.major >= 3) { + err = cudaSetDevice(i); + if (err != cudaSuccess) { + std::cout << "Could not select CUDA device " << i << std::endl; + } else { + return 0; + } + } + } + + std::cout << "Could not find a CUDA enabled card supporting compute >=3.0" + << std::endl; + + return 1; +} + +int main(int argc, char** argv) +{ + int ret = choose_cuda_device(); + if (ret) { + return 0; + } + + cudaError_t err; + file2_launch_kernel(42); + err = cudaGetLastError(); + if (err != cudaSuccess) { + std::cerr << "file2_launch_kernel: kernel launch failed: " + << cudaGetErrorString(err) << std::endl; + return 1; + } + + main_launch_kernel(1); + err = cudaGetLastError(); + if (err == cudaSuccess) { + // This kernel launch should fail as the file2_func was device linked + // into the static library and is not usable by the executable + std::cerr << "main_launch_kernel: kernel launch should have failed" + << std::endl; + return 1; + } + + return 0; +} |