Witam wszystkich!
Jak brzmi mój problem każdy wie już pewnie z samego tytułu, więc przejdę do szczegółów - w ramach rozwijania mojego programu szukającego liczb pierwszych postanowiłem , aby główne obliczenia wykonywane były na mojej karcie graficznej.
Dane:
IDE: VS2017
Cuda Toolkit: 9 RC
Karta graficzna: NVidia GeForce 920M 2GB lub 4GB (nie jestem pewien ile dokładnie)
Czas na kod (wytłumaczenie pod nim):
kernel.cuh
#pragma once
#include <cuda.h>
#include <cuda_runtime_api.h>
#include <device_launch_parameters.h>
#include "C_primes_search.h"
__global__ void primes_search_kernel( unsigned __int64 num_of_nums_to_check, unsigned __int64 divisors_arr_size, unsigned __int64 primes_arr_size, unsigned __int64 size_of_block, unsigned __int64 last_num, unsigned __int64 * divisors_arr, unsigned __int64 * primes_arr, unsigned __int64 * primes_num );
kernel.cu
__global__ void primes_search_kernel( unsigned __int64 num_of_nums_to_check, unsigned __int64 divisors_arr_size, unsigned __int64 primes_arr_size, unsigned __int64 size_of_block, unsigned __int64 last_num, unsigned __int64 * divisors_arr, unsigned __int64 * primes_arr, unsigned __int64 * primes_num ) {
int
id = blockIdx.x * size_of_block + threadIdx.x;
unsigned __int64
* div_arr = divisors_arr + sizeof( unsigned __int64 ) * blockIdx.x * divisors_arr_size,
* pri_arr = primes_arr + sizeof( unsigned __int64 ) * id * primes_arr_size,
& pri_num = *( primes_num + sizeof( unsigned __int64 ) * id ),
num_checking_limit = last_num + num_of_nums_to_check *( id + 1 ),
square_root,
reminder,
div_ord;
pri_num = 0;
for( unsigned __int64 num_in_checking = last_num + num_of_nums_to_check * id; num_in_checking < num_checking_limit; num_in_checking++ ) {
div_ord = 0;
square_root = __dsqrt_ru( static_cast < double >( num_in_checking ) );
do {
reminder = num_in_checking % div_arr[ div_ord ];
} while( reminder != 0 && div_arr[ div_ord++ ] <= square_root );
if( reminder != 0 ) {
pri_arr[ pri_num++ ] = num_in_checking;
}
}
}
C_primes_search.h
#pragma once
#include <fstream>
#include <iostream>
#include <time.h>
#include <cuda.h>
#include <math.h>
#include <cuda_runtime_api.h>
using namespace std;
class C_primes_search {
private:
time_t
beg_t, end_t;
fstream
data_flow;
public:
unsigned __int64
* divisors_arr = nullptr,
* primes_arr = nullptr,
* primes_num = nullptr,
num_of_cycles,
* divisors_buffer,
last_num,
primes_per_cycle;
const size_t
num_of_blocks = 16,
size_of_block = 16,
divisors_arr_size = 1000000,
num_of_nums_to_check = 5000000,
primes_arr_size = static_cast < size_t >( ceil( num_of_nums_to_check * 0.06457502562L ) + 1000 );
void
run(),
divisors_read(),
last_num_read(),
file_save(),
cuda_kernel_launch(),
nullptr_check( void *, int ),
operator delete[]( void * ),
check_for_cuda_error( cudaError_t, int );
void *
operator new[]( size_t );
C_primes_search();
~C_primes_search();
};
C_primes_search.cpp
#include "C_primes_search.h"
#include "kernel.cuh"
#include <vcruntime_new.h>
void C_primes_search::run() {
divisors_read();
last_num_read();
data_flow.seekp( 0, data_flow.end );
for( unsigned __int64 cycle = 0; cycle < num_of_cycles; cycle++ ) {
cout << "<----------------->" << endl;
cout << "Cykl numer " << cycle << endl;
beg_t = clock();
cuda_kernel_launch();
end_t = clock() - beg_t;
file_save();
cout << "Zakres: " << last_num << " - ";
last_num += num_of_blocks * size_of_block * 500000;
cout << last_num - 1 << endl;
cout << "Znalezionych liczb pierwszych: " << primes_per_cycle << endl;
cout << "Czas cyklu: " << static_cast < double >( end_t ) / CLOCKS_PER_SEC << endl;
cout << "Czas rzeczywisty: " << static_cast < double >( end_t ) * 1000000 /( primes_per_cycle * CLOCKS_PER_SEC ) << endl;
}
}
void C_primes_search::operator delete[]( void * ptr ) {
cudaFree( ptr );
}
void C_primes_search::nullptr_check( void * ptr, int line ) {
if( ptr == nullptr ) {
cout << "nullptr - linia: " << line << endl;
}
}
void * C_primes_search::operator new[]( size_t size ) {
void * ptr;
cudaMallocManaged(( void ** ) & ptr, size );
return ptr;
}
void C_primes_search::cuda_kernel_launch() {
check_for_cuda_error( cudaSetDevice( 0 ), __LINE__ );
check_for_cuda_error( cudaDeviceReset(), __LINE__ );
primes_search_kernel << < num_of_blocks, size_of_block >> >(
num_of_nums_to_check,
divisors_arr_size,
primes_arr_size,
size_of_block,
last_num,
divisors_arr,
primes_arr,
primes_num );
check_for_cuda_error( cudaDeviceSynchronize(), __LINE__ );
}
void C_primes_search::check_for_cuda_error( cudaError_t state, int line ) {
if( state != cudaSuccess ) {
cout << "Cuda error - linia " << line << " - " << cudaGetErrorString( state ) << endl;
}
}
void C_primes_search::file_save() {
primes_per_cycle = 0;
for( unsigned __int64 index = 0; index < num_of_blocks * size_of_block; index++ ) {
primes_per_cycle += primes_num[ index ];
data_flow.write(( char * ) primes_arr[ index * primes_arr_size ], primes_num[ index ] * sizeof( unsigned __int64 ) );
}
}
void C_primes_search::last_num_read() {
data_flow.seekg( - 8, data_flow.end );
data_flow.read(( char * ) & last_num, sizeof( last_num ) );
last_num++;
}
C_primes_search::C_primes_search() {
nullptr_check( primes_num = new unsigned __int64[ num_of_blocks * size_of_block ], __LINE__ );
nullptr_check( divisors_arr = new unsigned __int64[ num_of_blocks * divisors_arr_size ], __LINE__ );
nullptr_check( primes_arr = new unsigned __int64[ num_of_blocks * size_of_block * primes_arr_size ], __LINE__ );
data_flow.open( "pierwsze.bin", ios::binary | ios::in | ios::out | ios::app | ios::ate );
if( !data_flow.is_open() ) {
throw "Plik pierwsze.bin nie otworzyl sie";
}
cout << "<----------------->" << endl;
cout << "Podaj liczbe cykli: ";
cin >> num_of_cycles;
}
void C_primes_search::divisors_read() {
for( unsigned __int64 blockID = 0; blockID < num_of_blocks; blockID++ ) {
data_flow.seekg( 0, data_flow.beg );
data_flow.read(( char * ) & divisors_arr[ blockID * divisors_arr_size ], divisors_arr_size * sizeof( unsigned __int64 ) );
}
}
C_primes_search::~C_primes_search() {
data_flow.close();
delete[] divisors_arr;
delete[] primes_arr;
delete[] primes_num;
}
Program wykorzystuje najwydajniejszą wersję
metody naiwnej (wiem, że istniej sito erostenesa, ale wolę, żeby było jak jest).
Szczegóły działania:
Program pobiera do pamięci milion pierwszych liczb pierwszych, a potem każdy z wątków sprawdza po 5 milionów liczb, po czym pierwsze zostają zapisane do pliku (np. wątek1 sprawdza od 20000000 do 25000000, wątek 2 od 25000000 do 30000000 itd.). Na początku użytkownik podaje ile takich cykli, w których działa zestaw wątków, ma zostać wykonanych.
Program kompiluje się i działa do momentu dotarcia do zaznaczonego w kodzie punktu, po czym zaczyna wywalać błędy naruszenia dostępu do pamięci (najprawdopodobniej jest jakiś błąd w funkcji wykonywanej na GPU, ale nie mogę go znaleźć).
Starałem się także, aby kod był "samoopisujący się" (żeby nie było jakichś dziwnych jednoliterowych zmiennych, czy coś), lecz jeżeli coś jest niejasne to służę wyjaśnieniami. Kod jest pełny i nic nie było wycinane przeze mnie, więc jak ktoś chce przetestować to nie będzie z tym problemów.
Drogi użytkowniku jeżeli dotarłeś do tego momentu i nie jesteś zamęczony na śmierć to proszę o pomoc ;D
PS. Bym zapomniał -
https://drive.google.com/file/d/0BwLxpuSgHpzsOENET1BFcE9fU1E/view do pobrania pliku potrzebnego do działania.