diff --git a/OpenCL/inc_common.cl b/OpenCL/inc_common.cl
index a87649b3e..844e4ba12 100644
--- a/OpenCL/inc_common.cl
+++ b/OpenCL/inc_common.cl
@@ -18,7 +18,7 @@ DECLSPEC u8 v8a_from_v32_S (const u32 v32)
 
   v.v32 = v32;
 
-  return v.v8a;
+  return v.v8.a;
 }
 
 DECLSPEC u8 v8b_from_v32_S (const u32 v32)
@@ -27,7 +27,7 @@ DECLSPEC u8 v8b_from_v32_S (const u32 v32)
 
   v.v32 = v32;
 
-  return v.v8b;
+  return v.v8.b;
 }
 
 DECLSPEC u8 v8c_from_v32_S (const u32 v32)
@@ -36,7 +36,7 @@ DECLSPEC u8 v8c_from_v32_S (const u32 v32)
 
   v.v32 = v32;
 
-  return v.v8c;
+  return v.v8.c;
 }
 
 DECLSPEC u8 v8d_from_v32_S (const u32 v32)
@@ -45,7 +45,7 @@ DECLSPEC u8 v8d_from_v32_S (const u32 v32)
 
   v.v32 = v32;
 
-  return v.v8d;
+  return v.v8.d;
 }
 
 DECLSPEC u16 v16a_from_v32_S (const u32 v32)
@@ -54,7 +54,7 @@ DECLSPEC u16 v16a_from_v32_S (const u32 v32)
 
   v.v32 = v32;
 
-  return v.v16a;
+  return v.v16.a;
 }
 
 DECLSPEC u16 v16b_from_v32_S (const u32 v32)
@@ -63,15 +63,15 @@ DECLSPEC u16 v16b_from_v32_S (const u32 v32)
 
   v.v32 = v32;
 
-  return v.v16b;
+  return v.v16.b;
 }
 
 DECLSPEC u32 v32_from_v16ab_S (const u16 v16a, const u16 v16b)
 {
   vconv32_t v;
 
-  v.v16a = v16a;
-  v.v16b = v16b;
+  v.v16.a = v16a;
+  v.v16.b = v16b;
 
   return v.v32;
 }
@@ -82,7 +82,7 @@ DECLSPEC u32 v32a_from_v64_S (const u64 v64)
 
   v.v64 = v64;
 
-  return v.v32a;
+  return v.v32.a;
 }
 
 DECLSPEC u32 v32b_from_v64_S (const u64 v64)
@@ -91,15 +91,15 @@ DECLSPEC u32 v32b_from_v64_S (const u64 v64)
 
   v.v64 = v64;
 
-  return v.v32b;
+  return v.v32.b;
 }
 
 DECLSPEC u64 v64_from_v32ab_S (const u32 v32a, const u32 v32b)
 {
   vconv64_t v;
 
-  v.v32a = v32a;
-  v.v32b = v32b;
+  v.v32.a = v32a;
+  v.v32.b = v32b;
 
   return v.v64;
 }
diff --git a/OpenCL/inc_platform.cl b/OpenCL/inc_platform.cl
index 5c6d9780e..ceb12a4f1 100644
--- a/OpenCL/inc_platform.cl
+++ b/OpenCL/inc_platform.cl
@@ -8,6 +8,7 @@
 #include "inc_platform.h"
 
 #ifdef IS_NATIVE
+#define SYNC_THREADS()
 #endif
 
 #ifdef IS_CUDA
@@ -22,6 +23,11 @@ DECLSPEC u32 atomic_inc (u32 *p)
   return atomicAdd (p, 1);
 }
 
+DECLSPEC u32 atomic_or (u32 *p, u32 val)
+{
+  return atomicOr (p, val);
+}
+
 DECLSPEC size_t get_global_id  (const u32 dimindx __attribute__((unused)))
 {
   return blockDim.x * blockIdx.x + threadIdx.x;
@@ -38,7 +44,9 @@ DECLSPEC size_t get_local_size (const u32 dimindx __attribute__((unused)))
   return blockDim.x;
 }
 
+#define SYNC_THREADS() __syncthreads ()
 #endif
 
 #ifdef IS_OPENCL
+#define SYNC_THREADS() barrier (CLK_LOCAL_MEM_FENCE)
 #endif
