IDE для CUDA

Tagged:  

Хочу выбрать 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, правда не знаю, насколько он сейчас самостоятелен. Критериями является быстродействие и удобство работы с матрицами, наличие простых и сложных операций и функций.

Comments

вроде бы собрал, на основе 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 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? Или как там ставить брейкпоинт? Из сухого мануала, что я нашел, мало понятного.

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

В 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.

Насчет отладчика не знаю, но функция 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?

Пока что хочу разобраться без оптимизации. Оказалось, что функция работает, однако только для маленьких изображений, т.е. проверял на нескольких, в итоге на изображении в 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. }

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

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

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

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

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

Под "обычной редукцией" вы понимаете каскадную схему суммирования?

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

я открыл код 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. }

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

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

:-D

:-D

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

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

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

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

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

Надо бы пренести то что не касается IDE в другую тему.

Ну да.

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

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

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

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

> однако функция умножения таким образом не запускается
Какая ошибка выводится?
>Я чувствую подвох в этом, поэтому выставил 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 потоков). Но это уже дело оптимизации

3x3 можно и руками расписать

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

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

да, пока 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 и на другую ОС перейти не могу.

Copyright © 2008-2011 Alex Tutubalin