W moim głównym programie mam dużą tablicę znaków, którą kopiuję porcjami do pamięci urządzenia. W swoim programie uruchamiam około 500 000 wątków, a każdy wątek uzyskuje dostęp do 2000 znaków. Więc przenoszę 500 000 * 2000 = 1 GB bajtów na raz z kodem

err = cudaMemcpy (dev_database, adjusted_database[k], JOBS * 2000 * sizeof(char), cudaMemcpyHostToDevice);
if(err != cudaSuccess) { printf("CUDA error: %s\n", cudaGetErrorString(err)); exit(EXIT_FAILURE); }

W moim jądrze definiuję również trzy współdzielone tablice

//__shared__ char dev_query[200];
__shared__ float dev_scores[200*5];
__shared__ int dev_index[26];

I zainicjuj je za pomocą

if(threadIdx.x == 0) { 
  //for(i = 0; i < 200; i++){ dev_query[i] = dev_query_constant[i]; }
  for(i = 0; i < 200 * 5; i++){ dev_scores[i] = dev_scores_constant[i]; }
  for(i = 0; i < 26; i++){ dev_index[i] = dev_index_constant[i]; }
}
__syncthreads(); 

Jeśli uruchomię program z dwoma wierszami zakomentowanymi, moje jądro zwróci dziwne wartości, a po skopiowaniu drugiego fragmentu tablicy znaków pojawia się błąd

Błąd CUDA: nieokreślony błąd uruchamiania

Jeśli odkomentuję wiersze w kodzie powyżej, wszystko działa dobrze. Jeśli skopiuję mniejsze fragmenty tablicy, takie jak 100 MB zamiast 1 GB, działa dobrze, dopóki nie dotrę do szóstego fragmentu, w którym otrzymuję te same błędy, co powyżej.

To bardzo dziwne zachowanie i chciałbym zrozumieć, dlaczego tak się dzieje. Czy jest gdzieś błąd, który to powoduje? Trudno to wskazać, ponieważ program działa dobrze, jeśli przeniosę mały fragment (np. 100 MB) i zignoruję pozostałe. Działa również dobrze, jeśli odkomentuję linie związane ze współdzielonymi zmiennymi lub zmienię współdzielone zmienne na stałe. Każda pomoc byłaby bardzo mile widziana. Dziękuję!

EDYTOWAĆ: Oto moje jądro. Podsumowując, obliczam wynik podobieństwa dwóch ciągi, porównując ich i-ty znak dla wszystkich i między 0 a ich długością. Poniższy kod spowoduje powyższy błąd chyba że odkomentujesz linię natychmiast po if(threadIdx.x == 0) {. Lub jeśli ty zamień poniższe wspólne tablice na stałe, wtedy również działa dobrze.

__global__ void assign7(int jobs_todo, char* database, float* results, int flag) {
unsigned int id = threadIdx.x + blockIdx.x * blockDim.x;

if(id < jobs_todo) {
__shared__ char dev_query[200];
__shared__ float dev_pos_specific_scores[200*5];
__shared__ int dev_subst_index[26];

int j_, i, p, stop, k; //stop2;
float score=0, max=0;
char ch; //ch1, ch2;

if(threadIdx.x == 0) {
//for(i = 0; i < 51; i++){ dev_query[i] = dev_query_constant[i]; }
  for(i = 0; i < 5 * 200; i++){ dev_pos_specific_scores[i] = dev_pos_specific_scores_constant[i]; }
  for(i = 0; i < 26; i++){ dev_subst_index[i] = dev_subst_index_constant[i]; }
}
__syncthreads(); 

for(i = 1; i <= 2000 - 51; i += 1){
  p = jobs_todo*(i-1);
  score = 0;
  stop = 51/1; stop = stop*1;
  for(j_ = 1; j_ <= stop; j_ += 1){
    k = (j_-1)*5;
    ch = database[p + id];
    score += dev_pos_specific_scores[k + dev_subst_index[ch - 'A']];
    if(score < 0) score = 0;
    if(score > max) max = score;                                      
    p += jobs_todo;
  }
}
results[id] = max;
}
}
0
Ross 28 luty 2012, 02:00

2 odpowiedzi

Najlepsza odpowiedź

W danych było kilka znaków, które spowodowały, że dev_index[ch-'A'] zwróciło -1. Ten utworzył indeks dev_scores -1, gdy k = 0. Uważam, że to było źródło błąd pamięci w moim kodzie. Komentowałem wszystko i stopniowo odkomentowałem fragmenty. Teraz działa dobrze. Dziękuję @talonmies, @harrism i @perreal za komentarze!

0
Ross 1 marzec 2012, 18:05

Poniższa część używa k bez jej inicjowania:

ch = database[p + id];
score += dev_scores[k + dev_index[ch - 'A']];

To nie ma znaczenia, ale ta część:

if(threadIdx.x == 0) { 
  //for(i = 0; i < 200; i++){ dev_query[i] = dev_query_constant[i]; }
  for(i = 0; i < 200 * 5; i++){ dev_scores[i] = dev_scores_constant[i]; }
  for(i = 0; i < 26; i++){ dev_index[i] = dev_index_constant[i]; }
}

Można zmienić na coś takiego:

if(threadIdx.x < 200) {
  // dev_query[i] = dev_query_constant[i];
}

if(threadIdx.x < 200 * 5) { // or iterate whole block 5 times..
  dev_scores[i] = dev_scores_constant[i];
}
...
2
perreal 28 luty 2012, 03:44