IDE для CUDA

Хочу выбрать IDE для CUDA, работаю под Kubuntu, больше всего нравится Eclipse CDT, однако два найденных в интернете руководства не смог осилить, делаю всё по инструкции, но не получается результат как у авторов, к тому же они привязываются к CUDA SDK, к его путям и make-файлу common.mk, где всё на нем завязано, а мне бы хотелось использовать свои.
С NetBeans тоже не получилось, там еще меньше инфы нашел. Есть реализация FindCUDA для CMake, но мне бы хотелось нормальную среду с эдитором, спейсингом и дебагом, а как в cmake дебажить, не знаю, консольный дебаг-не вариант.
Подскажите, кто на чем работает, есть ли работающие инструкции?
Собственно, задача - переписать существующее ПО цифровой обработки изображений под CUDA. Ранее использовалась библиотека GENIAL, удобная для работы с матрицами, можно было работать местами почти как в MATLAB, операции +,-,*,/ работали с матрицами как и с числами.
Однако посмотрев на классическую CUDA, понял, что там всё не так устроено, хотелось бы более удобный вариант, чтоб акцент был на функциональной части. CUBLAS - не подходит, мощные функции, а простые трудно реализовать. Есть "обертки" над CUBLAS. Я обратил внимание на две:
Thrust - обертка , аналогичная по интерфейсу STL, местами удобная, но немного странная как для меня
VSIPL - мощная библиотека, судя по описанию, использует параллельные вычисления, в том числе и CUDA, но явно не указано, в какой степени оно зависит от CUDA, функций великое множество, ближе всего по реализации к GENIAL.
Если кто-то ими интересовался или использовал, или знает другую библиотеку, подскажите, пожалуйста.
Или может имеет смысл перейти на OpenCL, правда не знаю, насколько он сейчас самостоятелен. Критериями является быстродействие и удобство работы с матрицами, наличие простых и сложных операций и функций.

Forums: 

да, пока 0 ответов. Все же на

да, пока 0 ответов. Все же на чем-то пишут. Я взял findCUDA-проект, правда ничего не понимаю в CMAke и по одной инструкции сделал связку CMake и Eclipse, где в Eclipse создан проект с внешним makefile - ом, который генерируется как раз с помощью CMake. Причем CMake был добавлен как external tool приложение, которое я сконфигурировал и запускаю, когда нужно переделать makefile, после того как я редактирую файл CMakeList.txt. Первое приложение я написал на GPU VSIPL, при этом всё прекрасно компилилось и билдилось, даже получил прирост в 5 раз по сравнению с обычным исполнением на CPU. По идее он должен был использовать CUDA, так заявляют разработчики, однако я решил перейти на чистую CUDA, так как мне не понравился такой слабый прирост GPU VSIPL, и тут столкнулся с проблемами. У меня выпадают вот такие ошибки при билде
/matutils.cu:112: ошибка: нет декларации blockIdx в этой области видимости
/matutils.cu:112: ошибка: нет декларации blockDim в этой области видимости
/matutils.cu:112: ошибка: нет декларации threadIdx в этой области видимости
/matutils.cu:117: ошибка: нет декларации __syncthreads в этой области видимости
и т.д. по поводу самых основных объектов CUDA, я подключил два хедера,
#include "/usr/local/cuda/include/cuda.h"
#include "/usr/local/cuda/include/cuda_runtime.h"
их должно быть достаточно, остальные должны были подключится из них, однако ни Eclipse, ни CMake похоже не видят эти хедеры, или же мне еще подсказали на буржуйском форуме, что я всё еще не использую nvcc для компиляции, но я не знаю, где он настраивается, в Eclipse в свойствах проекта ничего нет, так как он предполагает использование стороннего Makefile и там должен быть nvcc, в CMakeList.txt тоже не вижу упоминания про nvcc, не знаю как его туда (или куда) впихнуть. Подскажите, кто на чём работает? Вообще у меня Ubuntu 9.04 и на другую ОС перейти не могу.

На мой взгляд использование

На мой взгляд использование common.mk в своих make-файлах достаточно удобно под Linux. Eclipse использую только для раскраски текста). Но то что GNU-ый отладчик не самый удобный - соглашусь! Хотя это хотя бы что-то.

