WIP OpenCL

Signed-off-by: Bensuperpc <bensuperpc@gmail.com>

source/gta_cheat_finder/CMakeLists.txt

@@ -84,15 +84,15 @@ if (OpenMP_FOUND OR OpenMP_CXX_FOUND)
 endif()
 
 if (CUDAToolkit_FOUND)
-  target_include_directories(GTA_SA_cheat_finder_lib PRIVATE ${CUDAToolkit_INCLUDE_DIRS})
-  target_link_libraries(GTA_SA_cheat_finder_lib PRIVATE CUDA::cudart)
-  target_link_libraries(GTA_SA_cheat_finder_lib PRIVATE cuda_lib)
+  target_include_directories(GTA_SA_cheat_finder_lib PUBLIC ${CUDAToolkit_INCLUDE_DIRS})
+  target_link_libraries(GTA_SA_cheat_finder_lib PUBLIC CUDA::cudart)
+  target_link_libraries(GTA_SA_cheat_finder_lib PUBLIC cuda_lib)
 endif()
 
 if (OpenCL_FOUND)
-  target_include_directories(GTA_SA_cheat_finder_lib PRIVATE ${OpenCL_INCLUDE_DIRS})
-  target_link_libraries(GTA_SA_cheat_finder_lib PRIVATE ${OpenCL_LIBRARIES})
-  target_link_libraries(GTA_SA_cheat_finder_lib PRIVATE opencl_lib)
+  target_include_directories(GTA_SA_cheat_finder_lib PUBLIC ${OpenCL_INCLUDE_DIRS})
+  target_link_libraries(GTA_SA_cheat_finder_lib PUBLIC ${OpenCL_LIBRARIES})
+  target_link_libraries(GTA_SA_cheat_finder_lib PUBLIC opencl_lib)
 endif()
 
 set_target_properties(GTA_SA_cheat_finder_lib

source/gta_cheat_finder/opencl/kernel.cpp

@@ -110,10 +110,11 @@ std::string my::opencl::kernel::jamcrc1Byte() {
     return R"(
         __kernel void jamcrc1Byte(__global void* data, __global ulong *length, __global uint *previousCrc32, __global uint *result)
         {
-            //size_t i = get_global_id(0);
-            uint crc = ~previousCrc32[0];
+            __local uint crc;
+            crc = ~previousCrc32[0];
             uchar* current = (uchar*)data;
-            ulong len = length[0];
+            __local ulong len;
+            len = length[0];
             while (len--) {
                 crc = (crc >> 8) ^ crc32LookupTable[0][(crc & 0xFF) ^ *current++];
             }
@@ -126,9 +127,11 @@ std::string my::opencl::kernel::jamcrc4Byte() {
     return R"(
         __kernel void jamcrc4Byte(__global void* data, __global ulong *length, __global uint *previousCrc32, __global uint *result) {
             //size_t i = get_global_id(0);
-            uint crc = ~previousCrc32[0];
+            __local uint crc;
+            crc = ~previousCrc32[0];
             uint* current = (uint*)data;
-            ulong len = length[0];
+            __local ulong len;
+            len = length[0];
             while (len >= 4) {
                 uint one = *current++ ^ crc;
                 crc = crc32LookupTable[0][(one >> 24) & 0xFF]
@@ -150,7 +153,7 @@ std::string my::opencl::kernel::jamcrc4Byte() {
 std::string my::opencl::kernel::generateString() {
     return R"(
         __constant uchar alpha[27] = "ABCDEFGHIJKLMNOPQRSTUVWXYZ";
-        __kernel void GenerateStringKernel(__global uchar* array, __global ulong* n, __global ulong* terminatorIndex) {
+        __kernel void GenerateStringKernel(__local uchar* array, __local ulong* n, __local ulong* terminatorIndex) {
             ulong index = n[0];
             // If index < 27
             if (index < 26) {
@@ -169,3 +172,24 @@ std::string my::opencl::kernel::generateString() {
         }
     )";
 }
+
+std::string my::opencl::kernel::FindAlternativeCheat() {
+    return R"(
+        __kernel void FindAlternativeCheat(__global uint* crc_result, __global ulong* index_result, __global ulong* array_size, __global uint* arrayIndex, __global ulong* a, __global ulong* b) {
+            __local ulong index;
+            index = get_global_id(0) + a[0];
+            
+            if (index >= b[0]) {
+                return;
+            }
+
+            __local uchar array[29];
+
+            __local ulong terminatorIndex;
+            terminatorIndex = 0;
+
+            GenerateStringKernel(array, &index, &terminatorIndex);
+        }
+    )";
+}
+

source/gta_cheat_finder/opencl/kernel.hpp

@@ -8,6 +8,7 @@ namespace my::opencl::kernel
     std::string jamcrc1Byte();
     std::string jamcrc4Byte();
     std::string generateString();
+    std::string FindAlternativeCheat();
 }
 
 #endif  // JAMCRC_OPENCL_HPP

source/gta_cheat_finder/opencl/wrapper.cpp

@@ -165,5 +165,222 @@ void my::opencl::launchKernel(std::vector<uint32_t>& jamcrc_results,
                                 const uint64_t maxRange,
                                 const uint64_t cudaBlockSize) {
 
+    std::cout << "Launching kernel..." << std::endl;
+
+    uint64_t calcRange = maxRange - minRange;
+
+    uint64_t arrayIndexValue = 0;
+
+    // Calculate length of the array with maxRange and minRange (Estimate size of the array)
+    uint64_t arrayLength = static_cast<uint64_t>((calcRange / 20'000'000) + 128);
+
+	cl_platform_id platformId = NULL;
+	cl_device_id deviceID = NULL;
+	cl_uint retNumDevices;
+	cl_uint retNumPlatforms;
+	cl_int ret = clGetPlatformIDs(1, &platformId, &retNumPlatforms);
+    if (ret != CL_SUCCESS) {
+        std::cout << "Error: Failed to get platform id!" << std::endl;
+        return;
+    }
+
+	ret = clGetDeviceIDs(platformId, CL_DEVICE_TYPE_DEFAULT, 1, &deviceID, &retNumDevices);
+    if (ret != CL_SUCCESS) {
+        std::cout << "Error: Failed to get device id!" << std::endl;
+        return;
+    }
+
+    cl_context context = clCreateContext(NULL, 1, &deviceID, NULL, NULL,  &ret);
+
+    cl_command_queue commandQueue = clCreateCommandQueue(context, deviceID, 0, &ret);
+
+    cl_mem crcResult = clCreateBuffer(context, CL_MEM_WRITE_ONLY, arrayLength * sizeof(uint32_t), NULL, &ret);
+    if (ret != CL_SUCCESS) {
+        std::cout << "Error: Failed to allocate device memory: crcResult" << std::endl;
+        return;
+    }
+
+    cl_mem indexResult = clCreateBuffer(context, CL_MEM_WRITE_ONLY, arrayLength  * sizeof(uint64_t), NULL, &ret);
+    if (ret != CL_SUCCESS) {
+        std::cout << "Error: Failed to allocate device memory: indexResult" << std::endl;
+        return;
+    }
+
+    cl_mem arraySize = clCreateBuffer(context, CL_MEM_READ_ONLY, 1 * sizeof(uint64_t), NULL, &ret);
+    if (ret != CL_SUCCESS) {
+        std::cout << "Error: Failed to allocate device memory: openclLength" << std::endl;
+        return;
+    }
+
+    cl_mem arrayIndex = clCreateBuffer(context, CL_MEM_WRITE_ONLY, 1 * sizeof(uint64_t), NULL, &ret);
+    if (ret != CL_SUCCESS) {
+        std::cout << "Error: Failed to allocate device memory: openclLength" << std::endl;
+        return;
+    }
+
+    cl_mem a = clCreateBuffer(context, CL_MEM_READ_ONLY, 1 * sizeof(uint64_t), NULL, &ret);
+    if (ret != CL_SUCCESS) {
+        std::cout << "Error: Failed to allocate device memory: openclLength" << std::endl;
+        return;
+    }
+
+    cl_mem b = clCreateBuffer(context, CL_MEM_READ_ONLY, 1 * sizeof(uint64_t), NULL, &ret);
+    if (ret != CL_SUCCESS) {
+        std::cout << "Error: Failed to allocate device memory: openclLength" << std::endl;
+        return;
+    }
+
+    ret = clEnqueueWriteBuffer(commandQueue, arraySize, CL_TRUE, 0, 1 * sizeof(uint64_t), &arrayLength, 0, NULL, NULL);
+    if (ret != CL_SUCCESS) {
+        std::cout << "Error: Failed to write to arraySize!" << std::endl;
+        return;
+    }
+
+    ret = clEnqueueWriteBuffer(commandQueue, arrayIndex, CL_TRUE, 0, 1 * sizeof(uint64_t), &arrayIndexValue, 0, NULL, NULL);
+    if (ret != CL_SUCCESS) {
+        std::cout << "Error: Failed to write to arrayIndex!" << std::endl;
+        return;
+    }
+
+    ret = clEnqueueWriteBuffer(commandQueue, a, CL_TRUE, 0, 1 * sizeof(uint64_t), &minRange, 0, NULL, NULL);
+    if (ret != CL_SUCCESS) {
+        std::cout << "Error: Failed to write to a!" << std::endl;
+        return;
+    }
+
+    ret = clEnqueueWriteBuffer(commandQueue, b, CL_TRUE, 0, 1 * sizeof(uint64_t), &maxRange, 0, NULL, NULL);
+    if (ret != CL_SUCCESS) {
+        std::cout << "Error: Failed to write to b!" << std::endl;
+        return;
+    }
+
+    static std::string kernelSourceStr = my::opencl::kernel::jamcrc_table() + my::opencl::kernel::jamcrc1Byte() + my::opencl::kernel::generateString() + my::opencl::kernel::FindAlternativeCheat();
+
+    const char *kernelSource = kernelSourceStr.c_str();
+    size_t kernelSize = kernelSourceStr.size();
+
+    cl_program program = clCreateProgramWithSource(context, 1, (const char **)&kernelSource, (const size_t *)&kernelSize, &ret);
+
+
+    char buildOptions[] = "-cl-std=CL3.0";
+
+    ret = clBuildProgram(program, 1, &deviceID, (const char*)&buildOptions, NULL, NULL);
+
+    if (ret != CL_SUCCESS) {
+        std::cout << "Error: Failed to build program!" << std::endl;
+
+        size_t len;
+        char buffer[8192];
+        clGetProgramBuildInfo(program, deviceID, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);
+        std::cout << std::string(buffer, len) << std::endl;
+        return;
+    }
+
+    cl_kernel kernel = clCreateKernel(program, "FindAlternativeCheat", &ret);
+
+    ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), static_cast<void*>(&crcResult));
+    if (ret != CL_SUCCESS) {
+        std::cout << "Error: Failed to set kernel arguments!" << std::endl;
+        return;
+    }
+
+    ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), static_cast<void*>(&indexResult));
+    if (ret != CL_SUCCESS) {
+        std::cout << "Error: Failed to set kernel arguments!" << std::endl;
+        return;
+    }
+
+    ret = clSetKernelArg(kernel, 2, sizeof(cl_mem), static_cast<void*>(&arraySize));
+    if (ret != CL_SUCCESS) {
+        std::cout << "Error: Failed to set kernel arguments!" << std::endl;
+        return;
+    }
+
+    ret = clSetKernelArg(kernel, 3, sizeof(cl_mem), static_cast<void*>(&arrayIndex));
+    if (ret != CL_SUCCESS) {
+        std::cout << "Error: Failed to set kernel arguments!" << std::endl;
+        return;
+    }
+
+    ret = clSetKernelArg(kernel, 4, sizeof(cl_mem), static_cast<void*>(&a));
+    if (ret != CL_SUCCESS) {
+        std::cout << "Error: Failed to set kernel arguments!" << std::endl;
+        return;
+    }
+
+    ret = clSetKernelArg(kernel, 5, sizeof(cl_mem), static_cast<void*>(&b));
+    if (ret != CL_SUCCESS) {
+        std::cout << "Error: Failed to set kernel arguments!" << std::endl;
+        return;
+    }
+
+    size_t globalItemSize = 1;
+    size_t localItemSize = 1;
+
+    ret = clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, &globalItemSize, &localItemSize, 0, NULL, NULL);
+    if (ret != CL_SUCCESS) {
+        std::cout << "Error: Failed to execute kernel!" << std::endl;
+        return;
+    }
+
+    ret = clFinish(commandQueue);
+    if (ret != CL_SUCCESS) {
+        std::cout << "Error: Failed to finish!" << std::endl;
+        if (ret == CL_OUT_OF_RESOURCES) {
+            std::cout << "CL_OUT_OF_RESOURCES" << std::endl;
+        } else if (ret == CL_OUT_OF_HOST_MEMORY) {
+            std::cout << "CL_OUT_OF_HOST_MEMORY" << std::endl;
+        } else {
+            std::cout << "Unknown error: " << ret << std::endl;
+        }
+        return;
+    }
+
+    uint32_t* jamcrc_results_ptr = new uint32_t[arrayLength];
+    uint64_t* index_results_ptr = new uint64_t[arrayLength];
+
+    ret = clEnqueueReadBuffer(commandQueue, crcResult, CL_TRUE, 0, arrayLength, jamcrc_results_ptr, 0, NULL, NULL);
+    if (ret != CL_SUCCESS) {
+        std::cout << "Error: Failed to read crcResult!" << std::endl;
+        return;
+    }
+
+
+    ret = clEnqueueReadBuffer(commandQueue, indexResult, CL_TRUE, 0, arrayLength, index_results_ptr, 0, NULL, NULL);
+    if (ret != CL_SUCCESS) {
+        std::cout << "Error: Failed to read indexResult!" << std::endl;
+        return;
+    }
+
+
+    ret = clEnqueueReadBuffer(commandQueue, arrayIndex, CL_TRUE, 0, 1 * sizeof(uint64_t), &arrayIndexValue, 0, NULL, NULL);
+    if (ret != CL_SUCCESS) {
+        std::cout << "Error: Failed to read arrayIndex!" << std::endl;
+        return;
+    }
+
+    if (arrayIndexValue > arrayLength) {
+        arrayIndexValue = arrayLength;
+        std::cout << "Warning: arrayIndexValue > arrayLength" << std::endl;
+    }
+
+    for (uint64_t i = 0; i < arrayIndexValue; i++) {
+        jamcrc_results.push_back(jamcrc_results_ptr[i]);
+        index_results.push_back(index_results_ptr[i]);
+    }
+
+    delete[] jamcrc_results_ptr;
+    delete[] index_results_ptr;
 
-    }
+    ret = clFlush(commandQueue);
+    ret = clReleaseCommandQueue(commandQueue);
+    ret = clReleaseKernel(kernel);
+    ret = clReleaseProgram(program);
+    ret = clReleaseMemObject(crcResult);
+    ret = clReleaseMemObject(indexResult);
+    ret = clReleaseMemObject(arraySize);
+    ret = clReleaseMemObject(arrayIndex);
+    ret = clReleaseMemObject(a);
+    ret = clReleaseMemObject(b);
+    ret = clReleaseContext(context);
+}

source/gta_cheat_finder/state/GTA_SA_cheat_finder_virtual.hpp

@@ -32,6 +32,12 @@
 #ifndef BUILD_WITH_CUDA
 #define BUILD_WITH_CUDA
 #endif
+#else 
+#if _MSC_VER && !__INTEL_COMPILER
+#pragma message("CUDA disabled.")
+#else
+#warning CUDA disabled.
+#endif
 #endif
 
 #if defined(BUILD_WITH_CUDA)
@@ -42,6 +48,12 @@
 #ifndef BUILD_WITH_OPENCL
 #define BUILD_WITH_OPENCL
 #endif
+#else
+#if _MSC_VER && !__INTEL_COMPILER
+#pragma message("OpenCL disabled.")
+#else
+#warning OpenCL disabled.
+#endif
 #endif
 
 #if defined(BUILD_WITH_OPENCL)