From ae0c1e72f694f98fadb27e35b86dce1ce868d420 Mon Sep 17 00:00:00 2001 From: zkhatami Date: Fri, 13 Jan 2023 13:19:18 -0800 Subject: [PATCH 1/6] Thrust: providing the error messages about the lack of GPU or a GPU with an incompatible architecture --- thrust/system/cuda/detail/core/util.h | 30 ++++++++++++++++++++++++++- 1 file changed, 29 insertions(+), 1 deletion(-) diff --git a/thrust/system/cuda/detail/core/util.h b/thrust/system/cuda/detail/core/util.h index 11efc0858..65a0fbd59 100644 --- a/thrust/system/cuda/detail/core/util.h +++ b/thrust/system/cuda/detail/core/util.h @@ -618,7 +618,35 @@ namespace core { inline cuda_optional get_ptx_version() { int ptx_version = 0; - cudaError_t status = cub::PtxVersion(ptx_version); + cudaError_t status = cudaGetDevice(&device); + if (status != cudaSuccess) + { + throw thrust::system_error(status, thrust::cuda_category(), "No GPU is available\n"); + } + + status = cub::PtxVersion(ptx_version); + + // Any failure means the provided device binary does not match the generated function code + if (status != cudaSuccess) + { + int major = 0, minor = 0; + cudaError_t attr_status; + + attr_status = cudaDeviceGetAttribute(&major, cudaDevAttrComputeCapabilityMajor, device); + cuda_cub::throw_on_error(attr_status, + "get_ptx_version :" + "failed to get major CUDA device compute capability version."); + + attr_status = cudaDeviceGetAttribute(&minor, cudaDevAttrComputeCapabilityMinor, device); + cuda_cub::throw_on_error(attr_status, + "get_ptx_version :" + "failed to get minor CUDA device compute capability version."); + + throw thrust::system_error(status, thrust::cuda_category(), + "Incompatible GPU: you are trying to run this program on sm_%d%d, " + "different from the one that it was compiled for\n", + major, minor); + } return cuda_optional(ptx_version, status); } From f2a1d5fd8d820fd05cb359a6c5f0bbc977cbf1c0 Mon Sep 17 00:00:00 2001 From: zkhatami Date: Fri, 13 Jan 2023 13:39:04 -0800 Subject: [PATCH 2/6] renaming device to dev_id --- thrust/system/cuda/detail/core/util.h | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/thrust/system/cuda/detail/core/util.h b/thrust/system/cuda/detail/core/util.h index 65a0fbd59..6d639f9ba 100644 --- a/thrust/system/cuda/detail/core/util.h +++ b/thrust/system/cuda/detail/core/util.h @@ -618,7 +618,8 @@ namespace core { inline cuda_optional get_ptx_version() { int ptx_version = 0; - cudaError_t status = cudaGetDevice(&device); + int dev_id = 0; + cudaError_t status = cudaGetDevice(&dev_id); if (status != cudaSuccess) { throw thrust::system_error(status, thrust::cuda_category(), "No GPU is available\n"); @@ -632,12 +633,12 @@ namespace core { int major = 0, minor = 0; cudaError_t attr_status; - attr_status = cudaDeviceGetAttribute(&major, cudaDevAttrComputeCapabilityMajor, device); + attr_status = cudaDeviceGetAttribute(&major, cudaDevAttrComputeCapabilityMajor, dev_id); cuda_cub::throw_on_error(attr_status, "get_ptx_version :" "failed to get major CUDA device compute capability version."); - attr_status = cudaDeviceGetAttribute(&minor, cudaDevAttrComputeCapabilityMinor, device); + attr_status = cudaDeviceGetAttribute(&minor, cudaDevAttrComputeCapabilityMinor, dev_id); cuda_cub::throw_on_error(attr_status, "get_ptx_version :" "failed to get minor CUDA device compute capability version."); From 78b17fe7e2ad71abf523b09cd56734c00dfd5e01 Mon Sep 17 00:00:00 2001 From: Georgy Evtushenko Date: Thu, 26 Jan 2023 12:30:48 +0400 Subject: [PATCH 3/6] Adjust error message about lack of GPU Co-authored-by: Michael Schellenberger Costa --- .../system/cuda/detail/core/agent_launcher.h | 3 +- thrust/system/cuda/detail/core/util.h | 47 +++++++++++++------ 2 files changed, 33 insertions(+), 17 deletions(-) diff --git a/thrust/system/cuda/detail/core/agent_launcher.h b/thrust/system/cuda/detail/core/agent_launcher.h index 2140c2e63..ca3656993 100644 --- a/thrust/system/cuda/detail/core/agent_launcher.h +++ b/thrust/system/cuda/detail/core/agent_launcher.h @@ -491,8 +491,7 @@ namespace core { typename core::get_plan::type static get_plan(cudaStream_t , void* d_ptr = 0) { THRUST_UNUSED_VAR(d_ptr); - core::cuda_optional ptx_version = core::get_ptx_version(); - return get_agent_plan(ptx_version); + return get_agent_plan(core::get_ptx_version()); } THRUST_RUNTIME_FUNCTION diff --git a/thrust/system/cuda/detail/core/util.h b/thrust/system/cuda/detail/core/util.h index 6d639f9ba..6e35edb16 100644 --- a/thrust/system/cuda/detail/core/util.h +++ b/thrust/system/cuda/detail/core/util.h @@ -30,6 +30,7 @@ #include #include #include +#include #include #include @@ -615,40 +616,56 @@ namespace core { }; THRUST_RUNTIME_FUNCTION - inline cuda_optional get_ptx_version() + inline int get_ptx_version() { int ptx_version = 0; - int dev_id = 0; - cudaError_t status = cudaGetDevice(&dev_id); - if (status != cudaSuccess) + const int current_device = cub::CurrentDevice(); + + if (current_device < 0) { - throw thrust::system_error(status, thrust::cuda_category(), "No GPU is available\n"); + cuda_cub::throw_on_error(cudaErrorNoDevice, "No GPU is available\n"); } - status = cub::PtxVersion(ptx_version); - // Any failure means the provided device binary does not match the generated function code - if (status != cudaSuccess) + if (cub::PtxVersion(ptx_version) != cudaSuccess) { int major = 0, minor = 0; cudaError_t attr_status; - attr_status = cudaDeviceGetAttribute(&major, cudaDevAttrComputeCapabilityMajor, dev_id); + attr_status = cudaDeviceGetAttribute(&major, cudaDevAttrComputeCapabilityMajor, current_device); cuda_cub::throw_on_error(attr_status, "get_ptx_version :" "failed to get major CUDA device compute capability version."); - attr_status = cudaDeviceGetAttribute(&minor, cudaDevAttrComputeCapabilityMinor, dev_id); + attr_status = cudaDeviceGetAttribute(&minor, cudaDevAttrComputeCapabilityMinor, current_device); cuda_cub::throw_on_error(attr_status, "get_ptx_version :" "failed to get minor CUDA device compute capability version."); + + // Index from which SM code has to start in the message below + int code_offset = 37; + char str[] = "This program was not compiled for SM \n"; + + auto print_1_helper = [&](int v) { + str[code_offset] = v + '0'; + code_offset++; + }; - throw thrust::system_error(status, thrust::cuda_category(), - "Incompatible GPU: you are trying to run this program on sm_%d%d, " - "different from the one that it was compiled for\n", - major, minor); + // Assume two digits will be enough + auto print_2_helper = [&](int v) { + if (v / 10 != 0) { + print_1_helper(v / 10); + } + print_1_helper(v % 10); + }; + + print_2_helper(major); + print_2_helper(minor); + + cuda_cub::throw_on_error(cudaErrorInvalidDevice, str); } - return cuda_optional(ptx_version, status); + + return ptx_version; } THRUST_RUNTIME_FUNCTION From 22ed101aaac95cc7c48eeffb3233edc056b8eb1b Mon Sep 17 00:00:00 2001 From: Georgy Evtushenko Date: Thu, 26 Jan 2023 16:52:31 +0400 Subject: [PATCH 4/6] Remove optional usage around ptx version --- thrust/system/cuda/detail/core/agent_launcher.h | 6 ++---- thrust/system/cuda/detail/core/util.h | 3 +-- 2 files changed, 3 insertions(+), 6 deletions(-) diff --git a/thrust/system/cuda/detail/core/agent_launcher.h b/thrust/system/cuda/detail/core/agent_launcher.h index ca3656993..dbb26f33f 100644 --- a/thrust/system/cuda/detail/core/agent_launcher.h +++ b/thrust/system/cuda/detail/core/agent_launcher.h @@ -475,9 +475,7 @@ namespace core { #ifdef __CUDACC_RDC__ return core::get_agent_plan(s, d_ptr); #else - core::cuda_optional ptx_version = core::get_ptx_version(); - //CUDA_CUB_RET_IF_FAIL(ptx_version.status()); - return get_agent_plan(ptx_version); + return get_agent_plan(core::get_ptx_version()); #endif } THRUST_RUNTIME_FUNCTION @@ -527,7 +525,7 @@ namespace core { { #if THRUST_DEBUG_SYNC_FLAG cuda_optional occ = max_sm_occupancy(k); - core::cuda_optional ptx_version = core::get_ptx_version(); + const int ptx_version = core::get_ptx_version(); if (count > 0) { _CubLog("Invoking %s<<<%u, %d, %d, %lld>>>(), %llu items total, %d items per thread, %d SM occupancy, %d vshmem size, %d ptx_version \n", diff --git a/thrust/system/cuda/detail/core/util.h b/thrust/system/cuda/detail/core/util.h index 6e35edb16..d9478c254 100644 --- a/thrust/system/cuda/detail/core/util.h +++ b/thrust/system/cuda/detail/core/util.h @@ -618,15 +618,14 @@ namespace core { THRUST_RUNTIME_FUNCTION inline int get_ptx_version() { - int ptx_version = 0; const int current_device = cub::CurrentDevice(); - if (current_device < 0) { cuda_cub::throw_on_error(cudaErrorNoDevice, "No GPU is available\n"); } // Any failure means the provided device binary does not match the generated function code + int ptx_version = 0; if (cub::PtxVersion(ptx_version) != cudaSuccess) { int major = 0, minor = 0; From 65dff885acb947e2b32e2ed8b5a7d9456cf013be Mon Sep 17 00:00:00 2001 From: Georgy Evtushenko Date: Fri, 27 Jan 2023 14:53:25 +0400 Subject: [PATCH 5/6] Reduce number of API calls in PTX check --- thrust/system/cuda/detail/core/util.h | 16 +++++++++------- 1 file changed, 9 insertions(+), 7 deletions(-) diff --git a/thrust/system/cuda/detail/core/util.h b/thrust/system/cuda/detail/core/util.h index d9478c254..39d5b743e 100644 --- a/thrust/system/cuda/detail/core/util.h +++ b/thrust/system/cuda/detail/core/util.h @@ -618,16 +618,18 @@ namespace core { THRUST_RUNTIME_FUNCTION inline int get_ptx_version() { - const int current_device = cub::CurrentDevice(); - if (current_device < 0) - { - cuda_cub::throw_on_error(cudaErrorNoDevice, "No GPU is available\n"); - } - - // Any failure means the provided device binary does not match the generated function code int ptx_version = 0; if (cub::PtxVersion(ptx_version) != cudaSuccess) { + // Failure might mean that there's no device found + const int current_device = cub::CurrentDevice(); + if (current_device < 0) + { + cuda_cub::throw_on_error(cudaErrorNoDevice, "No GPU is available\n"); + } + + // Any subsequent failure means the provided device binary does not match + // the generated function code int major = 0, minor = 0; cudaError_t attr_status; From e636580d6539a032d260a2eb04db0e3e98614678 Mon Sep 17 00:00:00 2001 From: Georgy Evtushenko Date: Sun, 29 Jan 2023 08:02:46 +0400 Subject: [PATCH 6/6] Silence MSVC int/char conversion warning --- thrust/system/cuda/detail/core/util.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thrust/system/cuda/detail/core/util.h b/thrust/system/cuda/detail/core/util.h index 39d5b743e..e2f5f8299 100644 --- a/thrust/system/cuda/detail/core/util.h +++ b/thrust/system/cuda/detail/core/util.h @@ -648,7 +648,7 @@ namespace core { char str[] = "This program was not compiled for SM \n"; auto print_1_helper = [&](int v) { - str[code_offset] = v + '0'; + str[code_offset] = static_cast(v) + '0'; code_offset++; };