меня и так тоже устроит,

меня и так тоже устроит, только не знаю, как работать с common.mk. Он используется для make? Я не имел опыта работы с make напрямую. Все привык делать с визуальными IDE. Может расскажете, как её использовать и настроить для CUDA? А с gdb тоже работаете в консоли? Ведь он не позволяет отлаживать код функций CUDA или позволяет? У куда ведь был свой отладчик.

вроде бы собрал, на основе

вроде бы собрал, на основе template проекта из SDK, всё получилось, однако мне кроме CUDA надо подключить и обычный код для загрузки PNG файлов, в частности хедер и статик-либрари, однако пока что сыпяться ошибки. Вот мой makefile

EXECUTABLE := ../../../src/projectc/build/ptest
CUFILES := projectc.cu
CU_DEPS := matutils.cu
CCFILES := image.cpp
include ../../common/common.mk
CXXFLAGS += -I/usr/include -I/usr/local/include -I. -I..
CXXFLAGS += -L/usr/lib -L/usr/local/lib -L/usr/local/cuda/lib/ -O3
LDFLAGS += -lpng

где project.cu - основной файл с main функцией, matutils.cu - файл с функциями CUDA, image.cpp - файл с обычными функциями по загрузке PNG изображения. Только пока то, что в image.cpp ошибки, так как он не видит хедеры stdio.h png.h и т.д., которые я вставил в projectc.cu , а именно:

#include "stdlib.h"
#include "stdio.h"
#include "string.h"
#include "math.h"
#include "time.h"
#include "png.h"
using namespace std
#include
#include
#include
#include

С чем это связано? Файлы с расширением .cu полностью отдаются на компиляцию в CUDA или там можно писать код обычного С++?

так вроде разобрался. КУДе

так вроде разобрался. КУДе кудово, если файл с .cu, то хедеры из этого файла не транслируются в файлы cpp и они их не видят, я прав? Теперь не могу отладить функцию матричного умножения (не работает, но ошибок не выдает), как и другие, написанные для видеокарты. Кто какими методами дебага пользуется?

Во-первых: при компелции в

Во-первых: при компелции в режиме эмуляции(для компиляции с использованием common.mk параметр emu=1) можно использовать например printf даже в __device__ функции. Можно использовать расширение гнушного отладчика cuda-gdb.

компилирую (make emu=1

компилирую (make emu=1 dbg=1), однако в функциях __global__ printf не работает, ничего не выдает, пример
__global__ void __sadd(float * a, float b, float * c, long N, long M)
{
long i = blockIdx.x * blockDim.x + threadIdx.x;
long j = blockIdx.y * blockDim.y + threadIdx.y;
long index = i + j * N;
printf("%d %d\n",i,j);
if( i < N && j < M)
c[index] = a[index] + b;
__syncthreads();
}
а отладчик cuda-gdb можно прикрепить к Eclipse? Или как там ставить брейкпоинт? Из сухого мануала, что я нашел, мало понятного.

Насчет отладчика не знаю, но

Насчет отладчика не знаю, но функция printf, насколько мне известно, в gpu коде и не должна работать. Все что в кернеле компилируется под GPU.

да, заработал. Однако всё

да, заработал. Однако всё равно не могу вычислить , почему не работает матричное умножение. Использую код с сайт NVidia, Моя задача - умножить матрицу изображения, представленную как 3 канала Х Высота*Ширина, каждый канал на R,G и B соответственно, вытянуты они в одну строку, на другую матрицу константных чисел, размером всего лишь 3Х3 для преобразования в YIQ формат. По правилу необходимо умножить так T * ImageRGB = ImageYIQ, где T[3X3] и ImageRGB[3XHeight*Width], чтобы размер ImageYIQ остался таким же, как ImageRGB. однако функция умножения таким образом не запускается (пишу в ней printf для проверки входа). А когда проверяю умножения на других размерах, например [3*HeightXWidth], то умножение работает, однако мне надо именно [3XHeight*Width]. Исходя из кода функции я понял, что умножение разбивается на блоки-подматрицы размером BLOCK_SIZE X BLOCKSIZE, чтоб использовать в каждом блоке shared memory. Я чувствую подвох в этом, поэтому выставил BLOCK_SIZE = 1 вместо 16, чтоб не получилось 3/BLOCK_SIZE = 0. Однако это не помогло, а разобраться в расчете индексов входных и выходной матриц и их подматриц, принятом в функции пока что не получается. У кого-нибудь есть идеи или другой пример матричного умножения на CUDA?

> однако функция умножения

> однако функция умножения таким образом не запускается
Какая ошибка выводится?
>Я чувствую подвох в этом, поэтому выставил BLOCK_SIZE = 1 вместо 16
Никто же не запрещает сделать 2 переменные BLOCK_SIZE_X = 3 BLOCK_SIZE_Y = 16(или другие значения).

  1. int nGridY = lny/BLOCK_SIDE_Y;
  2. if(lny%BLOCK_SIDE_Y !=0){
  3.   nGridY++;
  4. }

А в кернале поспаить проверку что бы лишние значения не считались как то так например:

  1.        
  2. __global__ void foo(int x_end, int y_end){
  3.   const unsigned int x =  blockIdx.x*blockDim.x + threadIdx.x;
  4.   const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
  5.   if((x_end > x) && (y_end > y) ){
  6.     ..............
  7.   }
  8. }

Где x_end и y_end размер ImageRGB

Примичания: 1) Конечно не стоит забывать что потоки не попавшие в if будут простивать 2) Блок потоков размером 3Х16 - плохой вариант. Если уж BLOCK_SIZE_Х = 3 - жёстко то BLOCK_SIZE_Y я бы сделал равным 85(3*85 = 255 потоков). Но это уже дело оптимизации

