Update cuda kernel

Signed-off-by: Bensuperpc <[email protected]>
bensuperpc · Dec 22, 2024 · 35c5b7c · 35c5b7c
1 parent ad2333e
commit 35c5b7c
Show file tree

Hide file tree

Showing 6 changed files with 44 additions and 32 deletions.
diff --git a/source/gta_cheat_finder/cuda/CMakeLists.txt b/source/gta_cheat_finder/cuda/CMakeLists.txt
@@ -39,12 +39,14 @@ set(SRCS
     kernel.cu
     wrapper.cu
     jamcrc.cu
+    generateString.cu
 	)
 
 set(HEADERS
     kernel.cuh
     wrapper.hpp
     jamcrc.cuh
+    generateString.cuh
     )
 
 add_library(cuda_lib ${SRCS} ${HEADERS})

diff --git a/source/gta_cheat_finder/cuda/generateString.cu b/source/gta_cheat_finder/cuda/generateString.cu
@@ -0,0 +1,23 @@
+#include "generateString.cuh"
+
+__device__ void generateStringKernel(uint8_t* array, uint64_t n, uint64_t* terminatorIndex) {
+    // If n > 27
+    uint64_t i = 0;
+    while (n > 0) {
+        array[i] = alpha[(--n) % 26];
+        n /= 26;
+        ++i;
+    }
+    *terminatorIndex = i;
+}
+
+__device__ void generateStringKernelV2(uint8_t* array, uint64_t n, uint64_t* terminatorIndex) {
+    uint64_t i = 0;
+    do {
+        array[i++] = alpha[--n % 26];
+        n /= 26;
+    } while (n > 0);
+
+    *terminatorIndex = i;
+}
+
diff --git a/source/gta_cheat_finder/cuda/generateString.cuh b/source/gta_cheat_finder/cuda/generateString.cuh
@@ -0,0 +1,13 @@
+#ifndef CUDA_GENERATESTRING_CUH
+#define CUDA_GENERATESTRING_CUH
+
+#include <cuda.h>
+#include <cuda_runtime.h>
+#include <stdio.h>
+
+__device__ void generateStringKernel(uint8_t* array, uint64_t n, uint64_t* terminatorIndex);
+__device__ void generateStringKernelV2(uint8_t* array, uint64_t n, uint64_t* terminatorIndex);
+
+__device__ const uint8_t alpha[27] = {"ABCDEFGHIJKLMNOPQRSTUVWXYZ"};
+
+#endif
diff --git a/source/gta_cheat_finder/cuda/jamcrc.cuh b/source/gta_cheat_finder/cuda/jamcrc.cuh
@@ -1,5 +1,5 @@
-#ifndef CUDA_JAMCRC_H
-#define CUDA_JAMCRC_H
+#ifndef CUDA_JAMCRC_CUH
+#define CUDA_JAMCRC_CUH
 
 #include <cuda.h>
 #include <cuda_runtime.h>

diff --git a/source/gta_cheat_finder/cuda/kernel.cu b/source/gta_cheat_finder/cuda/kernel.cu
@@ -48,16 +48,12 @@ __global__ void findAlternativeCheatKernel(uint64_t* index_result,
     a = id + a;
 
     if (a <= b) {
-        //  Allocate memory for the array
         uint8_t array[29] = {0};
-
         uint64_t size = 0;
-        // Generate the array from index (a)
+
         generateStringKernel(array, a, &size);
 
-        // Calculate the JAMCRC
         const uint32_t result = jamcrc1Byte(array, size, 0);
-        // printf("id: %llu, size: %llu, array: %s, crc: 0x%x\n", id, size, array, result);
 
         bool found = false;
         for (uint8_t i = 0; i < 87; i += 3) {
@@ -70,6 +66,7 @@ __global__ void findAlternativeCheatKernel(uint64_t* index_result,
         if (!found) {
             return;
         }
+
 #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ < 600)
         uint32_t localArrayIndex = atomicAdd(arrayIndex, 1);
 #else
@@ -79,29 +76,9 @@ __global__ void findAlternativeCheatKernel(uint64_t* index_result,
         if (localArrayIndex >= array_size) {
             return;
         }
+
         crc_result[localArrayIndex] = result;
         index_result[localArrayIndex] = a;
     }
     //__syncthreads();
 }
-
-__device__ void generateStringKernel(uint8_t* array, uint64_t n, uint64_t* terminatorIndex) {
-    // If n > 27
-    uint64_t i = 0;
-    while (n > 0) {
-        array[i] = alpha[(--n) % 26];
-        n /= 26;
-        ++i;
-    }
-    *terminatorIndex = i;
-}
-
-__device__ void generateStringKernelV2(uint8_t* array, uint64_t n, uint64_t* terminatorIndex) {
-    uint64_t i = 0;
-    do {
-        array[i++] = alpha[--n % 26];
-        n /= 26;
-    } while (n > 0);
-
-    *terminatorIndex = i;
-}
diff --git a/source/gta_cheat_finder/cuda/kernel.cuh b/source/gta_cheat_finder/cuda/kernel.cuh
@@ -30,12 +30,10 @@
 #include <stdio.h>
 
 #include "jamcrc.cuh"
+#include "generateString.cuh"
 
 __global__ void jamcrcKernelWrapper(const void* data, uint32_t* result, uint64_t length, const uint32_t previousCrc32);
 
-__device__ void generateStringKernel(uint8_t* array, uint64_t n, uint64_t* terminatorIndex);
-__device__ void generateStringKernelV2(uint8_t* array, uint64_t n, uint64_t* terminatorIndex);
-
 __global__ void findAlternativeCheatKernel(uint64_t* index_result,
                                            uint32_t* crc_result,
                                            uint64_t array_size,
@@ -53,6 +51,5 @@ __device__ const uint32_t cheatList[87] = {
     0xF53EF5A5, 0xF2AA0C1D, 0xF36345A8, 0x8990D5E1, 0xB7013B1B, 0xCAEC94EE, 0x31F0C3CC, 0xB3B3E72A, 0xC25CDBFF, 0xD5CF4EFF, 0x680416B1,
     0xCF5FDA18, 0xF01286E9, 0xA841CC0A, 0x31EA09CF, 0xE958788A, 0x02C83A7C, 0xE49C3ED4, 0x171BA8CC, 0x86988DAE, 0x2BDD2FA1};
 
-__device__ const uint8_t alpha[27] = {"ABCDEFGHIJKLMNOPQRSTUVWXYZ"};
 
 #endif