diff --git a/OpenCL/inc_platform.h b/OpenCL/inc_platform.h
index 7d27852d9..a8ce27fef 100644
--- a/OpenCL/inc_platform.h
+++ b/OpenCL/inc_platform.h
@@ -9,6 +9,7 @@
 #ifdef IS_CUDA
 DECLSPEC u32    atomic_dec      (u32 *p);
 DECLSPEC u32    atomic_inc      (u32 *p);
+DECLSPEC u32    atomic_or       (u32 *p, u32 val);
 DECLSPEC size_t get_global_id   (const u32 dimindx __attribute__((unused)));
 DECLSPEC size_t get_local_id    (const u32 dimindx __attribute__((unused)));
 DECLSPEC size_t get_local_size  (const u32 dimindx __attribute__((unused)));
diff --git a/OpenCL/inc_types.h b/OpenCL/inc_types.h
index 0e254c24a..867bb728e 100644
--- a/OpenCL/inc_types.h
+++ b/OpenCL/inc_types.h
@@ -7,7 +7,10 @@
 #define _INC_TYPES_H
 
 #ifdef IS_CUDA
-typedef unsigned char uchar;
+typedef unsigned char  uchar;
+typedef unsigned short ushort;
+typedef unsigned int   uint;
+typedef unsigned long  ulong;
 #endif
 
 #ifdef KERNEL_STATIC
@@ -62,17 +65,19 @@ typedef union vconv32
 
   struct
   {
-    u16 v16a;
-    u16 v16b;
-  };
+    u16 a;
+    u16 b;
+
+  } v16;
 
   struct
   {
-    u8 v8a;
-    u8 v8b;
-    u8 v8c;
-    u8 v8d;
-  };
+    u8 a;
+    u8 b;
+    u8 c;
+    u8 d;
+
+  } v8;
 
 } vconv32_t;
 
@@ -82,29 +87,32 @@ typedef union vconv64
 
   struct
   {
-    u32 v32a;
-    u32 v32b;
-  };
+    u32 a;
+    u32 b;
+
+  } v32;
 
   struct
   {
-    u16 v16a;
-    u16 v16b;
-    u16 v16c;
-    u16 v16d;
-  };
+    u16 a;
+    u16 b;
+    u16 c;
+    u16 d;
+
+  } v16;
 
   struct
   {
-    u8 v8a;
-    u8 v8b;
-    u8 v8c;
-    u8 v8d;
-    u8 v8e;
-    u8 v8f;
-    u8 v8g;
-    u8 v8h;
-  };
+    u8 a;
+    u8 b;
+    u8 c;
+    u8 d;
+    u8 e;
+    u8 f;
+    u8 g;
+    u8 h;
+
+  } v8;
 
 } vconv64_t;
 
diff --git a/OpenCL/inc_vendor.h b/OpenCL/inc_vendor.h
index b2bbd9037..f8def9bd2 100644
--- a/OpenCL/inc_vendor.h
+++ b/OpenCL/inc_vendor.h
@@ -103,15 +103,11 @@
 #if defined IS_CPU
 #define DECLSPEC inline
 #elif defined IS_GPU
-#if defined IS_CUDA
-#define DECLSPEC __device__
-#else
 #if defined IS_AMD
 #define DECLSPEC inline static
 #else
 #define DECLSPEC
 #endif
-#endif
 #else
 #define DECLSPEC
 #endif
diff --git a/include/backend.h b/include/backend.h
index 9b25a3496..7141288fe 100644
--- a/include/backend.h
+++ b/include/backend.h
@@ -22,11 +22,22 @@ static const char CL_VENDOR_MESA[]            = "Mesa";
 static const char CL_VENDOR_NV[]              = "NVIDIA Corporation";
 static const char CL_VENDOR_POCL[]            = "The pocl project";
 
-int  cuda_init  (hashcat_ctx_t *hashcat_ctx);
-void cuda_close (hashcat_ctx_t *hashcat_ctx);
+int  cuda_init   (hashcat_ctx_t *hashcat_ctx);
+void cuda_close  (hashcat_ctx_t *hashcat_ctx);
 
