Do some CUDA and NVRTC version checks on startup

include/backend.h

@@ -38,6 +38,7 @@ int hc_nvrtcGetProgramLogSize    (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog,
 int hc_nvrtcGetProgramLog        (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, char *log);
 int hc_nvrtcGetPTXSize           (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, size_t *ptxSizeRet);
 int hc_nvrtcGetPTX               (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, char *ptx);
+int hc_nvrtcVersion              (hashcat_ctx_t *hashcat_ctx, int *major, int *minor);
 
 int hc_cuCtxCreate               (hashcat_ctx_t *hashcat_ctx, CUcontext *pctx, unsigned int flags, CUdevice dev);
 int hc_cuCtxDestroy              (hashcat_ctx_t *hashcat_ctx, CUcontext ctx);

include/ext_nvrtc.h

@@ -60,6 +60,7 @@ typedef nvrtcResult  (NVRTC_API_CALL *NVRTC_NVRTCGETPTXSIZE)         (nvrtcProgr
 typedef nvrtcResult  (NVRTC_API_CALL *NVRTC_NVRTCGETPROGRAMLOG)      (nvrtcProgram, char *);
 typedef nvrtcResult  (NVRTC_API_CALL *NVRTC_NVRTCGETPROGRAMLOGSIZE)  (nvrtcProgram, size_t *);
 typedef const char * (NVRTC_API_CALL *NVRTC_NVRTCGETERRORSTRING)     (nvrtcResult);
+typedef nvrtcResult  (NVRTC_API_CALL *NVRTC_NVRTCVERSION)            (int *, int *);
 
 typedef struct hc_nvrtc_lib
 {
@@ -75,6 +76,7 @@ typedef struct hc_nvrtc_lib
   NVRTC_NVRTCGETPROGRAMLOG      nvrtcGetProgramLog;
   NVRTC_NVRTCGETPROGRAMLOGSIZE  nvrtcGetProgramLogSize;
   NVRTC_NVRTCGETERRORSTRING     nvrtcGetErrorString;
+  NVRTC_NVRTCVERSION            nvrtcVersion;
 
 } hc_nvrtc_lib_t;
 

include/types.h

@@ -1448,6 +1448,7 @@ typedef struct backend_ctx
 
   // cuda
 
+  int                 nvrtc_driver_version;
   int                 cuda_driver_version;
 
   // opencl

src/backend.c

@@ -695,7 +695,10 @@ int nvrtc_init (hashcat_ctx_t *hashcat_ctx)
 
     char dllname[100];
 
-    for (int major = 20; major >= 3; major--) // older than 3.x do not ship _v2 functions anyway
+    for (int major = 20; major >= 10; major--) // older than 3.x do not ship _v2 functions anyway
+                                               // older than 7.x does not support sm 5.x
+                                               // older than 8.x does not have documentation archive online, no way to check if nvrtc support whatever we need
+                                               // older than 10.x is just a theoretical limit since we define 10.1 as the minimum required version
     {
       for (int minor = 20; minor >= 0; minor--)
       {
@@ -737,6 +740,7 @@ int nvrtc_init (hashcat_ctx_t *hashcat_ctx)
   HC_LOAD_FUNC (nvrtc, nvrtcGetProgramLog,      NVRTC_NVRTCGETPROGRAMLOG,     NVRTC, 1);
   HC_LOAD_FUNC (nvrtc, nvrtcGetProgramLogSize,  NVRTC_NVRTCGETPROGRAMLOGSIZE, NVRTC, 1);
   HC_LOAD_FUNC (nvrtc, nvrtcGetErrorString,     NVRTC_NVRTCGETERRORSTRING,    NVRTC, 1);
+  HC_LOAD_FUNC (nvrtc, nvrtcVersion,            NVRTC_NVRTCVERSION,           NVRTC, 1);
 
   return 0;
 }
@@ -886,6 +890,24 @@ int hc_nvrtcGetPTX (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, char *ptx)
   return 0;
 }
 
+int hc_nvrtcVersion (hashcat_ctx_t *hashcat_ctx, int *major, int *minor)
+{
+  backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
+
+  NVRTC_PTR *nvrtc = backend_ctx->nvrtc;
+
+  const nvrtcResult NVRTC_err = nvrtc->nvrtcVersion (major, minor);
+
+  if (NVRTC_err != NVRTC_SUCCESS)
+  {
+    event_log_error (hashcat_ctx, "nvrtcVersion(): %s", nvrtc->nvrtcGetErrorString (NVRTC_err));
+
+    return -1;
+  }
+
+  return 0;
+}
+
 // CUDA
 
 int cuda_init (hashcat_ctx_t *hashcat_ctx)
@@ -4983,6 +5005,19 @@ int backend_ctx_init (hashcat_ctx_t *hashcat_ctx)
 
   if ((rc_cuda_init == 0) && (rc_nvrtc_init == 0))
   {
+    // nvrtc version
+
+    int nvrtc_major = 0;
+    int nvrtc_minor = 0;
+
+    const int rc_nvrtcVersion = hc_nvrtcVersion (hashcat_ctx, &nvrtc_major, &nvrtc_minor);
+
+    if (rc_nvrtcVersion == -1) return -1;
+
+    int nvrtc_driver_version = (nvrtc_major * 1000) + (nvrtc_minor * 10);
+
+    backend_ctx->nvrtc_driver_version = nvrtc_driver_version;
+
     // cuda version
 
     int cuda_driver_version = 0;
@@ -4992,6 +5027,18 @@ int backend_ctx_init (hashcat_ctx_t *hashcat_ctx)
     if (rc_cuDriverGetVersion == -1) return -1;
 
     backend_ctx->cuda_driver_version = cuda_driver_version;
+
+    // some pre-check
+
+    if ((nvrtc_driver_version < 10010) || (cuda_driver_version < 10010))
+    {
+      event_log_error (hashcat_ctx, "Outdated NVIDIA CUDA Toolkit version '%d' detected!", cuda_driver_version);
+
+      event_log_warning (hashcat_ctx, "See hashcat.net for officially supported NVIDIA CUDA Toolkit versions.");
+      event_log_warning (hashcat_ctx, NULL);
+
+      return -1;
+    }
   }
   else
   {
@@ -6287,8 +6334,8 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime)
                 {
                   event_log_error (hashcat_ctx, "* Device #%u: Outdated or broken Intel OpenCL runtime '%s' detected!", device_id + 1, device_param->opencl_driver_version);
 
-                  event_log_warning (hashcat_ctx, "You are STRONGLY encouraged to use the officially supported NVIDIA driver.");
+                  event_log_warning (hashcat_ctx, "You are STRONGLY encouraged to use the officially supported Intel OpenCL runtime.");
-                  event_log_warning (hashcat_ctx, "See hashcat.net for officially supported NVIDIA drivers.");
+                  event_log_warning (hashcat_ctx, "See hashcat.net for officially supported Intel OpenCL runtime.");
                   event_log_warning (hashcat_ctx, "See also: https://hashcat.net/faq/wrongdriver");
                   event_log_warning (hashcat_ctx, "You can use --force to override this, but do not report related errors.");
                   event_log_warning (hashcat_ctx, NULL);
@@ -7655,16 +7702,17 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx)
 
           if (rc_nvrtcCreateProgram == -1) return -1;
 
-          char **nvrtc_options = (char **) hccalloc (3 + strlen (build_options_module_buf) + 1, sizeof (char *)); // ...
+          char **nvrtc_options = (char **) hccalloc (4 + strlen (build_options_module_buf) + 1, sizeof (char *)); // ...
 
-          nvrtc_options[0] = "--device-as-default-execution-space";
+          nvrtc_options[0] = "--restrict";
-          nvrtc_options[1] = "--gpu-architecture";
+          nvrtc_options[1] = "--device-as-default-execution-space";
+          nvrtc_options[2] = "--gpu-architecture";
 
-          hc_asprintf (&nvrtc_options[2], "compute_%d%d", device_param->sm_major, device_param->sm_minor);
+          hc_asprintf (&nvrtc_options[3], "compute_%d%d", device_param->sm_major, device_param->sm_minor);
 
           char *nvrtc_options_string = hcstrdup (build_options_module_buf);
 
-          const int num_options = 3 + nvrtc_make_options_array_from_string (nvrtc_options_string, nvrtc_options + 3);
+          const int num_options = 4 + nvrtc_make_options_array_from_string (nvrtc_options_string, nvrtc_options + 4);
 
           const int rc_nvrtcCompileProgram = hc_nvrtcCompileProgram (hashcat_ctx, program, num_options, (const char * const *) nvrtc_options);
 
@@ -7908,16 +7956,17 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx)
 
             if (rc_nvrtcCreateProgram == -1) return -1;
 
-            char **nvrtc_options = (char **) hccalloc (3 + strlen (build_options_buf) + 1, sizeof (char *)); // ...
+            char **nvrtc_options = (char **) hccalloc (4 + strlen (build_options_buf) + 1, sizeof (char *)); // ...
 
-            nvrtc_options[0] = "--device-as-default-execution-space";
+            nvrtc_options[0] = "--restrict";
-            nvrtc_options[1] = "--gpu-architecture";
+            nvrtc_options[1] = "--device-as-default-execution-space";
+            nvrtc_options[2] = "--gpu-architecture";
 
-            hc_asprintf (&nvrtc_options[2], "compute_%d%d", device_param->sm_major, device_param->sm_minor);
+            hc_asprintf (&nvrtc_options[3], "compute_%d%d", device_param->sm_major, device_param->sm_minor);
 
             char *nvrtc_options_string = hcstrdup (build_options_buf);
 
-            const int num_options = 3 + nvrtc_make_options_array_from_string (nvrtc_options_string, nvrtc_options + 3);
+            const int num_options = 4 + nvrtc_make_options_array_from_string (nvrtc_options_string, nvrtc_options + 4);
 
             const int rc_nvrtcCompileProgram = hc_nvrtcCompileProgram (hashcat_ctx, program, num_options, (const char * const *) nvrtc_options);
 
@@ -8162,16 +8211,17 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx)
 
             if (rc_nvrtcCreateProgram == -1) return -1;
 
