Non functional UI

Signed-off-by: Gaurav Kukreja <gmkukreja@gmail.com>

miklos/project_integration_gui/camera.cpp

+#include "camera.h"
+
+#include <opencv2/imgproc/imgproc.hpp>
+
+Camera camera;
+
+bool Camera::init(int device)
+{
+    capture_.open(device);
+    if (!capture_.isOpened())
+        return false;
+
+    int camW = 640;
+    int camH = 480;
+  	capture_.set(CV_CAP_PROP_FRAME_WIDTH, camW);
+  	capture_.set(CV_CAP_PROP_FRAME_HEIGHT, camH);
+
+    capture();
+    width_  = frame_.cols;
+    height_ = frame_.rows;
+
+    return true;
+}
+
+void Camera::run()
+{
+    QMutexLocker locker(&mutex);
+    while (true) {  // TODO: never terminates
+        frameCopied.wait(&mutex);
+        capture();
+        emit newFrame();
+    }
+}
+
+void Camera::capture()
+{
+    capture_ >> frame_;
+    cvtColor(frame_, frame_, CV_BGR2GRAY);
+
+    frame_.convertTo(frame_, CV_32F);
+    frame_ /= 255.f;
+}

miklos/project_integration_gui/camera.h

+#ifndef CAMERA_H
+#define CAMERA_H
+
+#include <QMutex>
+#include <QThread>
+#include <QWaitCondition>
+
+#include <opencv2/highgui/highgui.hpp>
+
+class Camera : public QThread {
+    Q_OBJECT
+
+public:
+    bool init(int device);
+    
+    size_t width() const { return width_; }
+    size_t height() const { return height_; }
+
+    float *data() const {
+        return reinterpret_cast<float *>(frame_.data);
+    }
+
+    QMutex mutex;
+    QWaitCondition frameCopied;
+
+signals:
+    void newFrame();
+
+protected:
+    void run();
+
+private:
+    cv::VideoCapture capture_;
+    cv::Mat frame_;
+    size_t width_, height_;
+    
+    void capture();
+};
+
+extern Camera camera;
+
+#endif // CAMERA_H

miklos/project_integration_gui/glwidget.cpp

+#include "glwidget.h"
+#include "camera.h"
+#include "kernel.h"
+
+#include <QGLFunctions>
+#include <iostream>
+
+#include <cuda_runtime_api.h>
+#include <cuda_gl_interop.h>
+
+static GLuint pixelsVBO;
+static struct cudaGraphicsResource* pixelsVBO_CUDA;
+
+GlWidget::GlWidget(QWidget *parent)
+    : QGLWidget(QGLFormat(), parent), gl(context())
+{
+}
+
+GlWidget::~GlWidget()
+{
+	cudaFree(d_in);
+}
+
+QSize GlWidget::sizeHint() const
+{
+    return QSize(camera.width(), camera.height());
+}
+
+void GlWidget::initializeGL()
+{
+    // Explicitly set device 0
+    cudaGLSetGLDevice(0);
+
+    // Create buffer object and register it with CUDA
+    gl.glGenBuffers(1, &pixelsVBO);
+    gl.glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pixelsVBO);
+    size_t size = camera.width() * camera.height() * 4 * sizeof(unsigned char);
+    gl.glBufferData(GL_PIXEL_UNPACK_BUFFER, size, 0, GL_DYNAMIC_DRAW);
+    cudaGraphicsGLRegisterBuffer(&pixelsVBO_CUDA, pixelsVBO, cudaGraphicsMapFlagsWriteDiscard);
+
+    size_t inBytes = camera.width() * camera.height() * sizeof(float);
+    cudaMalloc((void **)&d_in, inBytes);
+}
+
+void GlWidget::paintGL()
+{
+    // Map buffer object for writing from CUDA
+    void *d_out;
+    cudaGraphicsMapResources(1, &pixelsVBO_CUDA, 0);
+    size_t size; 
+    cudaGraphicsResourceGetMappedPointer(&d_out, &size,  pixelsVBO_CUDA);
+
+    size_t inBytes = camera.width() * camera.height() * sizeof(float);
+    {
+        QMutexLocker locker(&camera.mutex);
+        cudaMemcpy(d_in, camera.data(), inBytes, cudaMemcpyHostToDevice);
+        cudaDeviceSynchronize();
+        camera.frameCopied.wakeAll();
+    }
+
+    // Execute kernel
+    executeKernel(d_in, d_out, camera.width(), camera.height());
+
+    // Unmap buffer object
+    cudaGraphicsUnmapResources(1, &pixelsVBO_CUDA, 0);
+
+    // Render from buffer object
+    glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
+    glDrawPixels(camera.width(), camera.height(), GL_RGBA, GL_UNSIGNED_BYTE, 0);
+}