-int  ocl_init   (hashcat_ctx_t *hashcat_ctx);
-void ocl_close  (hashcat_ctx_t *hashcat_ctx);
+int  nvrtc_init  (hashcat_ctx_t *hashcat_ctx);
+void nvrtc_close (hashcat_ctx_t *hashcat_ctx);
+
+int  ocl_init    (hashcat_ctx_t *hashcat_ctx);
+void ocl_close   (hashcat_ctx_t *hashcat_ctx);
+
+int hc_nvrtcCreateProgram        (hashcat_ctx_t *hashcat_ctx, nvrtcProgram *prog, const char *src, const char *name, int numHeaders, const char * const *headers, const char * const *includeNames);
+int hc_nvrtcDestroyProgram       (hashcat_ctx_t *hashcat_ctx, nvrtcProgram *prog);
+int hc_nvrtcCompileProgram       (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, int numOptions, const char * const *options);
+int hc_nvrtcGetProgramLogSize    (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, size_t *logSizeRet);
+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_clBuildProgram            (hashcat_ctx_t *hashcat_ctx, cl_program program, cl_uint num_devices, const cl_device_id *device_list, const char *options, void (CL_CALLBACK *pfn_notify) (cl_program program, void *user_data), void *user_data);
 int hc_clCreateBuffer            (hashcat_ctx_t *hashcat_ctx, cl_context context, cl_mem_flags flags, size_t size, void *host_ptr, cl_mem *mem);
diff --git a/include/ext_nvrtc.h b/include/ext_nvrtc.h
index 407170c16..7bbbbd15a 100644
--- a/include/ext_nvrtc.h
+++ b/include/ext_nvrtc.h
@@ -50,15 +50,16 @@ typedef struct _nvrtcProgram *nvrtcProgram;
 
 #define NVRTC_API_CALL NVRTCAPI
 
-typedef nvrtcResult (NVRTC_API_CALL *NVRTC_NVRTCADDNAMEEXPRESSION)  (nvrtcProgram, const char *);
-typedef nvrtcResult (NVRTC_API_CALL *NVRTC_NVRTCCOMPILEPROGRAM)     (nvrtcProgram, int, const char **);
-typedef nvrtcResult (NVRTC_API_CALL *NVRTC_NVRTCCREATEPROGRAM)      (nvrtcProgram *, const char *, const char *, int, const char **, const char **);
-typedef nvrtcResult (NVRTC_API_CALL *NVRTC_NVRTCDESTROYPROGRAM)     (nvrtcProgram *);
-typedef nvrtcResult (NVRTC_API_CALL *NVRTC_NVRTCGETLOWEREDNAME)     (nvrtcProgram, const char *, const char **);
-typedef nvrtcResult (NVRTC_API_CALL *NVRTC_NVRTCGETPTX)             (nvrtcProgram, char *);
-typedef nvrtcResult (NVRTC_API_CALL *NVRTC_NVRTCGETPTXSIZE)         (nvrtcProgram, size_t *);
-typedef nvrtcResult (NVRTC_API_CALL *NVRTC_NVRTCGETPROGRAMLOG)      (nvrtcProgram, char *);
-typedef nvrtcResult (NVRTC_API_CALL *NVRTC_NVRTCGETPROGRAMLOGSIZE)  (nvrtcProgram, size_t *);
+typedef nvrtcResult  (NVRTC_API_CALL *NVRTC_NVRTCADDNAMEEXPRESSION)  (nvrtcProgram, const char * const);
+typedef nvrtcResult  (NVRTC_API_CALL *NVRTC_NVRTCCOMPILEPROGRAM)     (nvrtcProgram, int, const char * const *);
+typedef nvrtcResult  (NVRTC_API_CALL *NVRTC_NVRTCCREATEPROGRAM)      (nvrtcProgram *, const char *, const char *, int, const char * const *, const char * const *);
+typedef nvrtcResult  (NVRTC_API_CALL *NVRTC_NVRTCDESTROYPROGRAM)     (nvrtcProgram *);
+typedef nvrtcResult  (NVRTC_API_CALL *NVRTC_NVRTCGETLOWEREDNAME)     (nvrtcProgram, const char * const, const char **);
+typedef nvrtcResult  (NVRTC_API_CALL *NVRTC_NVRTCGETPTX)             (nvrtcProgram, char *);
+typedef nvrtcResult  (NVRTC_API_CALL *NVRTC_NVRTCGETPTXSIZE)         (nvrtcProgram, size_t *);
+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 struct hc_nvrtc_lib
 {
@@ -73,9 +74,12 @@ typedef struct hc_nvrtc_lib
   NVRTC_NVRTCGETPTXSIZE         nvrtcGetPTXSize;
   NVRTC_NVRTCGETPROGRAMLOG      nvrtcGetProgramLog;
   NVRTC_NVRTCGETPROGRAMLOGSIZE  nvrtcGetProgramLogSize;
+  NVRTC_NVRTCGETERRORSTRING     nvrtcGetErrorString;
 
 } hc_nvrtc_lib_t;
 
 typedef hc_nvrtc_lib_t NVRTC_PTR;
 
