ay

Almond.pro

@@ -64,12 +64,12 @@ else:unix:QMAKE_LFLAGS+= -fopenmp
 LIBS += -fopenmp
 unix:LIBS += -lm -latomic
 
-win32:CONFIG(release, debug|release): LIBS += -L$$PWD/../libs/ffmpeg-4.1.1-win32-dev/lib/ -lavcodec
-else:win32:CONFIG(debug, debug|release): LIBS += -L$$PWD/../libs/ffmpeg-4.1.1-win32-dev/lib/ -lavcodec
-else:unix: LIBS += -L$$PWD/../libs/ffmpeg-4.1.1-win32-dev/lib/ -lavcodec
+#win32:CONFIG(release, debug|release): LIBS += -L$$PWD/../libs/ffmpeg-4.1.1-win32-dev/lib/ -lavcodec
+#else:win32:CONFIG(debug, debug|release): LIBS += -L$$PWD/../libs/ffmpeg-4.1.1-win32-dev/lib/ -lavcodec
+#else:unix: LIBS += -L$$PWD/../libs/ffmpeg-4.1.1-win32-dev/lib/ -lavcodec
 
-INCLUDEPATH += $$PWD/../libs/ffmpeg-4.1.1-win32-dev/include
-DEPENDPATH += $$PWD/../libs/ffmpeg-4.1.1-win32-dev/include
+#INCLUDEPATH += $$PWD/../libs/ffmpeg-4.1.1-win32-dev/include
+#DEPENDPATH += $$PWD/../libs/ffmpeg-4.1.1-win32-dev/include
 
 
 #win32:CONFIG(release, debug|release): LIBS += -L$$PWD/'../../../../../Program Files (x86)/AMD APP SDK/3.0/lib/x86/' -lOpenCL
@@ -83,28 +83,28 @@ win32:CONFIG(release, debug|release): LIBS += -L$$PWD/../libs/ffmpeg-4.1.1-win32
 else:win32:CONFIG(debug, debug|release): LIBS += -L$$PWD/../libs/ffmpeg-4.1.1-win32-dev/lib/ -lavformat
 else:unix: LIBS += -L$$PWD/../libs/ffmpeg-4.1.1-win32-dev/lib/ -lavformat
 
-INCLUDEPATH += $$PWD/../libs/ffmpeg-4.1.1-win32-dev/include
-DEPENDPATH += $$PWD/../libs/ffmpeg-4.1.1-win32-dev/include
+#INCLUDEPATH += $$PWD/../libs/ffmpeg-4.1.1-win32-dev/include
+#DEPENDPATH += $$PWD/../libs/ffmpeg-4.1.1-win32-dev/include
 
-unix|win32: LIBS += -L$$PWD/../libs/ffmpeg-4.1.1-win32-dev/lib/ -lavdevice
+#unix|win32: LIBS += -L$$PWD/../libs/ffmpeg-4.1.1-win32-dev/lib/ -lavdevice
 
-INCLUDEPATH += $$PWD/../libs/ffmpeg-4.1.1-win32-dev/include
-DEPENDPATH += $$PWD/../libs/ffmpeg-4.1.1-win32-dev/include
+#INCLUDEPATH += $$PWD/../libs/ffmpeg-4.1.1-win32-dev/include
+#DEPENDPATH += $$PWD/../libs/ffmpeg-4.1.1-win32-dev/include
 
-unix|win32: LIBS += -L$$PWD/../libs/ffmpeg-4.1.1-win32-dev/lib/ -lavfilter
+#unix|win32: LIBS += -L$$PWD/../libs/ffmpeg-4.1.1-win32-dev/lib/ -lavfilter
 
-INCLUDEPATH += $$PWD/../libs/ffmpeg-4.1.1-win32-dev/include
-DEPENDPATH += $$PWD/../libs/ffmpeg-4.1.1-win32-dev/include
+#INCLUDEPATH += $$PWD/../libs/ffmpeg-4.1.1-win32-dev/include
+#DEPENDPATH += $$PWD/../libs/ffmpeg-4.1.1-win32-dev/include
 
-unix|win32: LIBS += -L$$PWD/../libs/ffmpeg-4.1.1-win32-dev/lib/ -lavutil
+#unix|win32: LIBS += -L$$PWD/../libs/ffmpeg-4.1.1-win32-dev/lib/ -lavutil
 