miklos/project_integration_gui/glwidget.h

+#ifndef GLWIDGET_H
+#define GLWIDGET_H
+
+#include <QGLWidget>
+#include <QGLFunctions>
+
+class GlWidget : public QGLWidget
+{
+    Q_OBJECT
+
+public:
+    explicit GlWidget(QWidget *parent = 0);
+    ~GlWidget();
+    QSize sizeHint() const;
+    float *d_in;
+
+protected:
+    void initializeGL();
+    void paintGL();
+
+private:
+    QGLFunctions gl;
+};
+
+#endif // GLWIDGET_H

miklos/project_integration_gui/hello-opengl.pro

+QT += core gui opengl
+CONFIG += debug
+
+TARGET = hello-opengl
+
+TEMPLATE = app
+
+SOURCES += main.cpp glwidget.cpp camera.cpp \
+    mainwindow.cpp
+HEADERS += glwidget.h kernel.h camera.h \
+    mainwindow.h
+
+CUDA_SOURCES += kernel.cu
+
+LIBS += -lGLEW -lcudart -lcuda -lopencv_core -lopencv_highgui -lopencv_imgproc
+
+QMAKE_LFLAGS += -Wl,-rpath,/usr/lib/nvidia-319-updates -Wl,-rpath,/usr/lib/nvidia-319-updates/tls
+# Path to cuda toolkit install
+CUDA_DIR      = /usr/local/cuda-5.5
+
+# Path to header and libs files
+INCLUDEPATH  += $$CUDA_DIR/include
+QMAKE_LIBDIR += $$CUDA_DIR/lib64
+QMAKE_LIBDIR += /usr/lib/nvidia-319-updates /usr/lib/nvidia-319-updates/tls
+# GPU architecture
+CUDA_ARCH     = sm_10
+NVCCFLAGS     = --compiler-options -use_fast_math --ptxas-options=-v
+
+CUDA_INC = $$join(INCLUDEPATH,' -I','-I',' ')
+ 
+cuda.commands = $$CUDA_DIR/bin/nvcc -m64 -O3 -arch=$$CUDA_ARCH -c $$NVCCFLAGS $$CUDA_INC $$LIBS  ${QMAKE_FILE_NAME} -o ${QMAKE_FILE_OUT}
+cuda.dependency_type = TYPE_C
+cuda.depend_command = $$CUDA_DIR/bin/nvcc -O3 -M $$CUDA_INC $$NVCCFLAGS ${QMAKE_FILE_NAME}
+cuda.input = CUDA_SOURCES
+cuda.output = ${OBJECTS_DIR}${QMAKE_FILE_BASE}_cuda.o
+
+# Tell Qt that we want add more stuff to the Makefile
+QMAKE_EXTRA_COMPILERS += cuda
+
+FORMS += \
+    mainwindow.ui

miklos/project_integration_gui/kernel.cu