Пока что хочу разобраться без

Пока что хочу разобраться без оптимизации. Оказалось, что функция работает, однако только для маленьких изображений, т.е. проверял на нескольких, в итоге на изображении в 400Х400 и выше уже не работает, совсем не запускается, а на тех, что меньше, запускается и туда заходит. Как это связано с размером, не могу понять. Вот еще раз код функции

  1. __global__ void __prod(float *A, float *B, long wA, long wB, float *C)
  2. {
  3.         long bx = blockIdx.x;
  4.         long by = blockIdx.y;
  5.         long tx = threadIdx.x;
  6.         long ty = threadIdx.y;
  7.         long aBegin = wA * BLOCK_SIZE * by;
  8.         long aEnd = aBegin + wA - 1;
  9.         long aStep = BLOCK_SIZE;
  10.         long bBegin = BLOCK_SIZE * bx;
  11.         long bStep = BLOCK_SIZE * wB;
  12.         float Csub = 0;
  13.         for (long a = aBegin, b = bBegin; a <= aEnd; a += aStep, b += bStep)
  14.         {
  15.         __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
  16.         __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
  17.         As[ty][tx] = A[a + wA * ty + tx];
  18.         Bs[ty][tx] = B[b + wB * ty + tx];
  19.         __syncthreads();
  20.         for (long k = 0; k < BLOCK_SIZE; ++k) Csub += As[ty][k] * Bs[k][tx];
  21.         __syncthreads();
  22.         }
  23.         long c = wB * BLOCK_SIZE * by + BLOCK_SIZE * bx;
  24.         C[c + wB * ty + tx] = Csub;
  25. }
  26. void prod(float *&d_A, float *&d_B, float *&d_C, long wA, long hA, long wB)
  27. {
  28.         dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
  29.         dim3 dimGrid(wB / dimBlock.x, hA / dimBlock.y);
  30.         __prod<<<dimGrid, dimBlock>>>(d_A, d_B, wA, wB, d_C);
  31. }
  32.  
  33. вот так я делаю преобразование
  34. void rgb_yiq(float *&img_in, float *&img_out, long iH,long iW, bool dir = true )
  35. {
  36.     if (dir)
  37.     {
  38.         float RGBYIQH[3][3]={{0.299 , 0.587 ,0.114},{ 0.595716, -0.274453, -0.321263},{ 0.211456, -0.522591, 0.311135}};
  39.         float *RGBYIQ;
  40.         cudaMalloc((void**) &RGBYIQ,3*3*sizeof(float));
  41.         cudaMemcpy(RGBYIQ,RGBYIQH,3*3*sizeof(float),cudaMemcpyHostToDevice);
  42.         prod(RGBYIQ,img_in,img_out,3,3,iH*iW);
  43.         cudaFree(RGBYIQ);
  44.     }
  45.     else
  46.     {
  47.         float YIQRGBH[3][3] = {{1.0, 0.9563, 0.621},{ 1.0, -0.2721, -0.6474},{ 1.0, -1.107, 1.7046}};
  48.         float *YIQRGB;
  49.         cudaMalloc((void**) &YIQRGB,3*3*sizeof(float));
  50.         cudaMemcpy(YIQRGB,YIQRGBH,3*3*sizeof(float),cudaMemcpyHostToDevice);
  51.         prod(YIQRGB,img_in,img_out,3,3,iH*iW);
  52.         cudaFree(YIQRGB);
  53.     }
  54. }

