From d825187add7e32ec0ced2aecf30d3a5e46f8196c Mon Sep 17 00:00:00 2001 From: Bensuperpc Date: Sat, 16 Dec 2023 21:41:38 +0100 Subject: [PATCH] Fix Signed-off-by: Bensuperpc --- CMakePresets.json | 12 +- qml/responsive.qml | 2 +- source/application.cpp | 2 +- source/gta_cheat_finder/cuda/CMakeLists.txt | 2 +- source/gta_cheat_finder/cuda/kernel.cu | 125 ++++++------- source/gta_cheat_finder/cuda/kernel.cuh | 6 +- source/gta_cheat_finder/cuda/wrapper.cu | 169 +++++++++--------- .../state/GTA_SA_cheat_finder_virtual.cpp | 7 +- test/CMakeLists.txt | 2 +- test/source/test/gta_sa_test.cpp | 112 +++++++++--- 10 files changed, 252 insertions(+), 187 deletions(-) diff --git a/CMakePresets.json b/CMakePresets.json index 196e474..49db98f 100755 --- a/CMakePresets.json +++ b/CMakePresets.json @@ -54,7 +54,7 @@ "CMAKE_CXX_STANDARD": "20", "CMAKE_CXX_STANDARD_REQUIRED": "ON", "CMAKE_C_EXTENSIONS": "ON", - "CMAKE_C_STANDARD": "11", + "CMAKE_C_STANDARD": "17", "CMAKE_C_STANDARD_REQUIRED": "ON" } }, @@ -89,10 +89,10 @@ "name": "flags-cuda", "hidden": true, "cacheVariables": { - "CUDA_NVCC_FLAGS": "--default-stream per-thread -arch=all-major", - "CUDA_ARCHITECTURES": "50;52;53;60;61;62;70;72;75;80;86;87", + "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", - "CUDA_SEPARABLE_COMPILATION": "ON" + "CMAKE_CUDA_SEPARABLE_COMPILATION": "ON" } }, { @@ -117,7 +117,7 @@ "name": "ci-linux", "generator": "Unix Makefiles", "hidden": true, - "inherits": ["flags-gcc-clang", "ci-std"], + "inherits": ["flags-gcc-clang", "ci-std", "ci-cuda", "flags-cuda"], "cacheVariables": { "CMAKE_BUILD_TYPE": "Release" } @@ -135,7 +135,7 @@ "name": "ci-base", "generator": "Unix Makefiles", "hidden": true, - "inherits": ["ci-std", "flags-gcc-clang", "dev-mode"] + "inherits": ["ci-std", "flags-gcc-clang", "dev-mode", "flags-cuda", "ci-cuda"] }, { "name": "ci-win64", diff --git a/qml/responsive.qml b/qml/responsive.qml index e3bd00c..9f41043 100644 --- a/qml/responsive.qml +++ b/qml/responsive.qml @@ -42,7 +42,7 @@ Page { id: grid anchors.left: parent.left anchors.right: parent.right - columns: (window.width > 1000) ? 3 : (window.width > 680) ? 2 : 1; + columns: (window.width > 900) ? 3 : (window.width > 660) ? 2 : 1; ColumnLayout { id: leftBox diff --git a/source/application.cpp b/source/application.cpp index aae2747..0c3aaab 100644 --- a/source/application.cpp +++ b/source/application.cpp @@ -203,6 +203,6 @@ QString application::appVersion() { #endif } -QString application::appAuthor() { return QString("Benoît"); } +QString application::appAuthor() { return QString("Bensuperpc"); } QString application::appBusiness() { return QString("bensuperpc (bensuperpc.org)"); } diff --git a/source/gta_cheat_finder/cuda/CMakeLists.txt b/source/gta_cheat_finder/cuda/CMakeLists.txt index a50f6b0..1d98998 100644 --- a/source/gta_cheat_finder/cuda/CMakeLists.txt +++ b/source/gta_cheat_finder/cuda/CMakeLists.txt @@ -51,7 +51,7 @@ add_library(cuda_lib ${SRCS} ${HEADERS}) target_include_directories(cuda_lib PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}) target_include_directories(cuda_lib PRIVATE ${CUDAToolkit_INCLUDE_DIRS}) target_link_libraries(cuda_lib PRIVATE CUDA::cudart) -target_compile_features(cuda_lib PUBLIC cuda_std_17) +#target_compile_features(cuda_lib PUBLIC cuda_std_17) #endif() set_target_properties(cuda_lib diff --git a/source/gta_cheat_finder/cuda/kernel.cu b/source/gta_cheat_finder/cuda/kernel.cu index 5883c89..8808e42 100644 --- a/source/gta_cheat_finder/cuda/kernel.cu +++ b/source/gta_cheat_finder/cuda/kernel.cu @@ -26,87 +26,80 @@ #include "kernel.cuh" __global__ void jamcrc_kernel_wrapper(const void *data, uint32_t *result, const uint64_t length, const uint32_t previousCrc32) { - const uint64_t blockId = blockIdx.z * gridDim.x * gridDim.y + blockIdx.y * gridDim.x + blockIdx.x; - const uint64_t threadsPerBlock = blockDim.x; - uint64_t id = blockId * threadsPerBlock + threadIdx.x; + const uint64_t blockId = blockIdx.z * gridDim.x * gridDim.y + blockIdx.y * gridDim.x + blockIdx.x; + const uint64_t threadsPerBlock = blockDim.x; + uint64_t id = blockId * threadsPerBlock + threadIdx.x; - if (id == 0) { - *result = jamcrc_kernel(data, length, previousCrc32); - } + if (id == 0) { + *result = jamcrc_kernel(data, length, previousCrc32); + } } __device__ uint32_t jamcrc_kernel(const void *data, uint64_t length, const uint32_t previousCrc32) { - uint32_t crc = ~previousCrc32; - uint8_t *current = (uint8_t *)data; - while (length--) - crc = (crc >> 8) ^ crc32_lookup[(crc & 0xFF) ^ *current++]; - return crc; + uint32_t crc = ~previousCrc32; + uint8_t *current = (uint8_t *)data; + while (length--) + crc = (crc >> 8) ^ crc32_lookup[(crc & 0xFF) ^ *current++]; + return crc; } -__global__ void runner_kernel(uint32_t *crc_result, uint64_t *index_result, uint64_t array_size, uint64_t a, uint64_t b) { - const uint64_t blockId = blockIdx.z * gridDim.x * gridDim.y + blockIdx.y * gridDim.x + blockIdx.x; - const uint64_t threadsPerBlock = blockDim.x; - uint64_t id = blockId * threadsPerBlock + threadIdx.x; - - id = id + a; +__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) { + const uint64_t blockId = blockIdx.z * gridDim.x * gridDim.y + blockIdx.y * gridDim.x + blockIdx.x; + const uint64_t threadsPerBlock = blockDim.x; + uint64_t id = blockId * threadsPerBlock + threadIdx.x; - if (id >= a && id <= b) { - // printf("id: %llu a: %llu b: %llu\n", id, a, b); - // Allocate memory for the array - uint8_t array[29] = {0}; + id = id + a; - uint64_t size = 0; - // Generate the array from index (id) - find_string_inv_kernel(array, id, &size); + if (id >= a && id <= b) { + // printf("id: %llu a: %llu b: %llu\n", id, a, b); + // Allocate memory for the array + uint8_t array[29] = {0}; - // Calculate the JAMCRC - const uint32_t result = jamcrc_kernel(array, size, 0); - // printf("id: %llu, size: %llu, array: %s, crc: 0x%x\n", id, size, array, result); + uint64_t size = 0; + // Generate the array from index (id) + find_string_inv_kernel(array, id, &size); - bool found = false; - for (uint8_t i = 0; i < 87; i++) { - if (result == cheat_list[i]) { - found = true; - break; - } - } + // Calculate the JAMCRC + const uint32_t result = jamcrc_kernel(array, size, 0); + // printf("id: %llu, size: %llu, array: %s, crc: 0x%x\n", id, size, array, result); - if (!found) { - return; - } + bool found = false; + for (uint8_t i = 0; i < 87; i++) { + if (result == cheat_list[i]) { + found = true; + break; + } + } - // Todo: Avoid datarace - //__syncthreads(); + if (!found) { + return; + } - for (uint64_t i = 0; i < array_size; i++) { - if (crc_result[i] == 0 && index_result[i] == 0) { - crc_result[i] = result; - index_result[i] = id; - // printf("Found %d at %d\n", result, id); - break; - } + //__syncthreads(); + uint32_t local_array_index = atomicAdd(array_index, 1); + if (local_array_index >= array_size) { + return; + } + crc_result[local_array_index] = result; + index_result[local_array_index] = id; } - } } __device__ void find_string_inv_kernel(uint8_t *array, uint64_t n, uint64_t *terminator_index) { - const uint32_t string_size_alphabet = 27; - - const uint8_t alpha[string_size_alphabet] = {"ABCDEFGHIJKLMNOPQRSTUVWXYZ"}; - // If n < 27 - if (n < 26) { - array[0] = alpha[n]; - array[1] = '\0'; - *terminator_index = 1; - return; - } - // If n > 27 - uint64_t i = 0; - while (n > 0) { - array[i] = alpha[(--n) % 26]; - n /= 26; - ++i; - } - array[i] = '\0'; - *terminator_index = i; + // If n < 27 + if (n < 26) { + array[0] = alpha[n]; + array[1] = '\0'; + *terminator_index = 1; + return; + } + // If n > 27 + uint64_t i = 0; + while (n > 0) { + array[i] = alpha[(--n) % 26]; + n /= 26; + ++i; + } + array[i] = '\0'; + *terminator_index = i; } diff --git a/source/gta_cheat_finder/cuda/kernel.cuh b/source/gta_cheat_finder/cuda/kernel.cuh index 5a4c57b..2ff0757 100644 --- a/source/gta_cheat_finder/cuda/kernel.cuh +++ b/source/gta_cheat_finder/cuda/kernel.cuh @@ -34,7 +34,7 @@ __device__ uint32_t jamcrc_kernel(const void *data, uint64_t length, const uint3 __global__ void jamcrc_kernel_wrapper(const void *data, uint32_t *result, uint64_t length, const uint32_t previousCrc32); __device__ void find_string_inv_kernel(uint8_t *array, uint64_t n, uint64_t *terminator_index); -__global__ void runner_kernel(uint32_t *crc_result, uint64_t *index_result, uint64_t array_size, uint64_t a, uint64_t b); +__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); __device__ const uint32_t crc32_lookup[256] = { 0x00000000, 0x77073096, 0xEE0E612C, 0x990951BA, 0x076DC419, 0x706AF48F, 0xE963A535, 0x9E6495A3, 0x0EDB8832, 0x79DCB8A4, 0xE0D5E91E, 0x97D2D988, 0x09B64C2B, @@ -67,4 +67,8 @@ __device__ const uint32_t cheat_list[87] = { 0x9A629401, 0xF53EF5A5, 0xF2AA0C1D, 0xF36345A8, 0x8990D5E1, 0xB7013B1B, 0xCAEC94EE, 0x31F0C3CC, 0xB3B3E72A, 0xC25CDBFF, 0xD5CF4EFF, 0x680416B1, 0xCF5FDA18, 0xF01286E9, 0xA841CC0A, 0x31EA09CF, 0xE958788A, 0x02C83A7C, 0xE49C3ED4, 0x171BA8CC, 0x86988DAE, 0x2BDD2FA1}; +__device__ const uint32_t string_size_alphabet = 27; + +__device__ const uint8_t alpha[string_size_alphabet] = {"ABCDEFGHIJKLMNOPQRSTUVWXYZ"}; + #endif diff --git a/source/gta_cheat_finder/cuda/wrapper.cu b/source/gta_cheat_finder/cuda/wrapper.cu index c0d83c8..3c69696 100644 --- a/source/gta_cheat_finder/cuda/wrapper.cu +++ b/source/gta_cheat_finder/cuda/wrapper.cu @@ -27,125 +27,132 @@ __host__ void jamcrc_wrapper(dim3 *grid, dim3 *threads, cudaStream_t *stream, const int device, const void *data, const uint64_t length, uint32_t *result, const uint32_t previousCrc32) { - jamcrc_kernel_wrapper<<<*grid, *threads, device, *stream>>>(data, result, length, previousCrc32); + jamcrc_kernel_wrapper<<<*grid, *threads, device, *stream>>>(data, result, length, previousCrc32); } __host__ uint32_t my::cuda::jamcrc(const void *data, const uint64_t length, const uint32_t previousCrc32, const uint32_t cuda_block_size) { - int device = 0; - cudaGetDevice(&device); + int device = 0; + cudaGetDevice(&device); - cudaStream_t stream; - cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking); + cudaStream_t stream; + cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking); - // Calculate length of the array with max_range and min_range - uint64_t data_size = (length) * sizeof(char); - uint32_t *data_cuda = nullptr; + // Calculate length of the array with max_range and min_range + uint64_t data_size = (length) * sizeof(char); + uint32_t *data_cuda = nullptr; - uint64_t result_size = 1 * sizeof(uint32_t); - uint32_t *result_cuda = nullptr; + uint64_t result_size = 1 * sizeof(uint32_t); + uint32_t *result_cuda = nullptr; - cudaMallocManaged(&data_cuda, data_size, cudaMemAttachGlobal); - cudaMallocManaged(&result_cuda, result_size, cudaMemAttachGlobal); + cudaMallocManaged(&data_cuda, data_size, cudaMemAttachGlobal); + cudaMallocManaged(&result_cuda, result_size, cudaMemAttachGlobal); - cudaStreamAttachMemAsync(stream, &data_cuda); - cudaStreamAttachMemAsync(stream, &result_cuda); + cudaStreamAttachMemAsync(stream, &data_cuda); + cudaStreamAttachMemAsync(stream, &result_cuda); - cudaMemPrefetchAsync(data_cuda, data_size, device, stream); - cudaMemPrefetchAsync(result_cuda, result_size, device, stream); + cudaMemPrefetchAsync(data_cuda, data_size, device, stream); + cudaMemPrefetchAsync(result_cuda, result_size, device, stream); - memcpy(data_cuda, data, data_size); - *result_cuda = 0; + memcpy(data_cuda, data, data_size); + *result_cuda = 0; - uint64_t grid_size = static_cast(ceil(static_cast(1) / cuda_block_size)); + uint64_t grid_size = static_cast(ceil(static_cast(1) / cuda_block_size)); - dim3 threads(static_cast(cuda_block_size), 1, 1); - dim3 grid(static_cast(grid_size), 1, 1); + dim3 threads(static_cast(cuda_block_size), 1, 1); + dim3 grid(static_cast(grid_size), 1, 1); - jamcrc_kernel_wrapper<<>>(data_cuda, result_cuda, length, previousCrc32); + jamcrc_kernel_wrapper<<>>(data_cuda, result_cuda, length, previousCrc32); - cudaStreamSynchronize(stream); - cudaDeviceSynchronize(); + cudaStreamSynchronize(stream); + cudaDeviceSynchronize(); - // std::cout << "result_cuda: " << *result_cuda << std::endl; + // std::cout << "result_cuda: " << *result_cuda << std::endl; - cudaFree(data_cuda); - cudaStreamDestroy(stream); + cudaFree(data_cuda); + cudaStreamDestroy(stream); - return *result_cuda; + return *result_cuda; } __host__ void my::cuda::launch_kernel(std::vector &jamcrc_results, std::vector &index_results, const uint64_t min_range, const uint64_t max_range, const uint64_t cuda_block_size) { - std::cout << "Launching kernel..." << std::endl; - std::cout << "min_range: " << min_range << std::endl; - std::cout << "max_range: " << max_range << std::endl; - // int device = -1; - // cudaGetDevice(&device); + std::cout << "Launching kernel..." << std::endl; + std::cout << "min_range: " << std::dec << min_range << std::endl; + std::cout << "max_range: " << std::dec << max_range << std::endl; + // int device = -1; + // cudaGetDevice(&device); - int device = 0; - cudaGetDevice(&device); + int device = 0; + cudaGetDevice(&device); - /* - int priority_high, priority_low; - cudaDeviceGetStreamPriorityRange(&priority_low, &priority_high); - cudaStream_t st_high, st_low; - cudaStreamCreateWithPriority(&st_high, cudaStreamNonBlocking, priority_high); - cudaStreamCreateWithPriority(&st_low, cudaStreamNonBlocking, priority_low); - */ + /* + int priority_high, priority_low; + cudaDeviceGetStreamPriorityRange(&priority_low, &priority_high); + cudaStream_t st_high, st_low; + cudaStreamCreateWithPriority(&st_high, cudaStreamNonBlocking, priority_high); + cudaStreamCreateWithPriority(&st_low, cudaStreamNonBlocking, priority_low); + */ - cudaStream_t stream; - cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking); + cudaStream_t stream; + cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking); - // cudaDeviceSetLimit(cudaLimitMallocHeapSize, 128 * 1024 * 1024); + // cudaDeviceSetLimit(cudaLimitMallocHeapSize, 128 * 1024 * 1024); - // Calculate length of the array with max_range and min_range (Estimate size of the array) - uint64_t array_length = static_cast((max_range - min_range) / 20000000); - uint64_t jamcrc_results_size = array_length * sizeof(uint32_t); - uint64_t index_results_size = array_length * sizeof(uint64_t); + // Calculate length of the array with max_range and min_range (Estimate size of the array) + uint64_t array_length = static_cast(((max_range - min_range) / 20'000'000) + 100); + uint64_t jamcrc_results_size = array_length * sizeof(uint32_t); + uint64_t index_results_size = array_length * sizeof(uint64_t); - uint32_t *jamcrc_results_ptr = nullptr; - uint64_t *index_results_ptr = nullptr; + uint32_t *jamcrc_results_ptr = nullptr; + uint64_t *index_results_ptr = nullptr; - cudaMallocManaged(&jamcrc_results_ptr, jamcrc_results_size, cudaMemAttachGlobal); - cudaMallocManaged(&index_results_ptr, index_results_size, cudaMemAttachGlobal); + uint32_t *array_index = nullptr; - cudaStreamAttachMemAsync(stream, &jamcrc_results_ptr); - cudaStreamAttachMemAsync(stream, &index_results_size); + cudaMallocManaged(&jamcrc_results_ptr, jamcrc_results_size, cudaMemAttachGlobal); + cudaMallocManaged(&index_results_ptr, index_results_size, cudaMemAttachGlobal); + cudaMallocManaged(&array_index, 1 * sizeof(uint32_t), cudaMemAttachGlobal); - cudaMemPrefetchAsync(jamcrc_results_ptr, jamcrc_results_size, device, stream); - cudaMemPrefetchAsync(index_results_ptr, index_results_size, device, stream); + cudaStreamAttachMemAsync(stream, &jamcrc_results_ptr); + cudaStreamAttachMemAsync(stream, &index_results_size); + cudaStreamAttachMemAsync(stream, &array_index); - for (uint64_t i = 0; i < array_length; ++i) { - jamcrc_results_ptr[i] = 0; - index_results_ptr[i] = 0; - } + cudaMemPrefetchAsync(jamcrc_results_ptr, jamcrc_results_size, device, stream); + cudaMemPrefetchAsync(index_results_ptr, index_results_size, device, stream); + cudaMemPrefetchAsync(array_index, 1 * sizeof(uint32_t), device, stream); - uint64_t grid_size = static_cast(ceil(static_cast(max_range - min_range) / cuda_block_size)); - std::cout << "CUDA Grid size: " << grid_size << std::endl; - std::cout << "CUDA Block size: " << cuda_block_size << std::endl; + for (uint64_t i = 0; i < array_length; ++i) { + jamcrc_results_ptr[i] = 0; + index_results_ptr[i] = 0; + } + *array_index = 0; + + uint64_t grid_size = static_cast(ceil(static_cast(max_range - min_range) / cuda_block_size)); + std::cout << "CUDA Grid size: " << grid_size << std::endl; + std::cout << "CUDA Block size: " << cuda_block_size << std::endl; - dim3 threads(static_cast(cuda_block_size), 1, 1); - dim3 grid(static_cast(grid_size), 1, 1); + dim3 threads(static_cast(cuda_block_size), 1, 1); + dim3 grid(static_cast(grid_size), 1, 1); - runner_kernel<<>>(jamcrc_results_ptr, index_results_ptr, array_length, min_range, max_range); + runner_kernel<<>>(jamcrc_results_ptr, index_results_ptr, array_length, array_index, min_range, max_range); - jamcrc_results.reserve(array_length); - index_results.reserve(array_length); + jamcrc_results.reserve(array_length); + index_results.reserve(array_length); - cudaStreamSynchronize(stream); + cudaStreamSynchronize(stream); - for (uint64_t i = 0; i < array_length; ++i) { - if (jamcrc_results_ptr[i] != index_results_ptr[i]) { - jamcrc_results.emplace_back(jamcrc_results_ptr[i]); - index_results.emplace_back(index_results_ptr[i]); + for (uint64_t i = 0; i < *array_index; ++i) { + jamcrc_results.emplace_back(jamcrc_results_ptr[i]); + index_results.emplace_back(index_results_ptr[i]); } - } + + cudaDeviceSynchronize(); + cudaFree(jamcrc_results_ptr); + cudaFree(index_results_ptr); + cudaFree(array_index); - cudaDeviceSynchronize(); - cudaFree(jamcrc_results_ptr); - cudaFree(index_results_ptr); + cudaStreamDestroy(stream); + // cudaStreamDestroy(st_high); + // cudaStreamDestroy(st_low); - cudaStreamDestroy(stream); - // cudaStreamDestroy(st_high); - // cudaStreamDestroy(st_low); + std::cout << "CUDA Kernel finished" << std::endl; } diff --git a/source/gta_cheat_finder/state/GTA_SA_cheat_finder_virtual.cpp b/source/gta_cheat_finder/state/GTA_SA_cheat_finder_virtual.cpp index 27dfcac..6b179f7 100644 --- a/source/gta_cheat_finder/state/GTA_SA_cheat_finder_virtual.cpp +++ b/source/gta_cheat_finder/state/GTA_SA_cheat_finder_virtual.cpp @@ -52,15 +52,13 @@ uint32_t GTA_SA_Virtual::max_thread_support() { void GTA_SA_Virtual::printResult() const { std::cout << "" << std::endl; - // std::sort(results.begin(), results.end(), [](const result &a, const result &b) { return a.index < b.index; }); - constexpr auto display_val = 18; - std::cout << std::setw(display_val + 3) << "Iter. N°" << std::setw(display_val) << "Code" << std::setw(display_val + 5) << "JAMCRC value" << std::setw(display_val + 5) + std::cout << std::setw(display_val + 3) << "Iter. N°" << std::setw(display_val) << "Code" << std::setw(display_val + 10) << "JAMCRC value" << std::setw(display_val + 5) << "Associated code" << std::endl; for (auto &result : results) { - std::cout << std::setw(display_val + 3) << result.index << std::setw(display_val) << result.code << std::setw(display_val) << "0x" << std::hex + std::cout << std::setw(display_val + 2) << std::dec << result.index << std::setw(display_val + 5) << result.code << std::setw(display_val) << "0x" << std::hex << result.jamcrc << std::setw(display_val) << result.associated_code << std::endl; } std::cout << "Time: " << std::chrono::duration_cast>(end_time - begin_time).count() << " sec" << std::endl; // Display time @@ -69,4 +67,5 @@ void GTA_SA_Virtual::printResult() const { << (static_cast(max_range - min_range) / std::chrono::duration_cast>(end_time - begin_time).count()) / 1000000 << " MOps/sec" << std::endl; // Display perf std::cout << "" << std::endl; + std::cout << "Number of results: " << std::dec << results.size() << std::endl; } diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index d113dec..482eb0d 100755 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -61,7 +61,7 @@ if(NOT WIN32) # Add tests test_bench_generator(find_string_inv_test true true) test_bench_generator(jamcrc_test true true) - #add_test_fn(gta_sa_test) + test_bench_generator(gta_sa_test true true) # Add bench test_bench_generator(find_string_inv_bench false false) diff --git a/test/source/test/gta_sa_test.cpp b/test/source/test/gta_sa_test.cpp index cf471d3..100613c 100644 --- a/test/source/test/gta_sa_test.cpp +++ b/test/source/test/gta_sa_test.cpp @@ -23,53 +23,74 @@ 535721682 ASBHGRB 0xa7613f99 */ -TEST(GTA_SA_OPENMP, basic_calc_base_1) { - GTA_SA_OPENMP gta_sa; +TEST(GTA_SA_STDTHREAD, basic_calc_base_1) { + GTA_SA_STDTHREAD gta_sa; gta_sa.min_range = 0; gta_sa.max_range = 60000; gta_sa.run(); - EXPECT_EQ(gta_sa.results.results.size(), 0); + EXPECT_EQ(gta_sa.results.size(), 0); } -TEST(GTA_SA_OPENMP, basic_calc_base_2) { - GTA_SA_OPENMP gta_sa; +TEST(GTA_SA_STDTHREAD, basic_calc_base_2) { + GTA_SA_STDTHREAD gta_sa; gta_sa.min_range = 20810700; gta_sa.max_range = 20810800; gta_sa.run(); - EXPECT_EQ(gta_sa.results.results.size(), 1); - EXPECT_EQ(std::get<0>(gta_sa.results.results[0]), 20810792); - EXPECT_EQ(std::get<1>(gta_sa.results.results[0]), "ASNAEB"); - EXPECT_EQ(std::get<2>(gta_sa.results.results[0]), 0x555fc201); + EXPECT_EQ(gta_sa.results.size(), 1); + EXPECT_EQ(gta_sa.results[0].index, 20810792); + EXPECT_EQ(gta_sa.results[0].code, "ASNAEB"); + EXPECT_EQ(gta_sa.results[0].jamcrc, 0x555fc201); } -TEST(GTA_SA_OPENMP, basic_calc_base_3) { - GTA_SA_OPENMP gta_sa; +TEST(GTA_SA_STDTHREAD, basic_calc_base_3) { + GTA_SA_STDTHREAD gta_sa; gta_sa.min_range = 181961000; gta_sa.max_range = 181961100; gta_sa.run(); - EXPECT_EQ(gta_sa.results.results.size(), 1); - EXPECT_EQ(std::get<0>(gta_sa.results.results[0]), 181961057); - EXPECT_EQ(std::get<1>(gta_sa.results.results[0]), "OHDUDE"); - EXPECT_EQ(std::get<2>(gta_sa.results.results[0]), 0xe958788a); + EXPECT_EQ(gta_sa.results.size(), 1); + EXPECT_EQ(gta_sa.results[0].index, 181961057); + EXPECT_EQ(gta_sa.results[0].code, "OHDUDE"); + EXPECT_EQ(gta_sa.results[0].jamcrc, 0xe958788a); } -TEST(GTA_SA_OPENMP, basic_calc_base_4) { - GTA_SA_OPENMP gta_sa; +TEST(GTA_SA_STDTHREAD, basic_calc_base_4) { + GTA_SA_STDTHREAD gta_sa; gta_sa.min_range = 299376700; gta_sa.max_range = 299376800; gta_sa.run(); - EXPECT_EQ(gta_sa.results.results.size(), 1); - EXPECT_EQ(std::get<0>(gta_sa.results.results[0]), 299376767); - EXPECT_EQ(std::get<1>(gta_sa.results.results[0]), "YECGAA"); - EXPECT_EQ(std::get<2>(gta_sa.results.results[0]), 0x40cf761); + EXPECT_EQ(gta_sa.results.size(), 1); + EXPECT_EQ(gta_sa.results[0].index, 299376767); + EXPECT_EQ(gta_sa.results[0].code, "YECGAA"); + EXPECT_EQ(gta_sa.results[0].jamcrc, 0x40cf761); +} + +TEST(GTA_SA_STDTHREAD, basic_calc_base_5) { + GTA_SA_STDTHREAD gta_sa; + gta_sa.min_range = 20810700; + gta_sa.max_range = 147491500; + + gta_sa.run(); + + EXPECT_EQ(gta_sa.results.size(), 3); + EXPECT_EQ(gta_sa.results[0].index, 20810792); + EXPECT_EQ(gta_sa.results[0].code, "ASNAEB"); + EXPECT_EQ(gta_sa.results[0].jamcrc, 0x555fc201); + + EXPECT_EQ(gta_sa.results[1].index, 75396850); + EXPECT_EQ(gta_sa.results[1].code, "FHYSTV"); + EXPECT_EQ(gta_sa.results[1].jamcrc, 0x44b34866); + + EXPECT_EQ(gta_sa.results[2].index, 147491485); + EXPECT_EQ(gta_sa.results[2].code, "LJSPQK"); + EXPECT_EQ(gta_sa.results[2].jamcrc, 0xfeda77f7); } #if defined(_OPENMP) @@ -80,11 +101,12 @@ TEST(GTA_SA_OPENMP, basic_calc_mode_openmp) { gta_sa.run(); - EXPECT_EQ(gta_sa.results.results.size(), 0); + EXPECT_EQ(gta_sa.results.size(), 0); } #endif #if defined(BUILD_WITH_CUDA) + TEST(GTA_SA_CUDA, basic_calc_cuda_1) { GTA_SA_CUDA gta_sa; gta_sa.min_range = 0; @@ -92,11 +114,51 @@ TEST(GTA_SA_CUDA, basic_calc_cuda_1) { gta_sa.run(); - EXPECT_EQ(gta_sa.results.results.size(), 0); + EXPECT_EQ(gta_sa.results.size(), 0); +} + +TEST(GTA_SA_CUDA, basic_calc_cuda_2) { + GTA_SA_CUDA gta_sa; + gta_sa.min_range = 20810700; + gta_sa.max_range = 20810800; + + gta_sa.run(); + + EXPECT_EQ(gta_sa.results.size(), 1); + EXPECT_EQ(gta_sa.results[0].index, 20810792); + EXPECT_EQ(gta_sa.results[0].code, "ASNAEB"); + EXPECT_EQ(gta_sa.results[0].jamcrc, 0x555fc201); +} + +TEST(GTA_SA_CUDA, basic_calc_cuda_3) { + GTA_SA_CUDA gta_sa; + gta_sa.min_range = 181961000; + gta_sa.max_range = 181961100; + + gta_sa.run(); + + EXPECT_EQ(gta_sa.results.size(), 1); + EXPECT_EQ(gta_sa.results[0].index, 181961057); + EXPECT_EQ(gta_sa.results[0].code, "OHDUDE"); + EXPECT_EQ(gta_sa.results[0].jamcrc, 0xe958788a); } + +TEST(GTA_SA_CUDA, basic_calc_cuda_4) { + GTA_SA_CUDA gta_sa; + gta_sa.min_range = 299376700; + gta_sa.max_range = 299376800; + + gta_sa.run(); + + EXPECT_EQ(gta_sa.results.size(), 1); + EXPECT_EQ(gta_sa.results[0].index, 299376767); + EXPECT_EQ(gta_sa.results[0].code, "YECGAA"); + EXPECT_EQ(gta_sa.results[0].jamcrc, 0x40cf761); +} + #endif auto main(int argc, char **argv) -> int { - ::testing::InitGoogleTest(&argc, argv); - return RUN_ALL_TESTS(); + ::testing::InitGoogleTest(&argc, argv); + return RUN_ALL_TESTS(); }