+#include "kernel.h"
+#include "timer.h"
+
+#include <algorithm>
+#include <stdio.h>
+
+texture<float,2,cudaReadModeElementType> texRef_Xi;
+texture<float,2,cudaReadModeElementType> texRef_Xj;
+
+template<typename T>
+__device__ __host__ T min(T a, T b)
+{
+    return (a < b) ? a : b;
+}
+
+template<typename T>
+__device__ __host__ T max(T a, T b)
+{
+    return (a > b) ? a : b;
+}
+
+template<typename T>
+__device__ __host__ T clamp(T m, T x, T M)
+{
+    return max(m, min(x, M));
+}
+
+
+__global__ void calculate_F(float *U, float *F, int w, int h, float c1, float c2, float lambda)
+{
+    int x = threadIdx.x + blockDim.x * blockIdx.x;
+    int y = threadIdx.y + blockDim.y * blockIdx.y;
+    if (x < w && y < h) {
+        size_t i = x + (size_t)w*y;
+        float temp_ui = U[i];
+        F[i] = lambda * ((c1 - temp_ui)*(c1 - temp_ui) - (c2 - temp_ui)*(c2 - temp_ui));
+    }
+}
+
+__device__ float diff_i(float *M, int w, int h, int x, int y)
+{
+    size_t i = x + (size_t)w*y;
+    return (x+1 < w) ? (M[i + 1] - M[i]) : 0.f;
+}
+
+__device__ float diff_j(float *M, int w, int h, int x, int y)
+{
+    size_t i = x + (size_t)w*y;
+    return (y+1 < h) ? (M[i + w] - M[i]) : 0.f;
+}
+
+__global__ void update_Xij(float *Xi, float *Xj, float *T, float *U, int w, int h, float sigma)
+{
+    int x = threadIdx.x + blockDim.x * blockIdx.x;
+    int y = threadIdx.y + blockDim.y * blockIdx.y;
+    if (x < w && y < h) {
+		size_t i = x + (size_t) w * y;
+		float xi = Xi[i] - sigma * (2 * diff_i(U, w, h, x, y) - diff_i(T, w, h, x, y));
+		float xj = Xj[i] - sigma * (2 * diff_j(U, w, h, x, y) - diff_j(T, w, h, x, y));
+		float dn = max(1.f, sqrtf(xi * xi + xj * xj));
+		Xi[i] = xi / dn;
+		Xj[i] = xj / dn;
+    }
+}
+
+__device__ float divergence(float *X, float *Y, int w, int h, int x, int y)
+{
+    float dx_x = tex2D(texRef_Xi, x + 0.5f , y + 0.5f) - tex2D(texRef_Xi, x - 0.5f , y + 0.5f);
+    float dy_y = tex2D(texRef_Xj, x + 0.5f , y + 0.5f) - tex2D(texRef_Xj, x + 0.5f , y - 0.5f);
+
+    return dx_x + dy_y;
+}
+
+__global__ void update_U(float *T, float *Xi, float *Xj, float *F, float *U, int w, int h, float tau)
+{
+    int x = threadIdx.x + blockDim.x * blockIdx.x;
+    int y = threadIdx.y + blockDim.y * blockIdx.y;
+    if (x < w && y < h) {
+        size_t i = x + (size_t)w*y;
+        U[i] = clamp(0.f, T[i] - tau * (divergence(Xi, Xj, w, h, x, y) + F[i]), 1.f);
+    }
+}
+
+__global__ void update_Output(uchar4* output, float *U, int w, int h) {
+
+    int x = threadIdx.x + blockDim.x * blockIdx.x;
+    int y = threadIdx.y + blockDim.y * blockIdx.y;
+    size_t i = x + (size_t)w*y;
+    size_t idx = x + (size_t) w*(h-1 - y);
+    unsigned char temp_res = roundf((U[i] * 255.f));
+    output[idx].x = temp_res;
+    output[idx].y = temp_res;
+    output[idx].z = temp_res;
+    output[idx].w = 255;
+
+}
+
+inline int div_ceil(int n, int b) { return (n + b - 1) / b; }
+
+inline dim3 make_grid(dim3 whole, dim3 block)
+{
+    return dim3(div_ceil(whole.x, block.x),
+                div_ceil(whole.y, block.y),
+                div_ceil(whole.z, block.z));
+}
+
+static float *d_T, *d_F, *d_Xi, *d_Xj;
+
+
+void allocate_device_memory(float *d_in, size_t w, size_t h)
+{
+	size_t imageBytes = w*h*sizeof(float);
+	cudaMalloc(&d_T, imageBytes);
+	cudaMalloc(&d_F, imageBytes);
+	cudaMalloc(&d_Xi, imageBytes);
+	cudaMalloc(&d_Xj, imageBytes);
+
+	// Define texture attributes
+    texRef_Xi.addressMode[0] = cudaAddressModeClamp;             // clamp x to border
+    texRef_Xi.addressMode[1] = cudaAddressModeClamp;            // clamp y to border
+    texRef_Xi.filterMode = cudaFilterModeLinear;            // linear interpolation
+    texRef_Xi.normalized = false;
+    cudaChannelFormatDesc desc_Xi = cudaCreateChannelDesc<float>();
+    cudaBindTexture2D(NULL, &texRef_Xi, d_Xi, &desc_Xi, w, h, w*sizeof(d_Xi[0]));
+
+	// Define texture attributes
+    texRef_Xj.addressMode[0] = cudaAddressModeClamp;             // clamp x to border
+    texRef_Xj.addressMode[1] = cudaAddressModeClamp;            // clamp y to border
+    texRef_Xj.filterMode = cudaFilterModeLinear;            // linear interpolation
+    texRef_Xj.normalized = false;
+    cudaChannelFormatDesc desc_Xj = cudaCreateChannelDesc<float>();
+    cudaBindTexture2D(NULL, &texRef_Xj, d_Xj, &desc_Xj, w, h, w*sizeof(d_Xj[0]));
+
+}
+
+void executeKernel(float *d_U, void *d_out, size_t w, size_t h)
+{
+    // float *d_U = reinterpret_cast<float *>(d_in);
+    uchar4 *pixel = reinterpret_cast<uchar4 *>(d_out);
+
+    static Timer timer;
+    timer.end();
+    printf("time: %.2fms (%.2f FPS)\n", timer.get() * 1E3F, 1.F / timer.get());
+    timer.start();
+
+    dim3 dimBlock(32, 16);
+    dim3 dimGrid = make_grid(dim3(w, h, 1), dimBlock);
+
+    // set parameters manually here
+    float lambda = 1.0;
+    float sigma = 0.4;
+    float tau = 0.4;
+    int N = 160;
+    float c1 = 1.0;
+    float c2 = 0.00;
+
+    size_t imageBytes = w*h*sizeof(float);
+    cudaMemcpy(d_T, d_U, imageBytes, cudaMemcpyDeviceToDevice);
+    cudaMemset(d_Xi, 0, imageBytes);
+    cudaMemset(d_Xj, 0, imageBytes);
+
+    calculate_F<<< dimGrid, dimBlock >>>(d_U, d_F, w, h, c1, c2, lambda);
+
+    for (int n = 0; n < N; n++) {
+        update_Xij<<< dimGrid, dimBlock >>>(d_Xi, d_Xj, d_T, d_U, w, h, sigma);
+        std::swap(d_U, d_T);
+        update_U<<< dimGrid, dimBlock >>>(d_T, d_Xi, d_Xj, d_F, d_U, w, h, tau);
+    }
+    update_Output<<< dimGrid, dimBlock >>>(pixel, d_U, w, h);
+}