Ставьте тег code, а то

Ставьте тег code, а то больше-меньше пропадают.

А чему равен BLOCK_SIZE?

BLOCK_SIZE равен 1,

BLOCK_SIZE равен 1, изначально был равен 16, но я изменил, чтобы предотвратить 3/BLOCK_SIZE = 0.

Одна из ваших проблем связана

Одна из ваших проблем связана с тем, что размерность грида равна iH*iW т.е. размеру картинки, для картинок больше чем 256x256 вы превышаете максимальную размерность в одном измерении (65536 или около того)

да, действительно, проблема в

да, действительно, проблема в этом. Пока что представил каждый канал в виде отдельной матрицы [Height X Width] и всё заработало. Вот теперь не знаю, как организовать вычисление среднего арифметического с использование CUDA, там аккумулирующая сумма - одна и та же ячейка в памяти, каждый поток одновременно или через гонки будет пытаться получить к ней доступ, и что они там насуммируют - неизвестно, результат должен быть ошибочным. Задача мне кажется исключительно последовательной. Как её решать в GPU?

Посмотрите как сделаны

Посмотрите как сделаны гистограммы (в примерах CUDA там даже дока отдельная на это место есть), это если хочется совсем выпендриться.

Но вообще, все несложно. Как-то так:
- делаем не очень много блоков (скажем 128)
- в каждом делаем не очень много threads
- в хвосте thread обычной редукцией складываем среднее для блока и записываем в глобальную память (уникальное место для каждого блока)
На хосте усредняем оставшиеся 128 чисел.

Ну да. Посмотрите пример из

Ну да.

Посмотрите пример из CUDA Toolkit, так и называется reduction, там в doc/ лежит презентация Харриса "как мы ускорили редукцию в 30 раз"

я открыл код reduction из

я открыл код reduction из toolkit, однако такой код мне не очень понятен, я написал свой, видимо это тоже каскадное суммирование, он работает, однако мне кажется, что выдает неверные значения, так как есть реализация обычная для тех же данных (того же изображения), и эти ср. арифметические отличаются от тех. Я разбил задачу на две стадии, сделав предварительный массив сумм.

  1. __global__ void __meanStage1(float *A, float *Sum, long N, long M)
  2. {
  3.     long iS = blockIdx.y * gridDim.x + blockIdx.x;
  4.     for (long j= blockIdx.y * BLOCK_SIZE; j< (blockIdx.y+1) * BLOCK_SIZE ; j++)
  5.                 for (long i= blockIdx.x * BLOCK_SIZE; i< (blockIdx.x+1) * BLOCK_SIZE ; i++)
  6.             if( i < N && j < M)
  7.                 Sum[iS] = Sum[iS] + A[j * M + i];
  8.     __syncthreads();
  9. }
  10. __global__ void __meanStage2(float *a,float *mean, float *sum, long dX, long dY, long N, long M)
  11. {
  12.         *mean = 0;
  13.     for (long i=0;i<dX * dY;i++)
  14.         *mean = *mean + sum[i];
  15.     if (M*N!=0)
  16.     *mean = (*mean)/(N*M);
  17.     else
  18.     *mean = 0;
  19.     __syncthreads();
  20. }
  21. void mean(float *a,float &Means, long N, long M)
  22. {
  23.         float *sum,*mean;
  24.         long iS,jS;
  25.         iS = N / BLOCK_SIZE;
  26.         jS = M / BLOCK_SIZE;
  27.         cudaMalloc((void**) &mean , sizeof(float));
  28.         cudaMalloc((void**) &sum , iS * jS * sizeof(float));
  29.         cudaMemset(sum,0, iS * jS * sizeof(float));
  30.     dim3 dimBlock1(1, 1);
  31.     dim3 dimGrid1(N/BLOCK_SIZE, M/BLOCK_SIZE);
  32.     __meanStage1<<<dimGrid1, dimBlock1>>>(a, sum, N, M);
  33.     dim3 dimBlock2(1, 1);
  34.     dim3 dimGrid2(1, 1);
  35.     __meanStage2<<<dimGrid2, dimBlock2>>>(a, mean, sum, iS, jS, N, M);
  36.     cudaFree(sum);
  37.     cudaMemcpy(&Means, mean, sizeof(float),cudaMemcpyDeviceToHost);
  38.     cudaFree(mean);
  39. }

