Comments (12)
Thank you for reporting this @illsk1lls! We do not explicitly use noinline
anywhere. So my current best guess is it's an NVIDIA driver bug, which I guess they'll fix soon. However, if they later make a stable general release with the bug still not fixed, then we may want to do something about it from our side.
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.
from john.
If it were me, since (your) testing environment is very good, I would try patches like this:
You can safely try editing the .cl
file.
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.
from john.
Thanks for reporting.
This is a warning, not an error, john
is working and cracking for you. So, you really don't need to worry about it too much.
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 john --list=opencl-devices
output?
from john.
I only get the error with OpenCL, but for all formats and I did recently update the NVIDIA driver.
Output of: john --list=opencl-devices
Platform #0 name: NVIDIA CUDA, version: OpenCL 3.0 CUDA 12.4.125
Device #0 (1) name: NVIDIA GeForce RTX 2070
Device vendor: NVIDIA Corporation
Device type: GPU (LE)
Device version: OpenCL 3.0 CUDA
OpenCL version support: OpenCL C 1.2
Driver version: 551.86 [recommended]
Native vector widths: char 1, short 1, int 1, long 1
Preferred vector width: char 1, short 1, int 1, long 1
Global Memory: 8191 MiB
Global Memory Cache: 1152 KiB
Local Memory: 48 KiB (Local)
Constant Buffer size: 64 KiB
Max memory alloc. size: 2047 MiB
Max clock (MHz): 1440
Profiling timer res.: 1000 ns
Max Work Group Size: 1024
Parallel compute cores: 36
CUDA INT32 cores: 2304 (36 x 64)
Speed index: 3317760
Warp size: 32
Max. GPRs/work-group: 65536
Compute capability: 7.5 (sm_75)
Kernel exec. timeout: yes
PCI device topology: 01:00.0
Platform #1 name: Intel(R) OpenCL, version: OpenCL 2.1
Device #0 (2) name: Intel(R) UHD Graphics 630
Device vendor: Intel(R) Corporation
Device type: GPU (LE)
Device version: OpenCL 2.1 NEO
OpenCL version support: OpenCL C 2.0
Driver version: 26.20.100.6911
Native vector widths: char 16, short 8, int 4, long 1
Preferred vector width: char 16, short 8, int 4, long 1
Global Memory: 13047 MiB
Global Memory Cache: 512 KiB
Local Memory: 64 KiB (Local)
Constant Buffer size: 4095 MiB
Max memory alloc. size: 4095 MiB
Max clock (MHz): 1150
Profiling timer res.: 83 ns
Max Work Group Size: 256
Parallel compute cores: 24
Stream processors: 192 (24 x 8)
Speed index: 220800
Device #1 (3) name: Intel(R) Core(TM) i7-9750H CPU @ 2.60GHz
Device vendor: Intel(R) Corporation
Device type: CPU (LE)
Device version: OpenCL 2.1 (Build 0)
OpenCL version support: OpenCL C 2.0
Driver version: 7.6.0.0228
Native vector widths: char 32, short 16, int 8, long 4
Preferred vector width: char 1, short 1, int 1, long 1
Global Memory: 32617 MiB
Global Memory Cache: 256 KiB
Local Memory: 32 KiB (Global)
Constant Buffer size: 128 KiB
Max memory alloc. size: 8154 MiB
Max clock (MHz): 2600
Profiling timer res.: 100 ns
Max Work Group Size: 8192
Parallel compute cores: 12
Speed index: 124800
I am grabbing the latest Jumbo pragmatically at each run with this: https://github.com/illsk1lls/ZipRipper
John is doing it's job correctly despite the warning..
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?
from john.
@claudioandre-br I'll give it a shot, ty
@solardiz NVIDIA GameReady drivers are the current public stable release for GTX and RTX cards
from john.
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.
from john.
Another thought I have is that maybe - just maybe - this is somehow exposed by our usage of -cl-std=CL1.2
. I guess NVIDIA does more testing with their current defaults than with OpenCL spec version override options like that. If so, this does not necessarily "affects most uses of OpenCL, also by other projects."
It could be helpful to test with other OpenCL projects, or to try removing this option from ours. It's in opencl_common.c
function get_build_opts
, unfortunately probably cannot be overridden from john.conf
as we apply those options first.
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?
from john.
I have a feeling that the -cl-std=CL1.2
option is too invasive as it stands now. We can easily add a john.conf attribute to prevent it (or, on the contrary, enable it only if required).
We can't do much testing of NVIDIA anyway, so we (I) don't have much of an opinion on the matter.
from john.
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?= <[email protected]>
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é <[email protected]>
---
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.1
By all means, it's a good tool for testing.
from john.
If it were me, since (your) testing environment is very good, I would try patches like this:
You can safely try editing the file.
.cl
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.
1、kernel inline void zip_final is not work for me.
new error: 0: OpenCL CL_INVALID_KERNEL_NAME (-46) error in xxxx.c:294 - Error creating kernel
The reason is that the it cann't find right kernel name.
2、opencl_common.c => else if (options.verbosity > LOG_VERB
Yes, it works !!!
Thank you !
from john.
@claudioandre-br @magnumripper I see we define LOG_VERB
differently based on JTR_RELEASE_BUILD
. I suggest that we drop this logic - always define LOG_VERB
the way we did for JTR_RELEASE_BUILD
.
+++ 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?
from john.
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.
from john.
Related Issues (20)
- Possible to get the Core Generator project for the ZTEX source project? HOT 9
- Not ISSUE: GPG SYMMETRIC ENCRYPTION HOT 3
- Złamanie hasła HOT 1
- libreoffice2john.py: UnboundLocalError: cannot access local variable 'start_key_generation_name' HOT 12
- Broken lower-case mssql hash validation HOT 1
- PBKDF2 With salt Support HOT 4
- On recent NVIDIA drivers, OpenCL formats in 1.9.0-jumbo-1 fail with: error: passing '__generic uchar *' (aka '__generic unsigned char *') to parameter of type 'const uchar *' (aka 'const unsigned char *') changes address space of pointer HOT 2
- Dynamic memory leak when using a constant HOT 5
- GCC is again complaining about `use after [..]alloc` HOT 2
- Wrong usage of pcapngepb and pcapngpb in wpapcap2john.c HOT 5
- Omit the number 4 from the make -sj4 recommendation HOT 5
- Add Python `requirements.txt` file with pinned version numbers HOT 10
- Undefined Behavior Sanitizer "errors" HOT 12
- Support Keplr v2 wallets HOT 10
- 'make check' oddities HOT 10
- limit number of threads HOT 2
- mssql05 misses some cracks in AVX512BW builds HOT 15
- Test suite failures HOT 1
- crypt format occasionally fails a `--test-full=0 --format=cpu` run on Ubuntu 22 powerpc64le HOT 13
Recommend Projects
-
React
A declarative, efficient, and flexible JavaScript library for building user interfaces.
-
Vue.js
🖖 Vue.js is a progressive, incrementally-adoptable JavaScript framework for building UI on the web.
-
Typescript
TypeScript is a superset of JavaScript that compiles to clean JavaScript output.
-
TensorFlow
An Open Source Machine Learning Framework for Everyone
-
Django
The Web framework for perfectionists with deadlines.
-
Laravel
A PHP framework for web artisans
-
D3
Bring data to life with SVG, Canvas and HTML. 📊📈🎉
-
Recommend Topics
-
javascript
JavaScript (JS) is a lightweight interpreted programming language with first-class functions.
-
web
Some thing interesting about web. New door for the world.
-
server
A server is a program made to process requests and deliver data to clients.
-
Machine learning
Machine learning is a way of modeling and interpreting data that allows a piece of software to respond intelligently.
-
Visualization
Some thing interesting about visualization, use data art
-
Game
Some thing interesting about game, make everyone happy.
Recommend Org
-
Facebook
We are working to build community through open source technology. NB: members must have two-factor auth.
-
Microsoft
Open source projects and samples from Microsoft.
-
Google
Google ❤️ Open Source for everyone.
-
Alibaba
Alibaba Open Source for everyone
-
D3
Data-Driven Documents codes.
-
Tencent
China tencent open source team.
from john.