+int nvrtc_make_options_array_from_string (char *string, char **options);
+
 #endif // _EXT_NVRTC_H
diff --git a/src/backend.c b/src/backend.c
index e7187c94d..efb118614 100644
--- a/src/backend.c
+++ b/src/backend.c
@@ -587,6 +587,7 @@ int nvrtc_init (hashcat_ctx_t *hashcat_ctx)
   HC_LOAD_FUNC (nvrtc, nvrtcGetPTXSize,         NVRTC_NVRTCGETPTXSIZE,        NVRTC, 1);
   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);
 
   return 0;
 }
@@ -831,6 +832,132 @@ void ocl_close (hashcat_ctx_t *hashcat_ctx)
   }
 }
 
+int hc_nvrtcCreateProgram (hashcat_ctx_t *hashcat_ctx, nvrtcProgram *prog, const char *src, const char *name, int numHeaders, const char * const *headers, const char * const *includeNames)
+{
+  backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
+
+  NVRTC_PTR *nvrtc = backend_ctx->nvrtc;
+
+  const nvrtcResult NVRTC_err = nvrtc->nvrtcCreateProgram (prog, src, name, numHeaders, headers, includeNames);
+
+  if (NVRTC_err != NVRTC_SUCCESS)
+  {
+    event_log_error (hashcat_ctx, "nvrtcCreateProgram(): %s", nvrtc->nvrtcGetErrorString (NVRTC_err));
+
+    return -1;
+  }
+
+  return 0;
+}
+
+int hc_nvrtcDestroyProgram (hashcat_ctx_t *hashcat_ctx, nvrtcProgram *prog)
+{
+  backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
+
+  NVRTC_PTR *nvrtc = backend_ctx->nvrtc;
+
+  const nvrtcResult NVRTC_err = nvrtc->nvrtcDestroyProgram (prog);
+
+  if (NVRTC_err != NVRTC_SUCCESS)
+  {
+    event_log_error (hashcat_ctx, "nvrtcDestroyProgram(): %s", nvrtc->nvrtcGetErrorString (NVRTC_err));
+
+    return -1;
+  }
+
+  return 0;
+}
+
+int hc_nvrtcCompileProgram (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, int numOptions, const char * const *options)
+{
+  backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
+
+  NVRTC_PTR *nvrtc = backend_ctx->nvrtc;
+
+  const nvrtcResult NVRTC_err = nvrtc->nvrtcCompileProgram (prog, numOptions, options);
+
+  if (NVRTC_err != NVRTC_SUCCESS)
+  {
+    event_log_error (hashcat_ctx, "nvrtcCompileProgram(): %s", nvrtc->nvrtcGetErrorString (NVRTC_err));
+
+    return -1;
+  }
+
+  return 0;
+}
+
+int hc_nvrtcGetProgramLogSize (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, size_t *logSizeRet)
+{
+  backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
+
+  NVRTC_PTR *nvrtc = backend_ctx->nvrtc;
+
+  const nvrtcResult NVRTC_err = nvrtc->nvrtcGetProgramLogSize (prog, logSizeRet);
+
+  if (NVRTC_err != NVRTC_SUCCESS)
+  {
+    event_log_error (hashcat_ctx, "nvrtcGetProgramLogSize(): %s", nvrtc->nvrtcGetErrorString (NVRTC_err));
+
+    return -1;
+  }
+
+  return 0;
+}
+
+int hc_nvrtcGetProgramLog (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, char *log)
+{
+  backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
+
+  NVRTC_PTR *nvrtc = backend_ctx->nvrtc;
+
+  const nvrtcResult NVRTC_err = nvrtc->nvrtcGetProgramLog (prog, log);
+
+  if (NVRTC_err != NVRTC_SUCCESS)
+  {
+    event_log_error (hashcat_ctx, "nvrtcGetProgramLog(): %s", nvrtc->nvrtcGetErrorString (NVRTC_err));
+
+    return -1;
+  }
+
+  return 0;
+}
+
+int hc_nvrtcGetPTXSize (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, size_t *ptxSizeRet)
+{
+  backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
+
+  NVRTC_PTR *nvrtc = backend_ctx->nvrtc;
+
+  const nvrtcResult NVRTC_err = nvrtc->nvrtcGetPTXSize (prog, ptxSizeRet);
+
+  if (NVRTC_err != NVRTC_SUCCESS)
+  {
+    event_log_error (hashcat_ctx, "nvrtcGetPTXSize(): %s", nvrtc->nvrtcGetErrorString (NVRTC_err));
+
+    return -1;
+  }
+
+  return 0;
+}
+
+int hc_nvrtcGetPTX (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, char *ptx)
+{
+  backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
+
+  NVRTC_PTR *nvrtc = backend_ctx->nvrtc;
+
+  const nvrtcResult NVRTC_err = nvrtc->nvrtcGetPTX (prog, ptx);
+
+  if (NVRTC_err != NVRTC_SUCCESS)
+  {
+    event_log_error (hashcat_ctx, "nvrtcGetPTX(): %s", nvrtc->nvrtcGetErrorString (NVRTC_err));
+
+    return -1;
+  }
+
+  return 0;
+}
+
 int hc_clEnqueueNDRangeKernel (hashcat_ctx_t *hashcat_ctx, cl_command_queue command_queue, cl_kernel kernel, cl_uint work_dim, const size_t *global_work_offset, const size_t *global_work_size, const size_t *local_work_size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event)
 {
   backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
@@ -3180,7 +3307,6 @@ int backend_ctx_init (hashcat_ctx_t *hashcat_ctx)
     ocl_close (hashcat_ctx);
   }
 