всё проверил, похоже этот код

всё проверил, похоже этот код среднего арифметического правильный, но не совсем. Работает он корректно только в эмуляции, когда же пытаюсь исполнить на самой видеокарте, он выдаёт неадекватные числа, или вообще создаёт артефакты на экране. Тут как будто две проблемы, рассинхронизация потоков и что-то с памятью видеокарты. Но вот конкретно не могу понять, где же это возникает.

Боюсь огорчить, но обычно

Боюсь огорчить, но обычно людям не свойственно читать чужой код с экрана и пытаться выполнить его "в уме".

я прошу совета, а не

я прошу совета, а не компилировать и исполнять мой код у себя. Возможно у кого-то были подобные ситуации и ошибки. Я нашел презентацию Харриса по редукции, взял оттуда метод под номером 4 (Reduction #4: First Add During Load), однако он вызывает segmentation fault. Там два параметра - входная и выходная матрица. Пробовал туда передавать одну и ту же и две разных матрицы, всё равно вываливается.

Если работает в эмуляторе

Если работает в эмуляторе (особенно если эмулятор исполняется на одном ядре/процессоре), но не работает на устройстве - скорее всего у вас проблемы с синхронизацией, разные треды пишут в одну shared mem или разные треды/блоки - в одно место глобальной памяти.

За минуту я проблемы не увидел, но симптомы - такие.

Я бы попробовал __meanStage1

Я бы попробовал __meanStage1 и __meanStage2 выгрузить sum в режиме эмуляции и без него и сравнить их, тем самым проверить правельно ли работает __meanStage1. Так хоть можно будет чуть локализовать ошибку

Когда изображение не

Когда изображение не превышает (примерно) 400Х400, обе работают нормально в режиме эмуляции и без. А когда большие изображения, то неправильно работает __meanStage1, при эмуляции даёт неверные данные, при обычном режиме бывает неверные данные, а бывает экран покрывается "артефактами". Я пробовал разные "высота х ширина", но закономерность вижу пока только в увеличении размеров.

:-D

:-D

есть еще одно, у меня иногда

есть еще одно, у меня иногда проявляются артефакты на экране в виде "снега", видимо портится память в видеокарте. Как вообще она там выделяется при выполнении программы? возможно ли , что пересечение памяти под другие приложения (в том числе отрисовки оболочки) и памяти программы?

Я настоятельно рекомендую

Я настоятельно рекомендую прочитать про аппаратные особенности видокарты. Хотя бы внимательно прочитать слайды с МГУ-шной презентации.

В Eclipse у меня не

В Eclipse у меня не получилось интегрировать cuda-gdb.
Для использования cuda-gdb надо иметь представление о работе gdb. Вообще про него(gdb) вот тут очень хорошо написано http://www.linux.org.ru/books/GNU/gdb/gdb_toc.html. А в Nvidia_CUDA_GDB_User_Manual_2.3 есть пошаговый пример(но сначала GNU/gdb) - думаю его мысла нет тут переписывать :-)

Хотя проблема вроде бы не требует исползования cuda-gdb.

Да еще __syncthreads(); в

Да еще __syncthreads(); в нутри ветвления НЕ НАДО ставить только если ты уверен что все потоки в блоке пройдут по одномой ветке. Иначе у тебя пожет программа уйти в себя )). У барьера ожидаться будут потоки, которые в эту ветку никогда не зайдут!