-            char **nvrtc_options = (char **) hccalloc (3 + strlen (build_options_buf) + 1, sizeof (char *)); // ...
+            char **nvrtc_options = (char **) hccalloc (4 + strlen (build_options_buf) + 1, sizeof (char *)); // ...
 
-            nvrtc_options[0] = "--device-as-default-execution-space";
+            nvrtc_options[0] = "--restrict";
-            nvrtc_options[1] = "--gpu-architecture";
+            nvrtc_options[1] = "--device-as-default-execution-space";
+            nvrtc_options[2] = "--gpu-architecture";
 
-            hc_asprintf (&nvrtc_options[2], "compute_%d%d", device_param->sm_major, device_param->sm_minor);
+            hc_asprintf (&nvrtc_options[3], "compute_%d%d", device_param->sm_major, device_param->sm_minor);
 
             char *nvrtc_options_string = hcstrdup (build_options_buf);
 
-            const int num_options = 3 + nvrtc_make_options_array_from_string (nvrtc_options_string, nvrtc_options + 3);
+            const int num_options = 4 + nvrtc_make_options_array_from_string (nvrtc_options_string, nvrtc_options + 4);
 
             const int rc_nvrtcCompileProgram = hc_nvrtcCompileProgram (hashcat_ctx, program, num_options, (const char * const *) nvrtc_options);