NVIDIA driver 551.86 new "Warning: Function [...] is a kernel, so overriding noinline attribute" #5456
Comments
illsk1lls
changed the title
|
Thanks for reporting. This is a warning, not an error, Anyway, I haven't found anything recent that could be causing this, so the root cause is probably related to a driver update (or your previous version of John was kind of out of date). Can you see the same error when using other formats? Could you also add the |
|
I only get the error with OpenCL, but for all formats and I did recently update the NVIDIA driver. Output of: I am grabbing the latest Jumbo pragmatically at each run with this: https://github.com/illsk1lls/ZipRipper Is there a way to suppress that specific warning and still get other output? verbosity=2 suppresses it but also the password output, 2>nul suppresses it, but that suppresses all status output because they also use the STDERR stream. I didn't think of the driver, that's actually extremely likely. Most importantly, what does the warning mean? |
solardiz
changed the title
|
Thank you for reporting this @illsk1lls! We do not explicitly use I notice they released 551.86 as "Game Ready". I don't know what exactly this means, but it's probably kind of a public alpha or beta version, for public testing. They have a feedback thread here, which you may add to - https://www.nvidia.com/en-us/geforce/forums/game-ready-drivers/13/539523/geforce-grd-55186-feedback-thread-released-31924/ - although I suspect they mostly care about feedback on gaming issues there, given the target audience and specific features of this release. |
|
If it were me, since (your) testing environment is very good, I would try patches like this: You can safely try editing the diff --git a/run/opencl/zip_kernel.cl b/run/opencl/zip_kernel.cl
index d181eb2f3e..40b570eb35 100644
--- a/run/opencl/zip_kernel.cl
+++ b/run/opencl/zip_kernel.cl
@@ -238,7 +238,7 @@ inline uint prepare(__global const uchar *pwbuf, __global const uint *buf_idx, u
#define ITERATIONS 1000 /* salt->iterations */
-__kernel void zip(__global const uchar *pwbuf,
+__kernel inline void zip(__global const uchar *pwbuf,
__global const uint *buf_idx,
__constant zip_salt *salt,
volatile __global uint *crack_count_ret,
@@ -264,7 +264,7 @@ __kernel void zip(__global const uchar *pwbuf,
}
}
-kernel void zip_final(__global const uchar *pwbuf,
+kernel inline void zip_final(__global const uchar *pwbuf,
__global const uint *buf_idx,
__constant zip_salt *salt,
__global const uchar *saltdata,
Or (needs a new build): diff --git a/src/opencl_common.c b/src/opencl_common.c
index 0370c649cd..4f9c08f822 100644
--- a/src/opencl_common.c
+++ b/src/opencl_common.c
@@ -1297,7 +1297,7 @@ void opencl_build(int sequential_id, const char *opts, int save, const char *fil
HANDLE_CLERROR(build_code, "clBuildProgram");
}
// Nvidia may return a single '\n' that we ignore
- else if (options.verbosity >= LOG_VERB && strlen(build_log) > 1)
+ else if (options.verbosity > LOG_VERB && strlen(build_log) > 1)
fprintf(stderr, "Build log: %s\n", build_log);
MEM_FREE(build_log);
But the first needs to be well tested and the second is probably not the best choice. |
|
@claudioandre-br I'll give it a shot, ty @solardiz NVIDIA GameReady drivers are the current public stable release for GTX and RTX cards |
OK, thanks @illsk1lls. Then it's more of an issue. Did you install just this driver or also separately install CUDA? If you did not separately install CUDA, then can you please report this driver issue to NVIDIA, such as in the feedback thread I found? Chances are it affects most uses of OpenCL, also by other projects. |
|
Another thought I have is that maybe - just maybe - this is somehow exposed by our usage of It could be helpful to test with other OpenCL projects, or to try removing this option from ours. It's in Maybe that options order is something for us to change (regardless of this specific issue) - apply externally provided options last, not first. What do you think @claudioandre-br @magnumripper? When the same option (with different parameter values) is specified more than once, compilers typically use the last one - but I don't know if that's also guaranteed or at least is typical with OpenCL backends, is it? |
|
I have a feeling that the We can't do much testing of NVIDIA anyway, so we (I) don't have much of an opinion on the matter. |
|
It's simple as we can see below, but what will happen to the "array arguments" in the new driver version when we disable it, and what effect will the OpenCL 1.2 headers have on these new versions? From ce9493f558c81c1900c653f08eeaba05e2edc20f Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Claudio=20Andr=C3=A9?= <dev@claudioandre.slmail.me>
Date: Mon, 1 Apr 2024 14:56:14 -0300
Subject: [PATCH] OpenCL: Add a flag to control when to use OpenCL 1.2
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit
Signed-off-by: Claudio André <dev@claudioandre.slmail.me>
---
doc/CHANGES | 2 ++
doc/README-OPENCL | 3 +++
run/john.conf | 5 +++++
src/opencl_common.c | 9 ++++++++-
4 files changed, 18 insertions(+), 1 deletion(-)
diff --git a/doc/CHANGES b/doc/CHANGES
index 8ff614bd85..4383f13a55 100644
--- a/doc/CHANGES
+++ b/doc/CHANGES
@@ -383,3 +383,5 @@ OpenBSD/PowerPC, OpenBSD/PA-RISC, OpenBSD/VAX, NetBSD/VAX, Solaris/SPARC64,
Mac OS X (PowerPC and x86), SCO, BeOS.
* Bug and portability fixes, and new bugs.
* Bonus: "Strip" cracker included in the default john.conf (john.ini).
+* Add a global bool "ForceOpenCLVersion" in john.conf that requires OpenCL
+compilers to use the OpenCL version 1.2 standard in the compilation process.
diff --git a/doc/README-OPENCL b/doc/README-OPENCL
index 6ad572ba88..4c947d2911 100644
--- a/doc/README-OPENCL
+++ b/doc/README-OPENCL
@@ -222,3 +222,6 @@ achieve that is "make kernel-cache-clean"
AMD drivers are hopeless. Every bugfix implements two new bugs. Try using
the latest and greatest. nvidia drivers are a whole lot more stable but if
seeing problems, first thing to try is updating drivers.
+
+You can set ForceOpenCLVersion=Y in john.conf to force the compiler to use
+OpenCL 1.2 standards in the OpenCL compilation process.
\ No newline at end of file
diff --git a/run/john.conf b/run/john.conf
index d26c3558c4..7db7b522ca 100644
--- a/run/john.conf
+++ b/run/john.conf
@@ -457,6 +457,11 @@ ForceScalar = N
# concatenated to this.
GlobalBuildOpts = -cl-mad-enable
+# Global build option to use OpenCL 1.2 standard.
+# Some recent nvidia drivers complains about array arguments in kernel functions
+# unless explicitly reverting to OpenCL 1.2. Add -cl-std=CL1.2 when applicable.
+ForceOpenCLVersion = N
+
# Initial local work-size for auto-tune (CPU devices excepted).
# 0 means let the OpenCL implementation pick a suitable value.
# 1 means query for "best multiple" (usually corresponds to "warp size").
diff --git a/src/opencl_common.c b/src/opencl_common.c
index 0370c649cd..d93c7ccbfd 100644
--- a/src/opencl_common.c
+++ b/src/opencl_common.c
@@ -1137,6 +1137,13 @@ static void print_device_info(int sequential_id)
log_event("Device %d: %s%s", sequential_id + 1, device_name, board_name);
}
+static char *get_force_opencl_version(int sequential_id)
+{
+ if (cfg_get_bool(SECTION_OPTIONS, SUBSECTION_OPENCL, "ForceOpenCLVersion", 0))
+ return get_device_version(sequential_id) >= 200 ? "-cl-std=CL1.2 " : "";
+ return "";
+}
+
static char *get_build_opts(int sequential_id, const char *opts)
{
char *build_opts = mem_alloc(LINE_BUFFER_SIZE);
@@ -1150,7 +1157,7 @@ static char *get_build_opts(int sequential_id, const char *opts)
snprintf(build_opts, LINE_BUFFER_SIZE,
"-I opencl %s %s%s%s%s%s%d %s%d %s -D_OPENCL_COMPILER %s",
global_opts,
- get_device_version(sequential_id) >= 200 ? "-cl-std=CL1.2 " : "",
+ get_force_opencl_version(sequential_id),
#ifdef __APPLE__
"-D__OS_X__ ",
#else
--
2.40.1By all means, it's a good tool for testing. |
1、kernel inline void zip_final is not work for me. 2、opencl_common.c => else if (options.verbosity > LOG_VERB Thank you ! |
|
@claudioandre-br @magnumripper I see we define +++ b/src/opencl_common.c
@@ -61,13 +61,9 @@
#define LOG_SIZE 1024*16
-// If we are a release build, only output OpenCL build log if
-// there was a fatal error (or --verbosity was increased).
-#ifdef JTR_RELEASE_BUILD
+/* Only output OpenCL build log if there was a fatal error
+ * (or --verbosity was increased). */
#define LOG_VERB VERB_LEGACY
-#else
-#define LOG_VERB VERB_DEFAULT
-#endif
/* Common OpenCL variables */
int platform_id;
diff --git a/src/params.h b/src/params.h
index 83afed9eb..74726debd 100644
--- a/src/params.h
+++ b/src/params.h
@@ -30,9 +30,7 @@
/*
* Define this for release tarballs. It affects the version reporting (will
- * be the string above and below and never a Git hash) as well as some other
- * details. Eg. it mutes output of OpenCL run-time build log unless the build
- * failed.
+ * be the string above and below and never a Git hash).
*/
//#define JTR_RELEASE_BUILD 1
OK with you? |
Since bleeding is no longer a place for experimentation, that's okay. In fact, this even helps with a possible policy change to allow frequent (source only) releases. It's frustrating to see people using and complaining about old software. |
Getting an error displaying on screen on the latest Win11 with jumbo when working on 7z, zip (tested)
Win11 - 10.0.22631.3296
Error text:
Build log: (): Warning: Function zip is a kernel, so overriding noinline attribute. The function may be inlined when called. (): Warning: Function zip_final is a kernel, so overriding noinline attribute. The function may be inlined when called.Example:
Since the output for john.exe is using stderr instead of stdout for job status info, this information cannot be suppressed.
Command used:
john --wordlist="password.lst" --rules=single,all --format=ZIP-opencl "hashfile"Hashfile contents:
winzip-AES-256-testpassword#.zip/test.txt:$zip2$*0*3*0*e3bd6c1a4c4950d0c35c1b0ca2bd5e84*061f*4*9c276325*38408eb4a32fe6220bb6*$/zip2$:test.txt:winzip-AES-256-testpassword#.zip:C:\john-samples-main\ZIP\winzip-AES-256-testpassword#.zipThe text was updated successfully, but these errors were encountered: