Fix

Signed-off-by: Bensuperpc <[email protected]>
bensuperpc · Dec 17, 2023 · 76009d9 · 76009d9
1 parent 37aa66c
commit 76009d9
Show file tree

Hide file tree

Showing 3 changed files with 19 additions and 5 deletions.
diff --git a/CMakePresets.json b/CMakePresets.json
@@ -89,9 +89,9 @@
       "name": "flags-cuda",
       "hidden": true,
       "cacheVariables": {
-        "CMAKE_CUDA_FLAGS": "--default-stream per-thread -arch=all-major",
-        "CMAKE_CUDA_ARCHITECTURES": "50;52;53;60;61;62;70;72;75;80;86;87",
-        "CUDA_PROPAGATE_HOST_FLAGS": "ON",
+        "CMAKE_CUDA_FLAGS": "--default-stream per-thread -Xfatbin=-compress-all -arch=all-major -Xcompiler=-Wall,-Wextra,-Wconversion,-Wsign-conversion,-Wcast-qual,-Wundef,-Wshadow,-Wunused,-Wnull-dereference,-Wdouble-promotion,-Wimplicit-fallthrough,-Wextra-semi,-Woverloaded-virtual,-Wnon-virtual-dtor,-Wformat-security",
+        "CMAKE_CUDA_ARCHITECTURES": "50;52;53;60;61;62;70;72;75;80;86;87;89;90",
+        "CUDA_PROPAGATE_HOST_FLAGS": "OFF",
         "CMAKE_CUDA_SEPARABLE_COMPILATION": "ON"
       }
     },

diff --git a/source/gta_cheat_finder/cuda/kernel.cu b/source/gta_cheat_finder/cuda/kernel.cu
@@ -63,25 +63,38 @@ __global__ void FindAlternativeCheatKernel(uint32_t *crc_result, uint64_t *index
         // 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++) {
+        for (uint8_t i = 0; i < 87; i += 3) {
             if (result == cheatList[i]) {
                 found = true;
                 break;
             }
+            if (result == cheatList[i + 1]) {
+                found = true;
+                break;
+            }
+            if (result == cheatList[i + 2]) {
+                found = true;
+                break;
+            }
         }
 
         if (!found) {
             return;
         }
 
-        //__syncthreads();
+#if __CUDA_ARCH__ >= 600
+        uint32_t localArrayIndex = atomicAdd_system(arrayIndex, 1);
+#else
         uint32_t localArrayIndex = atomicAdd(arrayIndex, 1);
+#endif
+
         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) {

diff --git a/source/gta_cheat_finder/cuda/kernel.cuh b/source/gta_cheat_finder/cuda/kernel.cuh
@@ -40,6 +40,7 @@ __device__ uint32_t jamcrcKernel(const void *data, uint64_t length, const uint32
 __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);
+
 __global__ void FindAlternativeCheatKernel(uint32_t *crc_result, uint64_t *index_result, uint64_t array_size, uint32_t* arrayIndex, uint64_t a, uint64_t b);
 
 __device__ const uint32_t crc32_lookup[256] = {