mirror of https://github.com/hashcat/hashcat.git
You can not select more than 25 topics
Topics must start with a letter or number, can include dashes ('-') and can be up to 35 characters long.
1482 lines
36 KiB
1482 lines
36 KiB
/**
|
|
* Author......: See docs/credits.txt
|
|
* License.....: MIT
|
|
*/
|
|
|
|
#include "common.h"
|
|
#include "types.h"
|
|
#include "memory.h"
|
|
#include "event.h"
|
|
#include "timer.h"
|
|
#include "ext_metal.h"
|
|
|
|
#include <sys/sysctl.h>
|
|
|
|
#include <CoreFoundation/CoreFoundation.h>
|
|
#include <Foundation/Foundation.h>
|
|
#include <Metal/Metal.h>
|
|
|
|
/*
|
|
typedef NS_ENUM(NSUInteger, hc_mtlFeatureSet)
|
|
{
|
|
MTL_FEATURESET_MACOS_GPUFAMILY_1_V1 = 10000,
|
|
MTL_FEATURESET_MACOS_GPUFAMILY_1_V2 = 10001,
|
|
MTL_FEATURESET_MACOS_GPUFAMILY_1_V3 = 10003,
|
|
MTL_FEATURESET_MACOS_GPUFAMILY_1_V4 = 10004,
|
|
MTL_FEATURESET_MACOS_GPUFAMILY_2_V1 = 10005,
|
|
|
|
} metalDeviceFeatureSet_macOS_t;
|
|
*/
|
|
|
|
typedef NS_ENUM(NSUInteger, hc_mtlLanguageVersion)
|
|
{
|
|
MTL_LANGUAGEVERSION_1_0 = (1 << 16),
|
|
MTL_LANGUAGEVERSION_1_1 = (1 << 16) + 1,
|
|
MTL_LANGUAGEVERSION_1_2 = (1 << 16) + 2,
|
|
MTL_LANGUAGEVERSION_2_0 = (2 << 16),
|
|
MTL_LANGUAGEVERSION_2_1 = (2 << 16) + 1,
|
|
MTL_LANGUAGEVERSION_2_2 = (2 << 16) + 2,
|
|
MTL_LANGUAGEVERSION_2_3 = (2 << 16) + 3,
|
|
MTL_LANGUAGEVERSION_2_4 = (2 << 16) + 4,
|
|
|
|
} metalLanguageVersion_t;
|
|
|
|
static bool iokit_getGPUCore (void *hashcat_ctx, int *gpu_core)
|
|
{
|
|
bool rc = false;
|
|
|
|
CFMutableDictionaryRef matching = IOServiceMatching ("IOAccelerator");
|
|
|
|
io_service_t service = IOServiceGetMatchingService (hc_IOMasterPortDefault, matching);
|
|
|
|
if (!service)
|
|
{
|
|
event_log_error (hashcat_ctx, "IOServiceGetMatchingService(): %08x", service);
|
|
|
|
return rc;
|
|
}
|
|
|
|
// "gpu-core-count" is present only on Apple Silicon
|
|
|
|
CFNumberRef num = IORegistryEntryCreateCFProperty (service, CFSTR ("gpu-core-count"), kCFAllocatorDefault, 0);
|
|
|
|
int gc = 0;
|
|
|
|
if (num == nil || CFNumberGetValue (num, kCFNumberIntType, &gc) == false)
|
|
{
|
|
//event_log_error (hashcat_ctx, "IORegistryEntryCreateCFProperty(): 'gpu-core-count' entry not found");
|
|
}
|
|
else
|
|
{
|
|
*gpu_core = gc;
|
|
|
|
rc = true;
|
|
}
|
|
|
|
IOObjectRelease (service);
|
|
|
|
return rc;
|
|
}
|
|
|
|
static int hc_mtlInvocationHelper (id target, SEL selector, void *returnValue)
|
|
{
|
|
if (target == nil) return -1;
|
|
if (selector == nil) return -1;
|
|
|
|
if ([target respondsToSelector: selector])
|
|
{
|
|
NSMethodSignature *signature = [object_getClass (target) instanceMethodSignatureForSelector: selector];
|
|
|
|
if (signature == nil) return -1;
|
|
|
|
NSInvocation *invocation = [NSInvocation invocationWithMethodSignature: signature];
|
|
|
|
if (invocation == nil) return -1;
|
|
|
|
[invocation setTarget: target];
|
|
[invocation setSelector: selector];
|
|
[invocation invoke];
|
|
[invocation getReturnValue: returnValue];
|
|
|
|
return 0;
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
static int hc_mtlBuildOptionsToDict (void *hashcat_ctx, const char *build_options_buf, const char *include_path, NSMutableDictionary *build_options_dict)
|
|
{
|
|
if (build_options_buf == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): build_options_buf is NULL", __func__);
|
|
return -1;
|
|
}
|
|
|
|
if (build_options_dict == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): build_options_dict is NULL", __func__);
|
|
return -1;
|
|
}
|
|
|
|
// NSString from build_options_buf
|
|
|
|
NSString *options = [NSString stringWithCString: build_options_buf encoding: NSUTF8StringEncoding];
|
|
|
|
if (options == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): stringWithCString failed", __func__);
|
|
return -1;
|
|
}
|
|
|
|
// replace '-D ' to ''
|
|
|
|
options = [options stringByReplacingOccurrencesOfString:@"-D " withString:@""];
|
|
|
|
if (options == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): stringByReplacingOccurrencesOfString(-D) failed", __func__);
|
|
return -1;
|
|
}
|
|
|
|
// replace '-I OpenCL ' to ''
|
|
|
|
options = [options stringByReplacingOccurrencesOfString:@"-I OpenCL " withString:@""];
|
|
|
|
if (options == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): stringByReplacingOccurrencesOfString(-I OpenCL) failed", __func__);
|
|
return -1;
|
|
}
|
|
|
|
//NSLog(@"options: '%@'", options);
|
|
|
|
// creating NSDictionary from options
|
|
|
|
NSArray *lines = [options componentsSeparatedByCharactersInSet:[NSCharacterSet whitespaceCharacterSet]];
|
|
|
|
for (NSString *aKeyValue in lines)
|
|
{
|
|
NSArray *components = [aKeyValue componentsSeparatedByString:@"="];
|
|
|
|
NSString *key = [components[0] stringByTrimmingCharactersInSet:[NSCharacterSet whitespaceCharacterSet]];
|
|
NSString *value = nil;
|
|
|
|
if ([components count] != 2)
|
|
{
|
|
if ([key isEqualToString:[NSString stringWithUTF8String:"KERNEL_STATIC"]] ||
|
|
[key isEqualToString:[NSString stringWithUTF8String:"IS_APPLE_SILICON"]] ||
|
|
[key isEqualToString:[NSString stringWithUTF8String:"DYNAMIC_LOCAL"]] ||
|
|
[key isEqualToString:[NSString stringWithUTF8String:"_unroll"]] ||
|
|
[key isEqualToString:[NSString stringWithUTF8String:"NO_UNROLL"]] ||
|
|
[key isEqualToString:[NSString stringWithUTF8String:"FORCE_DISABLE_SHM"]])
|
|
{
|
|
value = @"1";
|
|
}
|
|
else
|
|
{
|
|
#ifdef DEBUG
|
|
const char *tmp = [key UTF8String];
|
|
|
|
if (tmp != nil && strlen (tmp) > 0)
|
|
{
|
|
event_log_warning (hashcat_ctx, "%s(): skipping malformed build option: '%s'", __func__, tmp);
|
|
}
|
|
#endif
|
|
|
|
continue;
|
|
}
|
|
}
|
|
else
|
|
{
|
|
value = [components[1] stringByTrimmingCharactersInSet:[NSCharacterSet whitespaceCharacterSet]];
|
|
}
|
|
|
|
[build_options_dict setObject:value forKey:key];
|
|
}
|
|
|
|
// if set, add INCLUDE_PATH to hack Apple kernel build from source limitation on -I usage
|
|
if (include_path != nil)
|
|
{
|
|
NSString *path_key = @"INCLUDE_PATH";
|
|
NSString *path_value = [NSString stringWithCString: include_path encoding: NSUTF8StringEncoding];
|
|
// Include path may contain spaces, escape them with a backslash
|
|
path_value = [path_value stringByReplacingOccurrencesOfString:@" " withString:@"\\ "];
|
|
|
|
[build_options_dict setObject:path_value forKey:path_key];
|
|
}
|
|
|
|
//NSLog(@"Dict:\n%@", build_options_dict);
|
|
|
|
return 0;
|
|
}
|
|
|
|
int mtl_init (void *hashcat_ctx)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
MTL_PTR *mtl = (MTL_PTR *) backend_ctx->mtl;
|
|
|
|
memset (mtl, 0, sizeof (MTL_PTR));
|
|
|
|
mtl->devices = nil;
|
|
|
|
if (MTLCreateSystemDefaultDevice () == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "Metal is not supported on this computer");
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
void mtl_close (void *hashcat_ctx)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
MTL_PTR *mtl = (MTL_PTR *) backend_ctx->mtl;
|
|
|
|
if (mtl)
|
|
{
|
|
if (mtl->devices)
|
|
{
|
|
int count = (int) CFArrayGetCount (mtl->devices);
|
|
for (int i = 0; i < count; i++)
|
|
{
|
|
mtl_device_id device = (mtl_device_id) CFArrayGetValueAtIndex (mtl->devices, i);
|
|
if (device != nil)
|
|
{
|
|
hc_mtlReleaseDevice (hashcat_ctx, device);
|
|
}
|
|
}
|
|
mtl->devices = nil;
|
|
}
|
|
|
|
hcfree (backend_ctx->mtl);
|
|
|
|
backend_ctx->mtl = NULL;
|
|
}
|
|
}
|
|
|
|
int hc_mtlDeviceGetCount (void *hashcat_ctx, int *count)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
MTL_PTR *mtl = (MTL_PTR *) backend_ctx->mtl;
|
|
|
|
if (mtl == nil) return -1;
|
|
|
|
CFArrayRef devices = (CFArrayRef) MTLCopyAllDevices ();
|
|
|
|
if (devices == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "metalDeviceGetCount(): empty device objects");
|
|
|
|
return -1;
|
|
}
|
|
|
|
mtl->devices = devices;
|
|
|
|
*count = CFArrayGetCount (devices);
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_mtlDeviceGet (void *hashcat_ctx, mtl_device_id *metal_device, int ordinal)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
MTL_PTR *mtl = (MTL_PTR *) backend_ctx->mtl;
|
|
|
|
if (mtl == nil) return -1;
|
|
|
|
if (mtl->devices == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid devices pointer", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
mtl_device_id device = (mtl_device_id) CFArrayGetValueAtIndex (mtl->devices, ordinal);
|
|
|
|
if (device == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "metalDeviceGet(): invalid index");
|
|
|
|
return -1;
|
|
}
|
|
|
|
*metal_device = device;
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_mtlDeviceGetName (void *hashcat_ctx, char *name, size_t len, mtl_device_id metal_device)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
MTL_PTR *mtl = (MTL_PTR *) backend_ctx->mtl;
|
|
|
|
if (mtl == NULL) return -1;
|
|
|
|
if (metal_device == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid device", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
if (len <= 0)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): buffer length", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
id device_name_ptr = [metal_device name];
|
|
|
|
if (device_name_ptr == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): failed to get device name", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
const char *device_name_str = [device_name_ptr UTF8String];
|
|
|
|
if (device_name_str == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): failed to get UTF8String from device name", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
const size_t device_name_len = strlen (device_name_str);
|
|
|
|
if (device_name_len <= 0)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid device name length", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
if (strncpy (name, device_name_str, (device_name_len > len) ? len : device_name_len) != name)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): strncpy failed", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_mtlDeviceGetAttribute (void *hashcat_ctx, int *pi, metalDeviceAttribute_t attrib, mtl_device_id metal_device)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
MTL_PTR *mtl = (MTL_PTR *) backend_ctx->mtl;
|
|
|
|
if (mtl == NULL) return -1;
|
|
|
|
if (metal_device == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid device", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
uint64_t val64 = 0;
|
|
bool valBool = false;
|
|
unsigned long valULong = 0;
|
|
|
|
switch (attrib)
|
|
{
|
|
case MTL_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT:
|
|
// works only with Apple Silicon
|
|
if (iokit_getGPUCore (hashcat_ctx, pi) == false) *pi = 1;
|
|
break;
|
|
|
|
case MTL_DEVICE_ATTRIBUTE_UNIFIED_MEMORY:
|
|
*pi = 0;
|
|
|
|
SEL hasUnifiedMemorySelector = NSSelectorFromString (@"hasUnifiedMemory");
|
|
|
|
hc_mtlInvocationHelper (metal_device, hasUnifiedMemorySelector, &valBool);
|
|
|
|
*pi = (valBool == true) ? 1 : 0;
|
|
|
|
break;
|
|
|
|
case MTL_DEVICE_ATTRIBUTE_WARP_SIZE:
|
|
// return a fake size of 32, it will be updated later
|
|
*pi = 32;
|
|
break;
|
|
|
|
/* unused and deprecated
|
|
case MTL_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR:
|
|
*pi = 0;
|
|
|
|
if (*pi == 0 && [metal_device supportsFeatureSet:MTL_FEATURESET_MACOS_GPUFAMILY_2_V1] == true) *pi = 2;
|
|
if (*pi == 0 && [metal_device supportsFeatureSet:MTL_FEATURESET_MACOS_GPUFAMILY_1_V4] == true) *pi = 1;
|
|
if (*pi == 0 && [metal_device supportsFeatureSet:MTL_FEATURESET_MACOS_GPUFAMILY_1_V3] == true) *pi = 1;
|
|
if (*pi == 0 && [metal_device supportsFeatureSet:MTL_FEATURESET_MACOS_GPUFAMILY_1_V2] == true) *pi = 1;
|
|
if (*pi == 0 && [metal_device supportsFeatureSet:MTL_FEATURESET_MACOS_GPUFAMILY_1_V1] == true) *pi = 1;
|
|
|
|
if (*pi == 0)
|
|
{
|
|
//event_log_error (hashcat_ctx, "%s(): no feature sets supported", __func__);
|
|
return -1;
|
|
}
|
|
|
|
break;
|
|
*/
|
|
|
|
/* unused and deprecated
|
|
case MTL_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR:
|
|
*pi = 0;
|
|
|
|
if (*pi == 0 && [metal_device supportsFeatureSet:MTL_FEATURESET_MACOS_GPUFAMILY_2_V1] == true) *pi = 1;
|
|
if (*pi == 0 && [metal_device supportsFeatureSet:MTL_FEATURESET_MACOS_GPUFAMILY_1_V4] == true) *pi = 4;
|
|
if (*pi == 0 && [metal_device supportsFeatureSet:MTL_FEATURESET_MACOS_GPUFAMILY_1_V3] == true) *pi = 3;
|
|
if (*pi == 0 && [metal_device supportsFeatureSet:MTL_FEATURESET_MACOS_GPUFAMILY_1_V2] == true) *pi = 2;
|
|
if (*pi == 0 && [metal_device supportsFeatureSet:MTL_FEATURESET_MACOS_GPUFAMILY_1_V1] == true) *pi = 1;
|
|
|
|
if (*pi == 0)
|
|
{
|
|
//event_log_error (hashcat_ctx, "%s(): no feature sets supported", __func__);
|
|
return -1;
|
|
}
|
|
|
|
break;
|
|
*/
|
|
|
|
case MTL_DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK:
|
|
// M1 max is 1024
|
|
// [MTLComputePipelineState maxTotalThreadsPerThreadgroup]
|
|
*pi = 1024;
|
|
break;
|
|
|
|
case MTL_DEVICE_ATTRIBUTE_CLOCK_RATE:
|
|
// unknown
|
|
*pi = 1000000;
|
|
break;
|
|
|
|
case MTL_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK:
|
|
// 32k
|
|
*pi = 32768;
|
|
break;
|
|
|
|
case MTL_DEVICE_ATTRIBUTE_TOTAL_CONSTANT_MEMORY:
|
|
// Maximum function memory allocation for a buffer in the constant address space
|
|
// 64k
|
|
*pi = 64 * 1024;
|
|
break;
|
|
|
|
case MTL_DEVICE_ATTRIBUTE_MAX_TRANSFER_RATE:
|
|
val64 = 0;
|
|
|
|
SEL maxTransferRateSelector = NSSelectorFromString (@"maxTransferRate");
|
|
|
|
hc_mtlInvocationHelper (metal_device, maxTransferRateSelector, &val64);
|
|
|
|
*pi = (val64 == 0) ? 0 : val64 / 125; // kb/s
|
|
|
|
break;
|
|
|
|
case MTL_DEVICE_ATTRIBUTE_HEADLESS:
|
|
valBool = [metal_device isHeadless];
|
|
*pi = (valBool == true) ? 1 : 0;
|
|
break;
|
|
|
|
case MTL_DEVICE_ATTRIBUTE_LOW_POWER:
|
|
valBool = [metal_device isLowPower];
|
|
*pi = (valBool == true) ? 1 : 0;
|
|
break;
|
|
|
|
case MTL_DEVICE_ATTRIBUTE_REMOVABLE:
|
|
valBool = [metal_device isRemovable];
|
|
*pi = (valBool == true) ? 1 : 0;
|
|
break;
|
|
|
|
case MTL_DEVICE_ATTRIBUTE_REGISTRY_ID:
|
|
*pi = (int) [metal_device registryID];
|
|
break;
|
|
|
|
case MTL_DEVICE_ATTRIBUTE_PHYSICAL_LOCATION:
|
|
*pi = 0;
|
|
|
|
SEL locationSelector = NSSelectorFromString (@"location");
|
|
valULong = 0;
|
|
|
|
hc_mtlInvocationHelper (metal_device, locationSelector, &valULong);
|
|
|
|
*pi = valULong;
|
|
|
|
break;
|
|
|
|
case MTL_DEVICE_ATTRIBUTE_LOCATION_NUMBER:
|
|
*pi = 0;
|
|
|
|
SEL locationNumberSelector = NSSelectorFromString (@"locationNumber");
|
|
|
|
valULong = 0;
|
|
hc_mtlInvocationHelper (metal_device, locationNumberSelector, &valULong);
|
|
|
|
*pi = valULong;
|
|
|
|
break;
|
|
|
|
default:
|
|
event_log_error (hashcat_ctx, "%s(): unknown attribute (%d)", __func__, attrib);
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_mtlMemGetInfo (void *hashcat_ctx, size_t *mem_free, size_t *mem_total)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
MTL_PTR *mtl = (MTL_PTR *) backend_ctx->mtl;
|
|
|
|
if (mtl == NULL) return -1;
|
|
|
|
struct vm_statistics64 vm_stats;
|
|
vm_size_t page_size = 0;
|
|
unsigned int count = HOST_VM_INFO64_COUNT;
|
|
|
|
mach_port_t port = mach_host_self ();
|
|
|
|
if (host_page_size (port, &page_size) != KERN_SUCCESS)
|
|
{
|
|
event_log_error (hashcat_ctx, "metalMemGetInfo(): cannot get page_size");
|
|
|
|
return -1;
|
|
}
|
|
|
|
if (host_statistics64 (port, HOST_VM_INFO64, (host_info64_t) &vm_stats, &count) != KERN_SUCCESS)
|
|
{
|
|
event_log_error (hashcat_ctx, "metalMemGetInfo(): cannot get vm_stats");
|
|
|
|
return -1;
|
|
}
|
|
|
|
uint64_t mem_free_tmp = (uint64_t) (vm_stats.free_count - vm_stats.speculative_count) * page_size;
|
|
uint64_t mem_used_tmp = (uint64_t) (vm_stats.active_count + vm_stats.inactive_count + vm_stats.wire_count) * page_size;
|
|
|
|
*mem_free = (size_t) mem_free_tmp;
|
|
*mem_total = (size_t) (mem_free_tmp + mem_used_tmp);
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_mtlDeviceMaxMemAlloc (void *hashcat_ctx, size_t *bytes, mtl_device_id metal_device)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
MTL_PTR *mtl = (MTL_PTR *) backend_ctx->mtl;
|
|
|
|
if (mtl == NULL) return -1;
|
|
|
|
if (metal_device == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid device", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
uint64_t memsize = 0;
|
|
|
|
SEL maxBufferLengthSelector = NSSelectorFromString (@"maxBufferLength");
|
|
|
|
if (hc_mtlInvocationHelper (metal_device, maxBufferLengthSelector, &memsize) == -1) return -1;
|
|
|
|
if (memsize == 0)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid maxBufferLength", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
*bytes = (size_t) memsize;
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_mtlDeviceTotalMem (void *hashcat_ctx, size_t *bytes, mtl_device_id metal_device)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
MTL_PTR *mtl = (MTL_PTR *) backend_ctx->mtl;
|
|
|
|
if (mtl == NULL) return -1;
|
|
|
|
if (metal_device == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid device", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
uint64_t memsize = 0;
|
|
|
|
if (true)
|
|
{
|
|
memsize = [metal_device recommendedMaxWorkingSetSize];
|
|
}
|
|
else
|
|
{
|
|
size_t len = sizeof (memsize);
|
|
|
|
if (sysctlbyname ("hw.memsize", &memsize, &len, NULL, 0) != 0)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): sysctlbyname(hw.memsize) failed", __func__);
|
|
|
|
return -1;
|
|
}
|
|
}
|
|
|
|
if (memsize == 0)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid memory size", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
*bytes = (size_t) memsize;
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_mtlCreateCommandQueue (void *hashcat_ctx, mtl_device_id metal_device, mtl_command_queue *command_queue)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
MTL_PTR *mtl = (MTL_PTR *) backend_ctx->mtl;
|
|
|
|
if (mtl == NULL) return -1;
|
|
|
|
if (metal_device == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid device", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
mtl_command_queue queue = [metal_device newCommandQueue];
|
|
|
|
if (queue == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): failed to create newCommandQueue", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
*command_queue = queue;
|
|
|
|
return 0;
|
|
|
|
}
|
|
|
|
int hc_mtlCreateKernel (void *hashcat_ctx, mtl_device_id metal_device, mtl_library metal_library, const char *func_name, mtl_function *metal_function, mtl_pipeline *metal_pipeline)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
user_options_t *user_options = ((hashcat_ctx_t *) hashcat_ctx)->user_options;
|
|
|
|
MTL_PTR *mtl = (MTL_PTR *) backend_ctx->mtl;
|
|
|
|
if (mtl == NULL) return -1;
|
|
|
|
if (metal_device == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid device", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
if (metal_library == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid library", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
if (func_name == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid function name", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
__block NSError *error = nil;
|
|
|
|
NSString *f_name = [NSString stringWithCString: func_name encoding: NSUTF8StringEncoding];
|
|
|
|
mtl_function mtl_func = [metal_library newFunctionWithName: f_name];
|
|
|
|
if (mtl_func == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): failed to create '%s' function", __func__, func_name);
|
|
|
|
return -1;
|
|
}
|
|
|
|
// workaround for MTLCompilerService 'Infinite Loop' bug
|
|
|
|
/*
|
|
mtl_pipeline mtl_pipe = [metal_device newComputePipelineStateWithFunction: mtl_func error: &error];
|
|
|
|
if (error != nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): failed to create '%s' pipeline, %s", __func__, func_name, [[error localizedDescription] UTF8String]);
|
|
|
|
return -1;
|
|
}
|
|
*/
|
|
|
|
error = nil;
|
|
|
|
__block mtl_pipeline mtl_pipe;
|
|
|
|
dispatch_group_t group = dispatch_group_create ();
|
|
dispatch_queue_t queue = dispatch_get_global_queue (DISPATCH_QUEUE_PRIORITY_DEFAULT, 0);
|
|
|
|
// if no user-defined runtime, set to METAL_COMPILER_RUNTIME
|
|
long timeout = (user_options->metal_compiler_runtime > 0) ? user_options->metal_compiler_runtime : METAL_COMPILER_RUNTIME;
|
|
|
|
dispatch_time_t when = dispatch_time (DISPATCH_TIME_NOW,NSEC_PER_SEC * timeout);
|
|
|
|
__block int rc_async_err = 0;
|
|
|
|
dispatch_group_async (group, queue, ^(void)
|
|
{
|
|
mtl_pipe = [metal_device newComputePipelineStateWithFunction: mtl_func error: &error];
|
|
|
|
if (error != nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): failed to create '%s' pipeline, %s", __func__, func_name, [[error localizedDescription] UTF8String]);
|
|
|
|
rc_async_err = -1;
|
|
}
|
|
});
|
|
|
|
long rc_queue = dispatch_group_wait (group, when);
|
|
|
|
dispatch_release (group);
|
|
|
|
if (rc_async_err != 0) return -1;
|
|
|
|
if (rc_queue != 0)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): failed to create '%s' pipeline, timeout reached (status %ld)", __func__, func_name, rc_queue);
|
|
|
|
return -1;
|
|
}
|
|
|
|
if (mtl_pipe == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): failed to create '%s' pipeline", __func__, func_name);
|
|
|
|
return -1;
|
|
}
|
|
|
|
*metal_function = mtl_func;
|
|
*metal_pipeline = mtl_pipe;
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_mtlGetMaxTotalThreadsPerThreadgroup (void *hashcat_ctx, mtl_pipeline metal_pipeline, unsigned int *maxTotalThreadsPerThreadgroup)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
MTL_PTR *mtl = (MTL_PTR *) backend_ctx->mtl;
|
|
|
|
if (mtl == NULL) return -1;
|
|
|
|
if (metal_pipeline == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid pipeline", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
*maxTotalThreadsPerThreadgroup = [metal_pipeline maxTotalThreadsPerThreadgroup];
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_mtlGetThreadExecutionWidth (void *hashcat_ctx, mtl_pipeline metal_pipeline, unsigned int *threadExecutionWidth)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
MTL_PTR *mtl = (MTL_PTR *) backend_ctx->mtl;
|
|
|
|
if (mtl == NULL) return -1;
|
|
|
|
if (metal_pipeline == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid pipeline", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
*threadExecutionWidth = [metal_pipeline threadExecutionWidth];
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_mtlCreateBuffer (void *hashcat_ctx, mtl_device_id metal_device, size_t size, void *ptr, mtl_mem *metal_buffer)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
MTL_PTR *mtl = (MTL_PTR *) backend_ctx->mtl;
|
|
|
|
if (mtl == NULL) return -1;
|
|
|
|
if (metal_device == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid device", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
mtl_mem buf = NULL;
|
|
|
|
MTLResourceOptions bufferOptions = MTLResourceStorageModeShared;
|
|
|
|
if (ptr == NULL)
|
|
{
|
|
buf = [metal_device newBufferWithLength:size options:bufferOptions];
|
|
}
|
|
else
|
|
{
|
|
buf = [metal_device newBufferWithBytes:ptr length:size options:bufferOptions];
|
|
}
|
|
|
|
if (buf == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): %s failed (size: %zu)", __func__, (ptr == NULL) ? "newBufferWithLength" : "newBufferWithBytes", size);
|
|
|
|
return -1;
|
|
}
|
|
|
|
*metal_buffer = buf;
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_mtlReleaseMemObject (void *hashcat_ctx, mtl_mem metal_buffer)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
MTL_PTR *mtl = (MTL_PTR *) backend_ctx->mtl;
|
|
|
|
if (mtl == NULL) return -1;
|
|
|
|
if (metal_buffer == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid metal buffer", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
[metal_buffer setPurgeableState:MTLPurgeableStateEmpty];
|
|
[metal_buffer release];
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_mtlReleaseFunction (void *hashcat_ctx, mtl_function metal_function)
|
|
{
|
|
if (metal_function == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid metal function", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
[metal_function release];
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_mtlReleaseLibrary (void *hashcat_ctx, mtl_library metal_library)
|
|
{
|
|
if (metal_library == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid metal library", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
[metal_library release];
|
|
metal_library = nil;
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_mtlReleaseCommandQueue (void *hashcat_ctx, mtl_command_queue command_queue)
|
|
{
|
|
if (command_queue == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid metal command queue", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
[command_queue release];
|
|
command_queue = nil;
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_mtlReleaseDevice (void *hashcat_ctx, mtl_device_id metal_device)
|
|
{
|
|
if (metal_device == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid metal device", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
[metal_device release];
|
|
metal_device = nil;
|
|
|
|
return 0;
|
|
}
|
|
|
|
// device to device
|
|
|
|
int hc_mtlMemcpyDtoD (void *hashcat_ctx, mtl_command_queue command_queue, mtl_mem buf_dst, size_t buf_dst_off, mtl_mem buf_src, size_t buf_src_off, size_t buf_size)
|
|
{
|
|
if (command_queue == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): metal command queue is invalid", __func__);
|
|
return -1;
|
|
}
|
|
|
|
if (buf_src == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): metal src buffer is invalid", __func__);
|
|
return -1;
|
|
}
|
|
|
|
if (buf_src_off < 0)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): src buffer offset is invalid", __func__);
|
|
return -1;
|
|
}
|
|
|
|
if (buf_dst == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): metal dst buffer is invalid", __func__);
|
|
return -1;
|
|
}
|
|
|
|
if (buf_dst_off < 0)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): dst buffer offset is invalid", __func__);
|
|
return -1;
|
|
}
|
|
|
|
if (buf_size <= 0)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): buffer size is invalid", __func__);
|
|
return -1;
|
|
}
|
|
|
|
id<MTLCommandBuffer> command_buffer = [command_queue commandBuffer];
|
|
|
|
if (command_buffer == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): failed to create a new command buffer", __func__);
|
|
return -1;
|
|
}
|
|
|
|
id<MTLBlitCommandEncoder> blit_encoder = [command_buffer blitCommandEncoder];
|
|
|
|
if (blit_encoder == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): failed to create a blit command encoder", __func__);
|
|
return -1;
|
|
}
|
|
|
|
// copy
|
|
|
|
[blit_encoder copyFromBuffer: buf_src sourceOffset: buf_src_off toBuffer: buf_dst destinationOffset: buf_dst_off size: buf_size];
|
|
|
|
// finish encoding and start the data transfer
|
|
|
|
[blit_encoder endEncoding];
|
|
[command_buffer commit];
|
|
|
|
// Wait for complete
|
|
|
|
[command_buffer waitUntilCompleted];
|
|
|
|
return 0;
|
|
}
|
|
|
|
// host to device
|
|
|
|
int hc_mtlMemcpyHtoD (void *hashcat_ctx, mtl_command_queue command_queue, mtl_mem buf_dst, size_t buf_dst_off, const void *buf_src, size_t buf_size)
|
|
{
|
|
if (command_queue == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): metal command queue is invalid", __func__);
|
|
return -1;
|
|
}
|
|
|
|
if (buf_src == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): metal src buffer is invalid", __func__);
|
|
return -1;
|
|
}
|
|
|
|
if (buf_dst == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): host dst buffer is invalid", __func__);
|
|
return -1;
|
|
}
|
|
|
|
if (buf_size <= 0)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): buffer size is invalid", __func__);
|
|
return -1;
|
|
}
|
|
|
|
if (buf_dst_off < 0)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): buffer dst offset is invalid", __func__);
|
|
return -1;
|
|
}
|
|
|
|
void *buf_dst_ptr = [buf_dst contents];
|
|
|
|
if (buf_dst_ptr == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): failed to get metal buffer data pointer", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
if (memcpy (buf_dst_ptr + buf_dst_off, buf_src, buf_size) != buf_dst_ptr + buf_dst_off)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): memcpy failed", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
[buf_dst didModifyRange: NSMakeRange (buf_dst_off, buf_size)];
|
|
|
|
return 0;
|
|
}
|
|
|
|
// device to host
|
|
|
|
int hc_mtlMemcpyDtoH (void *hashcat_ctx, mtl_command_queue command_queue, void *buf_dst, mtl_mem buf_src, size_t buf_src_off, size_t buf_size)
|
|
{
|
|
if (command_queue == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): metal command queue is invalid", __func__);
|
|
return -1;
|
|
}
|
|
|
|
if (buf_src == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): metal src buffer is invalid", __func__);
|
|
return -1;
|
|
}
|
|
|
|
if (buf_dst == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): host dst buffer is invalid", __func__);
|
|
return -1;
|
|
}
|
|
|
|
if (buf_size <= 0)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): buffer size is invalid", __func__);
|
|
return -1;
|
|
}
|
|
|
|
id<MTLCommandBuffer> command_buffer = [command_queue commandBuffer];
|
|
|
|
if (command_buffer == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): failed to create a new command buffer", __func__);
|
|
return -1;
|
|
}
|
|
|
|
id<MTLBlitCommandEncoder> blit_encoder = [command_buffer blitCommandEncoder];
|
|
|
|
[blit_encoder synchronizeResource: buf_src];
|
|
|
|
// Finish encoding and start the data transfer to the CPU
|
|
|
|
[blit_encoder endEncoding];
|
|
[command_buffer commit];
|
|
|
|
// Wait for complete
|
|
|
|
[command_buffer waitUntilCompleted];
|
|
|
|
// get src buf ptr
|
|
|
|
void *buf_src_ptr = [buf_src contents];
|
|
|
|
if (buf_src_ptr == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): failed to get metal buffer data pointer", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
if (memcpy (buf_dst, buf_src_ptr + buf_src_off, buf_size) != buf_dst)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): memcpy failed", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_mtlRuntimeGetVersionString (void *hashcat_ctx, char *runtimeVersion_str, size_t *size)
|
|
{
|
|
CFURLRef plist_url = CFURLCreateWithFileSystemPath (kCFAllocatorDefault, CFSTR ("/System/Library/Frameworks/Metal.framework/Versions/Current/Resources/version.plist"), kCFURLPOSIXPathStyle, false);
|
|
|
|
if (plist_url == NULL)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): CFURLCreateWithFileSystemPath() failed\n", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
CFReadStreamRef plist_stream = CFReadStreamCreateWithFile (NULL, plist_url);
|
|
|
|
if (plist_stream == NULL)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): CFReadStreamCreateWithFile() failed\n", __func__);
|
|
|
|
CFRelease (plist_url);
|
|
|
|
return -1;
|
|
}
|
|
|
|
if (CFReadStreamOpen (plist_stream) == false)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): CFReadStreamOpen() failed\n", __func__);
|
|
|
|
CFRelease (plist_stream);
|
|
CFRelease (plist_url);
|
|
|
|
return -1;
|
|
}
|
|
|
|
CFPropertyListRef plist_prop = CFPropertyListCreateWithStream (NULL, plist_stream, 0, kCFPropertyListImmutable, NULL, NULL);
|
|
|
|
if (plist_prop == NULL)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): CFPropertyListCreateWithStream() failed\n", __func__);
|
|
|
|
CFReadStreamClose (plist_stream);
|
|
|
|
CFRelease (plist_stream);
|
|
CFRelease (plist_url);
|
|
return -1;
|
|
}
|
|
|
|
CFStringRef runtime_version_str = CFRetain (CFDictionaryGetValue (plist_prop, CFSTR ("CFBundleVersion")));
|
|
|
|
if (runtime_version_str != NULL)
|
|
{
|
|
if (runtimeVersion_str == NULL)
|
|
{
|
|
CFIndex len = CFStringGetLength (runtime_version_str);
|
|
CFIndex maxSize = CFStringGetMaximumSizeForEncoding (len, kCFStringEncodingUTF8) + 1;
|
|
*size = maxSize;
|
|
return 0;
|
|
}
|
|
|
|
CFIndex maxSize = *size;
|
|
|
|
if (CFStringGetCString (runtime_version_str, runtimeVersion_str, maxSize, kCFStringEncodingUTF8) == false)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): CFStringGetCString() failed\n", __func__);
|
|
|
|
hcfree (runtimeVersion_str);
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
int hc_mtlEncodeComputeCommand_pre (void *hashcat_ctx, mtl_pipeline metal_pipeline, mtl_command_queue metal_command_queue, mtl_command_buffer *metal_command_buffer, mtl_command_encoder *metal_command_encoder)
|
|
{
|
|
if (metal_pipeline == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid metal_pipeline", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
if (metal_command_queue == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid metal_command_queue", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
id<MTLCommandBuffer> metal_commandBuffer = [metal_command_queue commandBuffer];
|
|
|
|
if (metal_commandBuffer == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid metal_commandBuffer", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
id<MTLComputeCommandEncoder> metal_commandEncoder = [metal_commandBuffer computeCommandEncoder];
|
|
|
|
if (metal_commandEncoder == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid metal_commandBuffer", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
[metal_commandEncoder setComputePipelineState: metal_pipeline];
|
|
|
|
*metal_command_buffer = metal_commandBuffer;
|
|
*metal_command_encoder = metal_commandEncoder;
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_mtlSetCommandEncoderArg (void *hashcat_ctx, mtl_command_encoder metal_command_encoder, size_t off, size_t idx, mtl_mem buf, void *host_data, size_t host_data_size)
|
|
{
|
|
if (metal_command_encoder == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid metal_command_encoder", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
if (buf == nil && host_data == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid buf/host_data", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
if (buf == nil)
|
|
{
|
|
if (host_data_size <= 0)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid host_data size", __func__);
|
|
|
|
return -1;
|
|
}
|
|
}
|
|
else
|
|
{
|
|
if (off < 0)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid buf off", __func__);
|
|
|
|
return -1;
|
|
}
|
|
}
|
|
|
|
if (idx < 0)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid buf/host_data idx", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
if (host_data == nil)
|
|
{
|
|
[metal_command_encoder setBuffer: buf offset: off atIndex: idx];
|
|
}
|
|
else
|
|
{
|
|
[metal_command_encoder setBytes: host_data length: host_data_size atIndex: idx];
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_mtlEncodeComputeCommand (void *hashcat_ctx, mtl_command_encoder metal_command_encoder, mtl_command_buffer metal_command_buffer, size_t global_work_size, size_t local_work_size, double *ms)
|
|
{
|
|
MTLSize numThreadgroups = {local_work_size, 1, 1};
|
|
MTLSize threadsGroup = {global_work_size, 1, 1};
|
|
|
|
if (metal_command_encoder == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid metal_command_encoder", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
if (metal_command_buffer == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid metal_command_buffer", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
[metal_command_encoder dispatchThreadgroups: threadsGroup threadsPerThreadgroup: numThreadgroups];
|
|
|
|
[metal_command_encoder endEncoding];
|
|
[metal_command_buffer commit];
|
|
[metal_command_buffer waitUntilCompleted];
|
|
|
|
CFTimeInterval myGPUStartTime = 0;
|
|
CFTimeInterval myGPUEndTime = 0;
|
|
|
|
SEL myGPUStartTimeSelector = NSSelectorFromString (@"GPUStartTime");
|
|
SEL myGPUEndTimeSelector = NSSelectorFromString (@"GPUEndTime");
|
|
|
|
if (hc_mtlInvocationHelper (metal_command_buffer, myGPUStartTimeSelector, &myGPUStartTime) == -1) return -1;
|
|
if (hc_mtlInvocationHelper (metal_command_buffer, myGPUEndTimeSelector, &myGPUEndTime) == -1) return -1;
|
|
|
|
CFTimeInterval elapsed = myGPUEndTime - myGPUStartTime;
|
|
|
|
*ms = (1000.0 * elapsed);
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_mtlCreateLibraryWithFile (void *hashcat_ctx, mtl_device_id metal_device, const char *cached_file, mtl_library *metal_library)
|
|
{
|
|
NSError *error = nil;
|
|
|
|
if (metal_device == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid metal device", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
if (cached_file == nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): invalid metallib", __func__);
|
|
|
|
return -1;
|
|
}
|
|
|
|
NSString *k_string = [NSString stringWithCString: cached_file encoding: NSUTF8StringEncoding];
|
|
|
|
if (k_string != nil)
|
|
{
|
|
id <MTLLibrary> r = [metal_device newLibraryWithFile: k_string error: &error];
|
|
|
|
if (error != nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): failed to create metal library from metallib, %s", __func__, [[error localizedDescription] UTF8String]);
|
|
return -1;
|
|
}
|
|
|
|
*metal_library = r;
|
|
|
|
return 0;
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
int hc_mtlCreateLibraryWithSource (void *hashcat_ctx, mtl_device_id metal_device, const char *kernel_sources, const char *build_options_buf, const char *cpath, mtl_library *metal_library)
|
|
{
|
|
NSError *error = nil;
|
|
|
|
NSString *k_string = [NSString stringWithCString: kernel_sources encoding: NSUTF8StringEncoding];
|
|
|
|
if (k_string != nil)
|
|
{
|
|
MTLCompileOptions *compileOptions = [MTLCompileOptions new];
|
|
|
|
NSMutableDictionary *build_options_dict = nil;
|
|
|
|
if (build_options_buf != nil)
|
|
{
|
|
//printf ("using build_opts from arg:\n%s\n", build_options_buf);
|
|
|
|
build_options_dict = [NSMutableDictionary dictionary]; //[[NSMutableDictionary alloc] init];
|
|
|
|
if (hc_mtlBuildOptionsToDict (hashcat_ctx, build_options_buf, cpath, build_options_dict) == -1)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): failed to build options dictionary", __func__);
|
|
|
|
[build_options_dict release];
|
|
return -1;
|
|
}
|
|
|
|
compileOptions.preprocessorMacros = build_options_dict;
|
|
}
|
|
|
|
// todo: detect current os version and choose the right
|
|
// compileOptions.languageVersion = MTL_LANGUAGEVERSION_2_3;
|
|
/*
|
|
if (@available(macOS 12.0, *))
|
|
{
|
|
compileOptions.languageVersion = MTL_LANGUAGEVERSION_2_4;
|
|
}
|
|
else if (@available(macOS 11.0, *))
|
|
{
|
|
compileOptions.languageVersion = MTL_LANGUAGEVERSION_2_3;
|
|
}
|
|
else if (@available(macOS 10.15, *))
|
|
{
|
|
compileOptions.languageVersion = MTL_LANGUAGEVERSION_2_2;
|
|
}
|
|
else if (@available(macOS 10.14, *))
|
|
{
|
|
compileOptions.languageVersion = MTL_LANGUAGEVERSION_2_1;
|
|
}
|
|
else if (@available(macOS 10.13, *))
|
|
{
|
|
compileOptions.languageVersion = MTL_LANGUAGEVERSION_2_0;
|
|
}
|
|
else if (@available(macOS 10.12, *))
|
|
{
|
|
compileOptions.languageVersion = MTL_LANGUAGEVERSION_1_2;
|
|
}
|
|
else if (@available(macOS 10.11, *))
|
|
{
|
|
compileOptions.languageVersion = MTL_LANGUAGEVERSION_1_1;
|
|
}
|
|
*/
|
|
id<MTLLibrary> r = [metal_device newLibraryWithSource: k_string options: compileOptions error: &error];
|
|
|
|
[compileOptions release];
|
|
compileOptions = nil;
|
|
|
|
if (build_options_dict != nil)
|
|
{
|
|
[build_options_dict release];
|
|
build_options_dict = nil;
|
|
}
|
|
|
|
if (error != nil)
|
|
{
|
|
event_log_error (hashcat_ctx, "%s(): failed to create metal library, %s", __func__, [[error localizedDescription] UTF8String]);
|
|
|
|
return -1;
|
|
}
|
|
|
|
*metal_library = r;
|
|
|
|
return 0;
|
|
}
|
|
|
|
return -1;
|
|
}
|