Skip to content

Commit d825187

Browse files
bensuperpcbensuperpc
committed
Fix
Signed-off-by: Bensuperpc <[email protected]>
1 parent bb04584 commit d825187

File tree

10 files changed

+252
-187
lines changed

10 files changed

+252
-187
lines changed

‎CMakePresets.json‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -54,7 +54,7 @@
5454
"CMAKE_CXX_STANDARD": "20",
5555
"CMAKE_CXX_STANDARD_REQUIRED": "ON",
5656
"CMAKE_C_EXTENSIONS": "ON",
57-
"CMAKE_C_STANDARD": "11",
57+
"CMAKE_C_STANDARD": "17",
5858
"CMAKE_C_STANDARD_REQUIRED": "ON"
5959
}
6060
},
@@ -89,10 +89,10 @@
8989
"name": "flags-cuda",
9090
"hidden": true,
9191
"cacheVariables": {
92-
"CUDA_NVCC_FLAGS": "--default-stream per-thread -arch=all-major",
93-
"CUDA_ARCHITECTURES": "50;52;53;60;61;62;70;72;75;80;86;87",
92+
"CMAKE_CUDA_FLAGS": "--default-stream per-thread -arch=all-major",
93+
"CMAKE_CUDA_ARCHITECTURES": "50;52;53;60;61;62;70;72;75;80;86;87",
9494
"CUDA_PROPAGATE_HOST_FLAGS": "ON",
95-
"CUDA_SEPARABLE_COMPILATION": "ON"
95+
"CMAKE_CUDA_SEPARABLE_COMPILATION": "ON"
9696
}
9797
},
9898
{
@@ -117,7 +117,7 @@
117117
"name": "ci-linux",
118118
"generator": "Unix Makefiles",
119119
"hidden": true,
120-
"inherits": ["flags-gcc-clang", "ci-std"],
120+
"inherits": ["flags-gcc-clang", "ci-std", "ci-cuda", "flags-cuda"],
121121
"cacheVariables": {
122122
"CMAKE_BUILD_TYPE": "Release"
123123
}
@@ -135,7 +135,7 @@
135135
"name": "ci-base",
136136
"generator": "Unix Makefiles",
137137
"hidden": true,
138-
"inherits": ["ci-std", "flags-gcc-clang", "dev-mode"]
138+
"inherits": ["ci-std", "flags-gcc-clang", "dev-mode", "flags-cuda", "ci-cuda"]
139139
},
140140
{
141141
"name": "ci-win64",

‎qml/responsive.qml‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -42,7 +42,7 @@ Page {
4242
id: grid
4343
anchors.left: parent.left
4444
anchors.right: parent.right
45-
columns: (window.width > 1000) ? 3 : (window.width > 680) ? 2 : 1;
45+
columns: (window.width > 900) ? 3 : (window.width > 660) ? 2 : 1;
4646

4747
ColumnLayout {
4848
id: leftBox

‎source/application.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -203,6 +203,6 @@ QString application::appVersion() {
203203
#endif
204204
}
205205

206-
QString application::appAuthor() { return QString("Benoît"); }
206+
QString application::appAuthor() { return QString("Bensuperpc"); }
207207

208208
QString application::appBusiness() { return QString("bensuperpc (bensuperpc.org)"); }

‎source/gta_cheat_finder/cuda/CMakeLists.txt‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -51,7 +51,7 @@ add_library(cuda_lib ${SRCS} ${HEADERS})
5151
target_include_directories(cuda_lib PRIVATE ${CMAKE_CURRENT_SOURCE_DIR})
5252
target_include_directories(cuda_lib PRIVATE ${CUDAToolkit_INCLUDE_DIRS})
5353
target_link_libraries(cuda_lib PRIVATE CUDA::cudart)
54-
target_compile_features(cuda_lib PUBLIC cuda_std_17)
54+
#target_compile_features(cuda_lib PUBLIC cuda_std_17)
5555
#endif()
5656

5757
set_target_properties(cuda_lib

‎source/gta_cheat_finder/cuda/kernel.cu‎

Lines changed: 59 additions & 66 deletions
Original file line numberDiff line numberDiff line change
@@ -26,87 +26,80 @@
2626
#include "kernel.cuh"
2727

2828
__global__ void jamcrc_kernel_wrapper(const void *data, uint32_t *result, const uint64_t length, const uint32_t previousCrc32) {
29-
const uint64_t blockId = blockIdx.z * gridDim.x * gridDim.y + blockIdx.y * gridDim.x + blockIdx.x;
30-
const uint64_t threadsPerBlock = blockDim.x;
31-
uint64_t id = blockId * threadsPerBlock + threadIdx.x;
29+
const uint64_t blockId = blockIdx.z * gridDim.x * gridDim.y + blockIdx.y * gridDim.x + blockIdx.x;
30+
const uint64_t threadsPerBlock = blockDim.x;
31+
uint64_t id = blockId * threadsPerBlock + threadIdx.x;
3232

33-
if (id == 0) {
34-
*result = jamcrc_kernel(data, length, previousCrc32);
35-
}
33+
if (id == 0) {
34+
*result = jamcrc_kernel(data, length, previousCrc32);
35+
}
3636
}
3737

3838
__device__ uint32_t jamcrc_kernel(const void *data, uint64_t length, const uint32_t previousCrc32) {
39-
uint32_t crc = ~previousCrc32;
40-
uint8_t *current = (uint8_t *)data;
41-
while (length--)
42-
crc = (crc >> 8) ^ crc32_lookup[(crc & 0xFF) ^ *current++];
43-
return crc;
39+
uint32_t crc = ~previousCrc32;
40+
uint8_t *current = (uint8_t *)data;
41+
while (length--)
42+
crc = (crc >> 8) ^ crc32_lookup[(crc & 0xFF) ^ *current++];
43+
return crc;
4444
}
4545

46-
__global__ void runner_kernel(uint32_t *crc_result, uint64_t *index_result, uint64_t array_size, uint64_t a, uint64_t b) {
47-
const uint64_t blockId = blockIdx.z * gridDim.x * gridDim.y + blockIdx.y * gridDim.x + blockIdx.x;
48-
const uint64_t threadsPerBlock = blockDim.x;
49-
uint64_t id = blockId * threadsPerBlock + threadIdx.x;
50-
51-
id = id + a;
46+
__global__ void runner_kernel(uint32_t *crc_result, uint64_t *index_result, uint64_t array_size, uint32_t* array_index, uint64_t a, uint64_t b) {
47+
const uint64_t blockId = blockIdx.z * gridDim.x * gridDim.y + blockIdx.y * gridDim.x + blockIdx.x;
48+
const uint64_t threadsPerBlock = blockDim.x;
49+
uint64_t id = blockId * threadsPerBlock + threadIdx.x;
5250

53-
if (id >= a && id <= b) {
54-
// printf("id: %llu a: %llu b: %llu\n", id, a, b);
55-
// Allocate memory for the array
56-
uint8_t array[29] = {0};
51+
id = id + a;
5752

58-
uint64_t size = 0;
59-
// Generate the array from index (id)
60-
find_string_inv_kernel(array, id, &size);
53+
if (id >= a && id <= b) {
54+
// printf("id: %llu a: %llu b: %llu\n", id, a, b);
55+
// Allocate memory for the array
56+
uint8_t array[29] = {0};
6157

62-
// Calculate the JAMCRC
63-
const uint32_t result = jamcrc_kernel(array, size, 0);
64-
// printf("id: %llu, size: %llu, array: %s, crc: 0x%x\n", id, size, array, result);
58+
uint64_t size = 0;
59+
// Generate the array from index (id)
60+
find_string_inv_kernel(array, id, &size);
6561

66-
bool found = false;
67-
for (uint8_t i = 0; i < 87; i++) {
68-
if (result == cheat_list[i]) {
69-
found = true;
70-
break;
71-
}
72-
}
62+
// Calculate the JAMCRC
63+
const uint32_t result = jamcrc_kernel(array, size, 0);
64+
// printf("id: %llu, size: %llu, array: %s, crc: 0x%x\n", id, size, array, result);
7365

74-
if (!found) {
75-
return;
76-
}
66+
bool found = false;
67+
for (uint8_t i = 0; i < 87; i++) {
68+
if (result == cheat_list[i]) {
69+
found = true;
70+
break;
71+
}
72+
}
7773

78-
// Todo: Avoid datarace
79-
//__syncthreads();
74+
if (!found) {
75+
return;
76+
}
8077

81-
for (uint64_t i = 0; i < array_size; i++) {
82-
if (crc_result[i] == 0 && index_result[i] == 0) {
83-
crc_result[i] = result;
84-
index_result[i] = id;
85-
// printf("Found %d at %d\n", result, id);
86-
break;
87-
}
78+
//__syncthreads();
79+
uint32_t local_array_index = atomicAdd(array_index, 1);
80+
if (local_array_index >= array_size) {
81+
return;
82+
}
83+
crc_result[local_array_index] = result;
84+
index_result[local_array_index] = id;
8885
}
89-
}
9086
}
9187

9288
__device__ void find_string_inv_kernel(uint8_t *array, uint64_t n, uint64_t *terminator_index) {
93-
const uint32_t string_size_alphabet = 27;
94-
95-
const uint8_t alpha[string_size_alphabet] = {"ABCDEFGHIJKLMNOPQRSTUVWXYZ"};
96-
// If n < 27
97-
if (n < 26) {
98-
array[0] = alpha[n];
99-
array[1] = '\0';
100-
*terminator_index = 1;
101-
return;
102-
}
103-
// If n > 27
104-
uint64_t i = 0;
105-
while (n > 0) {
106-
array[i] = alpha[(--n) % 26];
107-
n /= 26;
108-
++i;
109-
}
110-
array[i] = '\0';
111-
*terminator_index = i;
89+
// If n < 27
90+
if (n < 26) {
91+
array[0] = alpha[n];
92+
array[1] = '\0';
93+
*terminator_index = 1;
94+
return;
95+
}
96+
// If n > 27
97+
uint64_t i = 0;
98+
while (n > 0) {
99+
array[i] = alpha[(--n) % 26];
100+
n /= 26;
101+
++i;
102+
}
103+
array[i] = '\0';
104+
*terminator_index = i;
112105
}

‎source/gta_cheat_finder/cuda/kernel.cuh‎

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -34,7 +34,7 @@ __device__ uint32_t jamcrc_kernel(const void *data, uint64_t length, const uint3
3434
__global__ void jamcrc_kernel_wrapper(const void *data, uint32_t *result, uint64_t length, const uint32_t previousCrc32);
3535

3636
__device__ void find_string_inv_kernel(uint8_t *array, uint64_t n, uint64_t *terminator_index);
37-
__global__ void runner_kernel(uint32_t *crc_result, uint64_t *index_result, uint64_t array_size, uint64_t a, uint64_t b);
37+
__global__ void runner_kernel(uint32_t *crc_result, uint64_t *index_result, uint64_t array_size, uint32_t* array_index, uint64_t a, uint64_t b);
3838

3939
__device__ const uint32_t crc32_lookup[256] = {
4040
0x00000000, 0x77073096, 0xEE0E612C, 0x990951BA, 0x076DC419, 0x706AF48F, 0xE963A535, 0x9E6495A3, 0x0EDB8832, 0x79DCB8A4, 0xE0D5E91E, 0x97D2D988, 0x09B64C2B,
@@ -67,4 +67,8 @@ __device__ const uint32_t cheat_list[87] = {
6767
0x9A629401, 0xF53EF5A5, 0xF2AA0C1D, 0xF36345A8, 0x8990D5E1, 0xB7013B1B, 0xCAEC94EE, 0x31F0C3CC, 0xB3B3E72A, 0xC25CDBFF, 0xD5CF4EFF, 0x680416B1, 0xCF5FDA18,
6868
0xF01286E9, 0xA841CC0A, 0x31EA09CF, 0xE958788A, 0x02C83A7C, 0xE49C3ED4, 0x171BA8CC, 0x86988DAE, 0x2BDD2FA1};
6969

70+
__device__ const uint32_t string_size_alphabet = 27;
71+
72+
__device__ const uint8_t alpha[string_size_alphabet] = {"ABCDEFGHIJKLMNOPQRSTUVWXYZ"};
73+
7074
#endif

0 commit comments

Comments
 (0)