miklos/project_integration_gui/kernel.h

+#ifndef KERNEL_H
+#define KERNEL_H
+
+#include <stdlib.h>
+
+extern "C" void allocate_device_memory(float *d_in, size_t width, size_t height);
+extern "C" void executeKernel(float *d_in, void *d_out, size_t width, size_t height);
+
+#endif // KERNEL_H

miklos/project_integration_gui/main.cpp

+//#include "camera.h"
+#include "kernel.h"
+#include "mainwindow.h"
+//#include "glwidget.h"
+
+#include <iostream>
+#include <QApplication>
+
+int main(int argc, char *argv[])
+{
+    QApplication app(argc, argv);
+    MainWindow w;
+
+    w.init();
+    w.show();
+
+    return app.exec();
+}

miklos/project_integration_gui/mainwindow.cpp

+#include "mainwindow.h"
+#include "camera.h"
+#include "ui_mainwindow.h"
+#include "kernel.h"
+
+#include <iostream>
+
+MainWindow::MainWindow(QWidget *parent) :
+    QMainWindow(parent),
+    ui(new Ui::MainWindow)
+{
+    ui->setupUi(this);
+}
+
+MainWindow::~MainWindow()
+{
+    delete ui;
+}
+
+int MainWindow::init()
+{
+    // Init camera
+    if(!camera.init(0)) {
+        std::cerr << "ERROR: Could not open camera" << std::endl;
+        return 1;
+    }
+
+    QObject::connect(&camera, SIGNAL(newFrame()), ui->widget, SLOT(updateGL()));
+    allocate_device_memory(ui->widget->d_in, camera.width(), camera.height());
+
+    camera.start();
+
+    return 0;
+}