-
   /**
    * return if both CUDA and OpenCL initialization failed
    */
@@ -4911,6 +5037,13 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx)
 
     if (vector_width > 16) vector_width = 16;
 
+    // CUDA currently support only scalar types
+
+    if (backend_ctx->cuda)
+    {
+      vector_width = 1;
+    }
+
     device_param->vector_width = vector_width;
 
     /**
@@ -5349,65 +5482,145 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx)
 
         if (rc_read_kernel == false) return -1;
 
-        CL_rc = hc_clCreateProgramWithSource (hashcat_ctx, device_param->context, 1, (const char **) kernel_sources, NULL, &device_param->program);
+        if (backend_ctx->nvrtc)
+        {
+          nvrtcProgram program;
 
-        if (CL_rc == -1) return -1;
+          const int rc_nvrtcCreateProgram = hc_nvrtcCreateProgram (hashcat_ctx, &program, kernel_sources[0], "main_kernel", 0, NULL, NULL);
 
-        CL_rc = hc_clBuildProgram (hashcat_ctx, device_param->program, 1, &device_param->device, build_options_module_buf, NULL, NULL);
+          if (rc_nvrtcCreateProgram == -1) return -1;
 
-        //if (CL_rc == -1) return -1;
+          char **nvrtc_options = (char **) hccalloc (1 + strlen (build_options_module_buf) + 1, sizeof (char *)); // ...
 
-        size_t build_log_size = 0;
+          nvrtc_options[0] = "--device-as-default-execution-space";
 
-        hc_clGetProgramBuildInfo (hashcat_ctx, device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &build_log_size);
+          char *nvrtc_options_string = hcstrdup (build_options_module_buf);
 
-        //if (CL_rc == -1) return -1;
+          const int num_options = 1 + nvrtc_make_options_array_from_string (nvrtc_options_string, nvrtc_options + 1);
 
-        #if defined (DEBUG)
-        if ((build_log_size > 1) || (CL_rc == -1))
-        #else
-        if (CL_rc == -1)
-        #endif
-        {
-          char *build_log = (char *) hcmalloc (build_log_size + 1);
+          const int rc_nvrtcCompileProgram = hc_nvrtcCompileProgram (hashcat_ctx, program, num_options, (const char * const *) nvrtc_options);
 
-          int CL_rc_build = hc_clGetProgramBuildInfo (hashcat_ctx, device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, build_log_size, build_log, NULL);
+          size_t build_log_size = 0;
 
-          if (CL_rc_build == -1) return -1;
+          hc_nvrtcGetProgramLogSize (hashcat_ctx, program, &build_log_size);
 
-          puts (build_log);
+          #if defined (DEBUG)
+          if ((build_log_size > 1) || (rc_nvrtcCompileProgram == -1))
+          #else
+          if (rc_nvrtcCompileProgram == -1)
+          #endif
+          {
+            char *build_log = (char *) hcmalloc (build_log_size + 1);
 
-          hcfree (build_log);
-        }
+            const int rc_nvrtcGetProgramLog = hc_nvrtcGetProgramLog (hashcat_ctx, program, build_log);
 
-        if (CL_rc == -1)
-        {
-          device_param->skipped_warning = true;
+            if (rc_nvrtcGetProgramLog == -1) return -1;
 
-          event_log_error (hashcat_ctx, "* Device #%u: Kernel %s build failed - proceeding without this device.", device_id + 1, source_file);
+            puts (build_log);
 
-          continue;
+            hcfree (build_log);
+          }
+
+          if (rc_nvrtcCompileProgram == -1)
+          {
+            device_param->skipped_warning = true;
+
+            event_log_error (hashcat_ctx, "* Device #%u: Kernel %s build failed - proceeding without this device.", device_id + 1, source_file);
+
+            continue;
+          }
+
+          hcfree (nvrtc_options);
+          hcfree (nvrtc_options_string);
+
+          if (cache_disable == false)
+          {
+            size_t binary_size;
+
+            const int rc_nvrtcGetPTXSize = hc_nvrtcGetPTXSize (hashcat_ctx, program, &binary_size);
+
+            if (rc_nvrtcGetPTXSize == -1) return -1;
+
+            char *binary = (char *) hcmalloc (binary_size);
+
+            const int nvrtcGetPTX = hc_nvrtcGetPTX (hashcat_ctx, program, binary);
+
+            if (nvrtcGetPTX == -1) return -1;
+
+            const bool rc_write = write_kernel_binary (hashcat_ctx, cached_file, binary, binary_size);
+
+            if (rc_write == false) return -1;
+
+            hcfree (binary);
+          }
+
+          const int rc_nvrtcDestroyProgram = hc_nvrtcDestroyProgram (hashcat_ctx, &program);
+
+          if (rc_nvrtcDestroyProgram == -1) return -1;
         }
 
-        if (cache_disable == false)
+        if (1) // later just else
         {
-          size_t binary_size;
-
-          CL_rc = hc_clGetProgramInfo (hashcat_ctx, device_param->program, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
+          CL_rc = hc_clCreateProgramWithSource (hashcat_ctx, device_param->context, 1, (const char **) kernel_sources, NULL, &device_param->program);
 
           if (CL_rc == -1) return -1;
 
-          char *binary = (char *) hcmalloc (binary_size);
+          CL_rc = hc_clBuildProgram (hashcat_ctx, device_param->program, 1, &device_param->device, build_options_module_buf, NULL, NULL);
 
-          CL_rc = hc_clGetProgramInfo (hashcat_ctx, device_param->program, CL_PROGRAM_BINARIES, sizeof (char *), &binary, NULL);
+          //if (CL_rc == -1) return -1;
 
-          if (CL_rc == -1) return -1;
+          size_t build_log_size = 0;
+
+          hc_clGetProgramBuildInfo (hashcat_ctx, device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &build_log_size);
+
+          //if (CL_rc == -1) return -1;
+
+          #if defined (DEBUG)
+          if ((build_log_size > 1) || (CL_rc == -1))
+          #else
+          if (CL_rc == -1)
+          #endif
+          {
+            char *build_log = (char *) hcmalloc (build_log_size + 1);
+
+            int CL_rc_build = hc_clGetProgramBuildInfo (hashcat_ctx, device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, build_log_size, build_log, NULL);
+
+            if (CL_rc_build == -1) return -1;
+
+            puts (build_log);
+
+            hcfree (build_log);
+          }
+
+          if (CL_rc == -1)
+          {
+            device_param->skipped_warning = true;
+
+            event_log_error (hashcat_ctx, "* Device #%u: Kernel %s build failed - proceeding without this device.", device_id + 1, source_file);
+
+            continue;
+          }
+
+          if (cache_disable == false)
+          {
+            size_t binary_size;
+
+            CL_rc = hc_clGetProgramInfo (hashcat_ctx, device_param->program, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
+
+            if (CL_rc == -1) return -1;
+
+            char *binary = (char *) hcmalloc (binary_size);
+
+            CL_rc = hc_clGetProgramInfo (hashcat_ctx, device_param->program, CL_PROGRAM_BINARIES, sizeof (char *), &binary, NULL);
+
+            if (CL_rc == -1) return -1;
 
-          const bool rc_write = write_kernel_binary (hashcat_ctx, cached_file, binary, binary_size);
+            const bool rc_write = write_kernel_binary (hashcat_ctx, cached_file, binary, binary_size);
 
-          if (rc_write == false) return -1;
+            if (rc_write == false) return -1;
 
-          hcfree (binary);
+            hcfree (binary);
+          }
         }
       }
       else
diff --git a/src/convert.c b/src/convert.c
index c9426f617..bed989ae4 100644
--- a/src/convert.c
+++ b/src/convert.c
@@ -845,7 +845,7 @@ u8 v8a_from_v32 (const u32 v32)
 
   v.v32 = v32;
 
-  return v.v8a;
+  return v.v8.a;
 }
 
 u8 v8b_from_v32 (const u32 v32)
@@ -854,7 +854,7 @@ u8 v8b_from_v32 (const u32 v32)
 
   v.v32 = v32;
 
-  return v.v8b;
+  return v.v8.b;
 }
 
 u8 v8c_from_v32 (const u32 v32)
@@ -863,7 +863,7 @@ u8 v8c_from_v32 (const u32 v32)
 
   v.v32 = v32;
 
-  return v.v8c;
+  return v.v8.c;
 }
 
 u8 v8d_from_v32 (const u32 v32)
@@ -872,7 +872,7 @@ u8 v8d_from_v32 (const u32 v32)
 
   v.v32 = v32;
 
-  return v.v8d;
+  return v.v8.d;
 }
 
 u16 v16a_from_v32 (const u32 v32)
@@ -881,7 +881,7 @@ u16 v16a_from_v32 (const u32 v32)
 
   v.v32 = v32;
 
-  return v.v16a;
+  return v.v16.a;
 }
 
 u16 v16b_from_v32 (const u32 v32)
@@ -890,15 +890,15 @@ u16 v16b_from_v32 (const u32 v32)
 
   v.v32 = v32;
 
-  return v.v16b;
+  return v.v16.b;
 }
 
 u32 v32_from_v16ab (const u16 v16a, const u16 v16b)
 {
   vconv32_t v;
 
-  v.v16a = v16a;
-  v.v16b = v16b;
+  v.v16.a = v16a;
+  v.v16.b = v16b;
 
   return v.v32;
 }
@@ -909,7 +909,7 @@ u32 v32a_from_v64 (const u64 v64)
 
   v.v64 = v64;
 
-  return v.v32a;
+  return v.v32.a;
 }
 
 u32 v32b_from_v64 (const u64 v64)
@@ -918,15 +918,15 @@ u32 v32b_from_v64 (const u64 v64)
 
   v.v64 = v64;
 
-  return v.v32b;
+  return v.v32.b;
 }
 
 u64 v64_from_v32ab (const u32 v32a, const u32 v32b)
 {
   vconv64_t v;
 
-  v.v32a = v32a;
-  v.v32b = v32b;
+  v.v32.a = v32a;
+  v.v32.b = v32b;
 
   return v.v64;
 }
diff --git a/src/ext_nvrtc.c b/src/ext_nvrtc.c
index 17e6ff03b..634caac5d 100644
--- a/src/ext_nvrtc.c
+++ b/src/ext_nvrtc.c
@@ -6,3 +6,22 @@
 #include "common.h"
 #include "types.h"
 #include "ext_nvrtc.h"
+
+int nvrtc_make_options_array_from_string (char *string, char **options)
+{
+  char *saveptr = NULL;
+
+  char *next = strtok_r (string, " ", &saveptr);
+
+  int cnt = 0;
+
+  do
+  {
+    options[cnt] = next;
+
+    cnt++;
+
+  } while ((next = strtok_r ((char *) NULL, " ", &saveptr)) != NULL);
+
+  return cnt;
+}
\ No newline at end of file