-INCLUDEPATH += $$PWD/../libs/ffmpeg-4.1.1-win32-dev/include
-DEPENDPATH += $$PWD/../libs/ffmpeg-4.1.1-win32-dev/include
+#INCLUDEPATH += $$PWD/../libs/ffmpeg-4.1.1-win32-dev/include
+#DEPENDPATH += $$PWD/../libs/ffmpeg-4.1.1-win32-dev/include
 
-unix|win32: LIBS += -L$$PWD/../libs/ffmpeg-4.1.1-win32-dev/lib/ -lswscale
+#unix|win32: LIBS += -L$$PWD/../libs/ffmpeg-4.1.1-win32-dev/lib/ -lswscale
 
-INCLUDEPATH += $$PWD/../libs/ffmpeg-4.1.1-win32-dev/include
-DEPENDPATH += $$PWD/../libs/ffmpeg-4.1.1-win32-dev/include
+#INCLUDEPATH += $$PWD/../libs/ffmpeg-4.1.1-win32-dev/include
+#DEPENDPATH += $$PWD/../libs/ffmpeg-4.1.1-win32-dev/include
 
 RESOURCES += \
     Almond.qrc

MandelWidget.cpp

@@ -107,7 +107,7 @@ void MandelView::adaptViewport(const MandelViewport& vp)
 MandelWidget::MandelWidget(mnd::MandelContext& ctxt, QWidget* parent) :
     QGLWidget{ QGLFormat(QGL::SampleBuffers), parent },
     mndContext{ ctxt },
-    mv{ ctxt.getCpuGenerator128() }
+    mv{ *ctxt.getDevices()[0].getGenerator128() }
 {
     this->setContentsMargins(0, 0, 0, 0);
     this->setSizePolicy(QSizePolicy::Expanding,

VideoStream.cpp

@@ -1,3 +1,5 @@
+#ifdef FFMPEG_ENABLED
+
 #include "VideoStream.h"
 
 #include <iostream>
@@ -144,3 +146,5 @@ void VideoStream::addFrame(const Bitmap<RGBColor>& frame)
     /* encode the image */
     encode(codecContext, picture, pkt, file);
 }
+
+#endif // FFMPEG_ENABLED

VideoStream.h

@@ -1,4 +1,7 @@
 #pragma once
+#ifdef FFMPEG_ENABLED
+
+
 #ifndef VIDEO_STREAM_H_
 #define VIDEO_STREAM_H_
 
@@ -37,3 +40,5 @@ public:
 };
 
 #endif // VIDEO_STREAM_H_
+
+#endif // FFMPEG_ENABLED

benchmarkdialog.cpp

@@ -94,6 +94,9 @@ void Benchmarker::start(void)
         if (mnd::Generator* gpud; gpud = devices[i].getGeneratorDouble()) {
             nTests++;
         }
+        if (mnd::Generator* gpu128; gpu128 = devices[i].getGenerator128()) {
+            nTests++;
+        }
     }
 
     double progress = 90.0 / nTests;
@@ -128,6 +131,11 @@ void Benchmarker::start(void)
             br.percentage += progress;
             emit update(br);
         }
+        if (mnd::Generator* gpu128; gpu128 = devices[i].getGenerator128()) {
+            gpu.push_back(benchmarkResult(*gpu128));
+            br.percentage += progress;
+            emit update(br);
+        }
     }
     printf("benchmark finished\n");
     emit update(br);

libmandel/include/ClGenerators.h

@@ -14,6 +14,7 @@ namespace mnd
     class ClGenerator;
     class ClGeneratorFloat;
     class ClGeneratorDouble;
+    class ClGenerator128;
 }
 
 
@@ -57,4 +58,16 @@ protected:
     virtual std::string getKernelCode(void) const;
 };
 
+
+class mnd::ClGenerator128 : public ClGenerator
+{
+public:
+    ClGenerator128(cl::Device device);
+    virtual ~ClGenerator128(void) = default;
+
+    virtual void generate(const MandelInfo& info, float* data);
+protected:
+    virtual std::string getKernelCode(void) const;
+};
+
 #endif // MANDEL_CLGENERATORS_H

libmandel/include/Mandel.h

@@ -28,6 +28,7 @@ private:
 
     std::unique_ptr<Generator> floatGenerator;
     std::unique_ptr<Generator> doubleGenerator;
