К сожалению, ваш вывод о том, что код работает при компиляции для sm_35 и запуске на sm_30 GPU, неверен. Виновник вот в чем:
void cbow_cuda(long window, long negative, float alpha, long sentence_length,
int *sen, long layer1_size, float *syn0, long hs, float *syn1,
float *expTable, int *vocab_codelen, char *vocab_code,
int *vocab_point, int *table, long table_size,
long vocab_size, float *syn1neg){
int blockSize = 256;
int gridSize = (sentence_length)/(blockSize/32);
size_t smsize = (blockSize/32)*(2*layer1_size+3)*sizeof(float);
//printf("sm size is %d\n", smsize);
//fflush(stdout);
cbow_kernel<1><<<gridSize, blockSize, smsize>>>
(window, negative, alpha, sentence_length, sen,
layer1_size, syn0, syn1, expTable, vocab_codelen,
vocab_code, vocab_point, table, table_size,
vocab_size, syn1neg);
}
Этот код автоматически завершится ошибкой, если запуск ядра завершится ошибкой из-за неполной проверки ошибок API. И запуск ядра невозможен, если вы собираете для sm_35 и запускаете на sm_30. Если вы измените код этой функции на этот (добавив проверку ошибок запуска ядра):
void cbow_cuda(long window, long negative, float alpha, long sentence_length,
int *sen, long layer1_size, float *syn0, long hs, float *syn1,
float *expTable, int *vocab_codelen, char *vocab_code,
int *vocab_point, int *table, long table_size,
long vocab_size, float *syn1neg){
int blockSize = 256;
int gridSize = (sentence_length)/(blockSize/32);
size_t smsize = (blockSize/32)*(2*layer1_size+3)*sizeof(float);
//printf("sm size is %d\n", smsize);
//fflush(stdout);
cbow_kernel<1><<<gridSize, blockSize, smsize>>>
(window, negative, alpha, sentence_length, sen,
layer1_size, syn0, syn1, expTable, vocab_codelen,
vocab_code, vocab_point, table, table_size,
vocab_size, syn1neg);
checkCUDAError( cudaPeekAtLastError() );
}
и скомпилируйте и запустите его для sm_35, вы должны получить это на устройстве sm_30:
~/cbow/word2vec_cbow$ make
nvcc word2vec.cu -o word2vec -O3 -Xcompiler -march=native -w -Xptxas="-v" -arch=sm_35 -lineinfo
ptxas info : 0 bytes gmem
ptxas info : Compiling entry function '_Z11cbow_kernelILx1EEvllflPKilPVfS3_PKfS1_PKcS1_S1_llS3_' for 'sm_35'
ptxas info : Function properties for _Z11cbow_kernelILx1EEvllflPKilPVfS3_PKfS1_PKcS1_S1_llS3_
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 34 registers, 448 bytes cmem[0], 8 bytes cmem[2]
~/cbow/word2vec_cbow$ ./word2vec -train text8 -output vectors.bin -cbow 1 -size 200 -window 7 -negative 1 -hs 1 -sample 1e-3 -threads 1 -binary 1 -save-vocab voc #> out 2>&1
Starting training using file text8
Vocab size: 71290
Words in train file: 16718843
vocab size = 71290
cbow.cu 114 : invalid device function
т.е. запуск ядра завершился неудачно, так как в полезной нагрузке CUDA cubin вашего приложения не было найдено соответствующего кода устройства. Это также отвечает на ваш более ранний вопрос о том, почему вывод этого кода неверен. Ядро анализа просто никогда не запускается на вашем оборудовании, если оно построено с параметрами по умолчанию.
Если я соберу этот код для sm_30 и запущу его на GTX 670 с 2 ГБ памяти (вычислительная способность 3.0), я получу следующее:
~/cbow/word2vec_cbow$ make
nvcc word2vec.cu -o word2vec -O3 -Xcompiler -march=native -w -Xptxas="-v" -arch=sm_30 -lineinfo
ptxas info : 0 bytes gmem
ptxas info : Compiling entry function '_Z11cbow_kernelILx1EEvllflPKilPVfS3_PKfS1_PKcS1_S1_llS3_' for 'sm_30'
ptxas info : Function properties for _Z11cbow_kernelILx1EEvllflPKilPVfS3_PKfS1_PKcS1_S1_llS3_
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 34 registers, 448 bytes cmem[0], 12 bytes cmem[2]
~/cbow/word2vec_cbow$ ./word2vec -train text8 -output vectors.bin -cbow 1 -size 200 -window 7 -negative 1 -hs 1 -sample 1e-3 -threads 1 -binary 1 -save-vocab voc #> out 2>&1
Starting training using file text8
Vocab size: 71290
Words in train file: 16718843
vocab size = 71290
Alpha: 0.000009 Progress: 100.00% Words/thread/sec: 1217.45k
т.е. код работает правильно до завершения без каких-либо ошибок. Я не могу сказать вам, почему вы не можете запустить код на своем оборудовании, потому что я не могу воспроизвести вашу ошибку на своем оборудовании. Вам нужно будет выполнить некоторую отладку самостоятельно, чтобы найти основную причину этого.
person
Community
schedule
17.07.2015
-lineinfo
, чтобы заставить cuda-memcheck сообщать номера строк. Я не могу объяснить, не видя кода, и ваша работа — найти код. Я не буду анализировать для вас 1000+ строк кода, извините. - person talonmies   schedule 15.07.2015