C++/CUDA - vprašanje

Pri svojem delu v C++ veckrat pisem algoritme za obdelavo vec MB dolgih arrayev, ki jih optimiziram z SSE/AVX ukazi. Meritve z zanesljivimi orodji so pokazale, da tipična hitrost neparalelnega algoritma, ki jo s tem dosežem, znaša 2 ns / element arraya.

Na enem projektu bi moral v rednih casovnih intervalih taksen algoritem izvesti na (vsakic novih) 10 arrayih velikosti 5MB, torej za vse skupaj porabim priblizno 100ms casa. Profilerji kažejo, da se slaba polovica časa uporabi za store operacije, preostanek vzame pa sam algoritem oz. preostali ukazi.
Prekrasno bi bilo, če bi ta čas uspel zmanjšati za cca trikrat.

Ali se iz povedanega da ocenit, če se je vredno usmeriti v iskanje rešitev s programiranjem na grafični, ali pa se morda že v štartu vidi, da se to v tem primeru ne splača?

Algoritem izvaja interpolacijo s pomocjo koordinat; v vec knjiznicah temu recejo remapping (ne vem pa kaksen je splosen izraz). Psevdokoda:

void interpolation(vector const& source, vector const& coordinates, vector& output) 
{
for(i = 0; i < N; ++i)
{
  c = coordinates[i];
  low = floor(c);
  high = ceil(c);
  output[i] = (high - c)*source[low] + (c - low)*source[high];
}
}

Operacije torej SO neodvisne druga od druge (in to seveda ze izkoriscam v implementaciji z AVX). Morda v prvem postu nisem bil jasen - z neparalelnostjo algoritmov sem imel v mislih to, da tecejo le v enem threadu. Operacije pa se z AVX ukazi seveda izvajajo na SIMD (single instruction multiple data) nacin.

A je mozno vnaprej "cez palec" (red velikosti) oceniti kaksen delay prinese prenasanje podatkov na GPU in nazaj?

Hvala tudi za zadnji nasvet. Aja, pa se to: do estimatea 2 ns / element sem prisel tako, da sem med delovanjem programa racunal povprecen cas trajanja celotnega algoritma, nato pa povprecen cas delili s stevilom elementov. Nisem benchmarkal vsake operacije posebej. Sem pa meritve ponovil z arrayi razlicnih velikosti in prisel vedno do enake ocene.

Hvala za te nasvete, prisli bodo zelo prav.

PS: izvajanje algoritma v vec threadih sem v startu crtal, ker to ne bi resilo vseh tezav. Prevelik latency algoritma (ki je sicer primaren obstacle) ni edini problem; tezave so tudi zaradi tega ker bi uvedba tega algoritma porabila znaten prevelik delez procesorske moci glede na to, ki nam je trenutno na voljo.

Zal ne morem povedati kaj vec, ker sem se v pogodbi zavezal, da tega ne bom pocel. :)

Aja, to pa ja: racunalniski vid in obdelava signalov.

Glede na zgornjo psevdo kodo lahko spakiraš podatke v teksturo, jo uploadaš na GPU in uporabiš texture filtering hardware, ki ga imajo GPU-ji... Cela tistale tvoja koda se spremeni efektivno v eno vrstico:

__global__ void interpolation(cudaTextureObject_t source, float *coordinates, float *output)
{
    const int i = blockIdx.x * blockDim.x + threadIdx.x;
    output[i] = tex1Dfetch<float>(source, coordinates[i]); // 1D koordinate
}

Zadeva bo tolk hitra kot boš imel na razpolago bandwidtha na GPU-ju. Trivialen case v bisvu (še ALU-jev praktično ne rabiš).