+    std::unique_ptr<Generator> generator128;
     MandelDevice(void);
 public:
 
@@ -36,6 +37,7 @@ public:
 
     Generator* getGeneratorFloat(void) const;
     Generator* getGeneratorDouble(void) const;
+    Generator* getGenerator128(void) const;
 };
 
 

libmandel/src/ClGenerators.cpp

@@ -9,6 +9,7 @@ using namespace cl;
 using mnd::ClGenerator;
 using mnd::ClGeneratorFloat;
 using mnd::ClGeneratorDouble;
+using mnd::ClGenerator128;
 
 Platform getPlatform() {
     /* Returns the first platform found. */
@@ -273,4 +274,60 @@ std::string ClGeneratorDouble::getKernelCode(void) const
         //            "   A[index] = ((float)n) + 1 - (a * a + b * b - 16) / (256 - 16);\n"
         //        "   A[get_global_id(0)] = 5;"
         "}";
+}
+
+
+
+ClGenerator128::ClGenerator128(cl::Device device) :
+    ClGenerator{ device }
+{
+    context = Context{ device };
+    Program::Sources sources;
+
+    std::string kcode = this->getKernelCode();
+
+    sources.push_back({ kcode.c_str(), kcode.length() });
+
+    program = Program{ context, sources };
+    if (program.build({ device }) != CL_SUCCESS) {
+        throw std::string(program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(device));
+    }
+
+    queue = CommandQueue(context, device);
+}
+
+
+void ClGenerator128::generate(const mnd::MandelInfo& info, float* data)
+{
+    ::size_t bufferSize = info.bWidth * info.bHeight * sizeof(float);
+
+    Buffer buffer_A(context, CL_MEM_WRITE_ONLY, bufferSize);
+    float pixelScaleX = info.view.width / info.bWidth;
+    float pixelScaleY = info.view.height / info.bHeight;
+
+    Kernel iterate = Kernel(program, "iterate");
+    iterate.setArg(0, buffer_A);
+    iterate.setArg(1, int(info.bWidth));
+    iterate.setArg(2, double(info.view.x));
+    iterate.setArg(3, double(info.view.y));
+    iterate.setArg(4, double(pixelScaleX));
+    iterate.setArg(5, double(pixelScaleY));
+    iterate.setArg(6, int(info.maxIter));
+
+    queue.enqueueNDRangeKernel(iterate, 0, NDRange(info.bWidth * info.bHeight));
+    queue.enqueueReadBuffer(buffer_A, CL_TRUE, 0, bufferSize, data);
+}
+
+#include <string>
+#include <fstream>
+#include <streambuf>
+
+std::string ClGenerator128::getKernelCode(void) const
+{
+    //fprintf(stderr, "starting file read\n");
+    std::ifstream t("mandel128.cl");
+    std::string str((std::istreambuf_iterator<char>(t)),
+        std::istreambuf_iterator<char>());
+    //fprintf(stderr, "%s\n", str);
+    return str;
 }

libmandel/src/mandel.cpp

@@ -41,6 +41,15 @@ mnd::Generator* MandelDevice::getGeneratorDouble(void) const
 }
 
 
+mnd::Generator* MandelDevice::getGenerator128(void) const
+{
+    if (generator128)
+        return generator128.get();
+    else
+        return nullptr;
+}
+
+
 MandelContext::MandelContext(void)
 {
     if (cpuInfo.hasAvx()) {
@@ -68,7 +77,7 @@ std::vector<MandelDevice> MandelContext::createDevices(void)
 
     std::vector<cl::Platform> platforms;
     cl::Platform::get(&platforms);
-    //platforms.erase(platforms.begin() + 1);
+    platforms.erase(platforms.begin() + 1);
 
     for (auto& platform : platforms) {
         std::string name = platform.getInfo<CL_PLATFORM_NAME>();
@@ -109,6 +118,14 @@ std::vector<MandelDevice> MandelContext::createDevices(void)
                 catch (const std::string& err) {
                 }
             }
+
+            try {
+                md.generator128 = std::make_unique<ClGenerator128>(device);
+            }
+            catch (const std::string& err) {
+                fprintf(stderr, "error creating 128bit cl generator: %s\n", err.c_str());
+            }
+
             mandelDevices.push_back(std::move(md));
         }
     }