CUDA + Ubuntu 9.10

На работе появилась платка с поддержкой CUDA. Чтобы заставить корректно компилироваться код на C++ под Ubuntu 9.10 пришлось немного повозиться. Проблема в том, что на сегодня программное обеспечение для CUDA поддерживает только версии 8.10 и 9.04. Это связано с компилятором g++, который должен иметь версию 4.3, а в новой Ubuntu — 4.4.
Поэтому, при компиляции возникают ошибки:

desktop:~/cuda$ nvcc --host-compilation=c++ cap.cu -o cucap
/usr/include/c++/4.4/ext/atomicity.h(46): error: identifier "__sync_fetch_and_add" is undefined
/usr/include/c++/4.4/ext/atomicity.h(50): error: identifier "__sync_fetch_and_add" is undefined
2 errors detected in the compilation of "/tmp/tmpxft_00006255_00000000-4_cap.cpp1.ii".

Решением проблемы стала установка дополнительно g++ версии 4.3, которая, оказывается, умеет мирно сосуществовать с 4.4. Для этого достаточно выполнить команду:

sudo apt-get install g++-4.3

После этого компиляция заработала с добавлением ключа --compiler-bindir=/usr/bin/gcc-4.3:

nvcc --host-compilation=c++ cudaTest.cu --compiler-bindir=/usr/bin/gcc-4.3

Для того, чтобы проверить работоспособность CUDA, я написал небольшую программку, которая читает текст из файла, а, затем, преобразует все буквы в верхний регистр.

#include <iostream>
#include <iterator>
#include <fstream>

using namespace std;

__global__ void setCAPS(unsigned char *a) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if(a[i]>96 && a[i]<123) {
       a[i] ^= 32;
   }
}

int main() {
ifstream f("SonnetVII.txt", fstream::in|fstream::binary);
f >> noskipws;
f.seekg(0, ios::end);
int fileSize = f.tellg();
f.seekg(0, ios::beg);
unsigned char *a = new unsigned char[fileSize];
copy(istream_iterator<unsigned char> (f), istream_iterator<unsigned char>(), a);
copy(a, a + fileSize, ostream_iterator<unsigned char>(cout,""));
cout << endl;

unsigned char *devMem = NULL;
cudaMalloc((void **)&devMem, fileSize);
cudaMemcpy(devMem, a, fileSize, cudaMemcpyHostToDevice);

dim3 threads = dim3(512, 1);
dim3 blocks = dim3(fileSize / threads.x + 1, 1);
setCAPS<<<blocks, threads>>> (devMem);

cudaMemcpy(a, devMem, fileSize, cudaMemcpyDeviceToHost);
cudaFree(devMem);

copy(a, a + fileSize, ostream_iterator<unsigned char>(cout,""));
cout << endl;
delete[](a);
return 0;
}

Результат работы:

Sonnet VII

Lo! in the orient when the gracious light
Lifts up his burning head, each under eye
Doth homage to his new-appearing sight,
Serving with looks his sacred majesty;
And having climb'd the steep-up heavenly hill,
Resembling strong youth in his middle age,
yet mortal looks adore his beauty still,
Attending on his golden pilgrimage;
But when from highmost pitch, with weary car,
Like feeble age, he reeleth from the day,
The eyes, 'fore duteous, now converted are
From his low tract and look another way:
So thou, thyself out-going in thy noon,
Unlook'd on diest, unless thou get a son.

SONNET VII

LO! IN THE ORIENT WHEN THE GRACIOUS LIGHT
LIFTS UP HIS BURNING HEAD, EACH UNDER EYE
DOTH HOMAGE TO HIS NEW-APPEARING SIGHT,
SERVING WITH LOOKS HIS SACRED MAJESTY;
AND HAVING CLIMB'D THE STEEP-UP HEAVENLY HILL,
RESEMBLING STRONG YOUTH IN HIS MIDDLE AGE,
YET MORTAL LOOKS ADORE HIS BEAUTY STILL,
ATTENDING ON HIS GOLDEN PILGRIMAGE;
BUT WHEN FROM HIGHMOST PITCH, WITH WEARY CAR,
LIKE FEEBLE AGE, HE REELETH FROM THE DAY,
THE EYES, 'FORE DUTEOUS, NOW CONVERTED ARE
FROM HIS LOW TRACT AND LOOK ANOTHER WAY:
SO THOU, THYSELF OUT-GOING IN THY NOON,
UNLOOK'D ON DIEST, UNLESS THOU GET A SON.

Интересно, что по производительности CUDA и CPU-версии одинаковы, но CUDA забирает на себя дополнительные 100 мс. При увеличении размера файла, CPU-версия начинает заметно выигрывать, вероятно, за счет копирования данных с девайса на хост и обратно. А вот когда я модифицировал функцию преобразования регистра вот так (аналогично была модифицирована и CPU-версия):

__global__ void setCAPS(unsigned char *a) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if(a[i]>96 && a[i]<123) {
      for(int k = 0; k < 1000001; ++k) {
          a[i] ^= 32;
      }
  }
}

получились следующие результаты:

desktop:~/cuda$ time ./cudacaps
real 0m0.294s
user 0m0.068s
sys 0m0.224s
desktop:~/cuda$ time ./cpucaps
real 0m1.758s
user 0m1.732s
sys 0m0.000s

Т.е. использование CUDA эффективно в случае, когда большая часть вычислений выполняется на видеокарте без копирования данных в память компьютера.