miklos/project_integration_gui/mainwindow.h

+#ifndef MAINWINDOW_H
+#define MAINWINDOW_H
+
+#include <QMainWindow>
+
+namespace Ui {
+class MainWindow;
+}
+
+class MainWindow : public QMainWindow
+{
+    Q_OBJECT
+
+public:
+    explicit MainWindow(QWidget *parent = 0);
+    ~MainWindow();
+
+    int init();
+
+private:
+    Ui::MainWindow *ui;
+};
+
+#endif // MAINWINDOW_H

miklos/project_integration_gui/mainwindow.ui

+<?xml version="1.0" encoding="UTF-8"?>
+<ui version="4.0">
+ <class>MainWindow</class>
+ <widget class="QMainWindow" name="MainWindow">
+  <property name="geometry">
+   <rect>
+    <x>0</x>
+    <y>0</y>
+    <width>865</width>
+    <height>602</height>
+   </rect>
+  </property>
+  <property name="sizePolicy">
+   <sizepolicy hsizetype="Fixed" vsizetype="Fixed">
+    <horstretch>0</horstretch>
+    <verstretch>0</verstretch>
+   </sizepolicy>
+  </property>
+  <property name="minimumSize">
+   <size>
+    <width>865</width>
+    <height>602</height>
+   </size>
+  </property>
+  <property name="maximumSize">
+   <size>
+    <width>865</width>
+    <height>602</height>
+   </size>
+  </property>
+  <property name="windowTitle">
+   <string>MainWindow</string>
+  </property>
+  <widget class="QWidget" name="centralwidget">
+   <widget class="GlWidget" name="widget" native="true">
+    <property name="geometry">
+     <rect>
+      <x>30</x>
+      <y>50</y>
+      <width>640</width>
+      <height>480</height>
+     </rect>
+    </property>
+    <property name="sizePolicy">
+     <sizepolicy hsizetype="Preferred" vsizetype="Preferred">
+      <horstretch>0</horstretch>
+      <verstretch>0</verstretch>
+     </sizepolicy>
+    </property>
+    <property name="baseSize">
+     <size>
+      <width>640</width>
+      <height>480</height>
+     </size>
+    </property>
+   </widget>
+   <widget class="QWidget" name="layoutWidget">
+    <property name="geometry">
+     <rect>
+      <x>700</x>
+      <y>50</y>
+      <width>131</width>
+      <height>481</height>
+     </rect>
+    </property>
+    <layout class="QVBoxLayout" name="verticalLayout">
+     <item>
+      <spacer name="verticalSpacer">
+       <property name="orientation">
+        <enum>Qt::Vertical</enum>
+       </property>
+       <property name="sizeHint" stdset="0">
+        <size>
+         <width>20</width>
+         <height>40</height>
+        </size>
+       </property>
+      </spacer>
+     </item>
+     <item>
+      <widget class="QLabel" name="label_4">
+       <property name="text">
+        <string>FPS</string>
+       </property>
+       <property name="alignment">
+        <set>Qt::AlignCenter</set>
+       </property>
+      </widget>
+     </item>
+     <item>
+      <widget class="QLCDNumber" name="lcdNumber"/>
+     </item>
+     <item>
+      <widget class="QLabel" name="label_3">
+       <property name="text">
+        <string>Alpha</string>
+       </property>
+       <property name="alignment">
+        <set>Qt::AlignCenter</set>
+       </property>
+      </widget>
+     </item>
+     <item>
+      <widget class="QDoubleSpinBox" name="doubleSpinBox_3"/>
+     </item>
+     <item>
+      <widget class="QLabel" name="label_2">
+       <property name="text">
+        <string>Beta</string>
+       </property>
+       <property name="alignment">
+        <set>Qt::AlignCenter</set>
+       </property>
+      </widget>
+     </item>
+     <item>
+      <widget class="QDoubleSpinBox" name="doubleSpinBox_2"/>
+     </item>
+     <item>
+      <widget class="QLabel" name="label">
+       <property name="text">
+        <string>Gama</string>
+       </property>
+       <property name="alignment">
+        <set>Qt::AlignCenter</set>
+       </property>
+      </widget>
+     </item>
+     <item>
+      <widget class="QDoubleSpinBox" name="doubleSpinBox"/>
+     </item>
+     <item>
+      <widget class="QPushButton" name="pushButton_2">
+       <property name="text">
+        <string>&amp;Refresh</string>
+       </property>
+      </widget>
+     </item>
+     <item>
+      <widget class="QPushButton" name="pushButton">
+       <property name="text">
+        <string>&amp;Quit</string>
+       </property>
+      </widget>
+     </item>
+    </layout>
+   </widget>
+  </widget>
+  <widget class="QMenuBar" name="menubar">
+   <property name="geometry">
+    <rect>
+     <x>0</x>
+     <y>0</y>
+     <width>865</width>
+     <height>23</height>
+    </rect>
+   </property>
+  </widget>
+  <widget class="QStatusBar" name="statusbar"/>
+ </widget>
+ <customwidgets>
+  <customwidget>
+   <class>GlWidget</class>
+   <extends>QWidget</extends>
+   <header>glwidget.h</header>
+   <container>1</container>
+   <slots>
+    <signal>update_param()</signal>
+   </slots>
+  </customwidget>
+ </customwidgets>
+ <resources/>
+ <connections>
+  <connection>
+   <sender>pushButton</sender>
+   <signal>clicked()</signal>
+   <receiver>MainWindow</receiver>
+   <slot>close()</slot>
+   <hints>
+    <hint type="sourcelabel">
+     <x>742</x>
+     <y>543</y>
+    </hint>
+    <hint type="destinationlabel">
+     <x>713</x>
+     <y>569</y>
+    </hint>
+   </hints>
+  </connection>
+ </connections>
+</ui>

miklos/project_integration_gui/timer.h

+#ifndef TIMER_H
+#define TIMER_H
+
+#include <time.h>
+
+// measuring time
+class Timer
+{
+    public:
+	Timer() : running(false), sec(0)
+	{
+	}
+	void start()
+	{
+        clock_gettime(CLOCK_MONOTONIC, &tStart);
+		running = true;
+	}
+	void end()
+	{
+		if (!running) { sec = 0; return; }
+        struct timespec tEnd;
+        clock_gettime(CLOCK_MONOTONIC, &tEnd);
+        sec = (tEnd.tv_sec - tStart.tv_sec) + (tEnd.tv_nsec - tStart.tv_nsec) * 1E-9F;
+		running = false;
+	}
+	float get()
+	{
+		if (running) end();
+		return sec;
+	}
+    private:
+    struct timespec tStart;
+	bool running;
+	float sec;
+};
+
+#endif // TIMER_H