diff --git a/source/gta_cheat_finder/CMakeLists.txt b/source/gta_cheat_finder/CMakeLists.txt index f840076..2a87e5b 100644 --- a/source/gta_cheat_finder/CMakeLists.txt +++ b/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 diff --git a/source/gta_cheat_finder/opencl/kernel.cpp b/source/gta_cheat_finder/opencl/kernel.cpp index 1700c75..58cdc3e 100644 --- a/source/gta_cheat_finder/opencl/kernel.cpp +++ b/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); + } + )"; +} + diff --git a/source/gta_cheat_finder/opencl/kernel.hpp b/source/gta_cheat_finder/opencl/kernel.hpp index 0712f67..7824c9c 100644 --- a/source/gta_cheat_finder/opencl/kernel.hpp +++ b/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 diff --git a/source/gta_cheat_finder/opencl/wrapper.cpp b/source/gta_cheat_finder/opencl/wrapper.cpp index 68926a0..7c8ad24 100644 --- a/source/gta_cheat_finder/opencl/wrapper.cpp +++ b/source/gta_cheat_finder/opencl/wrapper.cpp @@ -165,5 +165,222 @@ void my::opencl::launchKernel(std::vector& 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((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(&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(&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(&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(&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(&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(&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; - } \ No newline at end of file + 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); +} \ No newline at end of file diff --git a/source/gta_cheat_finder/state/GTA_SA_cheat_finder_virtual.hpp b/source/gta_cheat_finder/state/GTA_SA_cheat_finder_virtual.hpp index 73cee0e..27a2c12 100644 --- a/source/gta_cheat_finder/state/GTA_SA_cheat_finder_virtual.hpp +++ b/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) diff --git a/test/source/test/gta_sa_test.cpp b/test/source/test/gta_sa_test.cpp index f904602..45d6114 100644 --- a/test/source/test/gta_sa_test.cpp +++ b/test/source/test/gta_sa_test.cpp @@ -1,10 +1,28 @@ #include +#include "GTA_SA_cheat_finder_virtual.hpp" #include "GTA_SA_cheat_finder_openmp.hpp" #include "GTA_SA_cheat_finder_stdthread.hpp" + #if defined(BUILD_WITH_CUDA) #include "GTA_SA_cheat_finder_cuda.hpp" +#else +#if _MSC_VER && !__INTEL_COMPILER +#pragma message("CUDA test disabled.") +#else +#warning CUDA test disabled. +#endif #endif // BUILD_WITH_CUDA +#if defined(BUILD_WITH_OPENCL) +#include "GTA_SA_cheat_finder_opencl.hpp" +#else +#if _MSC_VER && !__INTEL_COMPILER +#pragma message("OpenCL test disabled.") +#else +#warning OpenCL test disabled. +#endif +#endif // BUILD_WITH_OPENCL + #include "gtest/gtest.h" /* @@ -106,7 +124,6 @@ TEST(GTA_SA_OPENMP, basic_calc_mode_openmp) { #endif #if defined(BUILD_WITH_CUDA) - TEST(GTA_SA_CUDA, basic_calc_cuda_1) { GTA_SA_CUDA gtaSA; gtaSA.minRange = 0; @@ -186,7 +203,89 @@ TEST(GTA_SA_CUDA, basic_calc_cuda_5) { EXPECT_EQ(gtaSA.results[151].jamcrc, 0xe1ef01ea); } */ +#endif + +#if defined(BUILD_WITH_OPENCL) + +TEST(GTA_SA_OPENCL, basic_calc_opencl_1) { + GTA_SA_OPENCL gtaSA; + gtaSA.minRange = 0; + gtaSA.maxRange = 6000000; + + gtaSA.run(); + + EXPECT_EQ(gtaSA.results.size(), 0); +} +/* +TEST(GTA_SA_OPENCL, basic_calc_opencl_2) { + GTA_SA_OPENCL gtaSA; + gtaSA.minRange = 20810700; + gtaSA.maxRange = 20810800; + + gtaSA.run(); + + EXPECT_EQ(gtaSA.results.size(), 1); + EXPECT_EQ(gtaSA.results[0].index, 20810792); + EXPECT_EQ(gtaSA.results[0].code, "ASNAEB"); + EXPECT_EQ(gtaSA.results[0].jamcrc, 0x555fc201); +} + +TEST(GTA_SA_OPENCL, basic_calc_opencl_3) { + GTA_SA_OPENCL gtaSA; + gtaSA.minRange = 181961000; + gtaSA.maxRange = 181961100; + + gtaSA.run(); + + EXPECT_EQ(gtaSA.results.size(), 1); + EXPECT_EQ(gtaSA.results[0].index, 181961057); + EXPECT_EQ(gtaSA.results[0].code, "OHDUDE"); + EXPECT_EQ(gtaSA.results[0].jamcrc, 0xe958788a); +} + +TEST(GTA_SA_OPENCL, basic_calc_opencl_4) { + GTA_SA_OPENCL gtaSA; + gtaSA.minRange = 299376700; + gtaSA.maxRange = 299376800; + + gtaSA.run(); + + EXPECT_EQ(gtaSA.results.size(), 1); + EXPECT_EQ(gtaSA.results[0].index, 299376767); + EXPECT_EQ(gtaSA.results[0].code, "YECGAA"); + EXPECT_EQ(gtaSA.results[0].jamcrc, 0x40cf761); +} +*/ +/* +TEST(GTA_SA_OPENCL, basic_calc_cuda_5) { + GTA_SA_OPENCL gtaSA; + gtaSA.minRange = 0; + gtaSA.maxRange = 8'031'810'176; + + gtaSA.run(); + + EXPECT_EQ(gtaSA.results.size(), 152); + EXPECT_EQ(gtaSA.results[0].index, 20810792); + EXPECT_EQ(gtaSA.results[0].code, "ASNAEB"); + EXPECT_EQ(gtaSA.results[0].jamcrc, 0x555fc201); + + EXPECT_EQ(gtaSA.results[50].index, 2746380464); + EXPECT_EQ(gtaSA.results[50].code, "HWCWJYV"); + EXPECT_EQ(gtaSA.results[50].jamcrc, 0xb2Afe368); + + EXPECT_EQ(gtaSA.results[100].index, 5151099438); + EXPECT_EQ(gtaSA.results[100].code, "PQNCSOD"); + EXPECT_EQ(gtaSA.results[100].jamcrc, 0x680416b1); + EXPECT_EQ(gtaSA.results[125].index, 6668091073); + EXPECT_EQ(gtaSA.results[125].code, "UOETDAG"); + EXPECT_EQ(gtaSA.results[125].jamcrc, 0xd4966d59); + + EXPECT_EQ(gtaSA.results[151].index, 8003059504); + EXPECT_EQ(gtaSA.results[151].code, "YWOBEJX"); + EXPECT_EQ(gtaSA.results[151].jamcrc, 0xe1ef01ea); +} +*/ #endif auto main(int argc, char** argv) -> int { diff --git a/test/source/test/jamcrc_test.cpp b/test/source/test/jamcrc_test.cpp index a5b4218..3689b5a 100644 --- a/test/source/test/jamcrc_test.cpp +++ b/test/source/test/jamcrc_test.cpp @@ -1,12 +1,17 @@ #include #include "GTA_SA_cheat_finder_openmp.hpp" #include "GTA_SA_cheat_finder_stdthread.hpp" + #if defined(BUILD_WITH_CUDA) #include "GTA_SA_cheat_finder_cuda.hpp" #include "cuda/wrapper.hpp" #endif // BUILD_WITH_CUDA +#if defined(BUILD_WITH_OPENCL) +#include "GTA_SA_cheat_finder_opencl.hpp" #include "opencl/wrapper.hpp" +#endif // BUILD_WITH_OPENCL + #include "gtest/gtest.h" @@ -151,6 +156,7 @@ TEST(jamcrc, cuda_basic6) { #endif +#if defined(BUILD_WITH_OPENCL) TEST(jamcrc, opencl_basic1) { std::string str = ""; uint32_t expected_crc = 0xffffffff; @@ -222,6 +228,7 @@ TEST(jamcrc, opencl_basic6) { EXPECT_NE(0x0, crc); EXPECT_EQ(expected_crc, crc); } +#endif auto main(int argc, char** argv) -> int { ::testing::InitGoogleTest(&argc, argv);