Skip to content
Permalink

Comparing changes

Choose two branches to see what’s changed or to start a new pull request. If you need to, you can also or learn more about diff comparisons.

Open a pull request

Create a new pull request by comparing changes across two branches. If you need to, you can also . Learn more about diff comparisons here.
base repository: ngscopeclient/scopehal
Failed to load repositories. Confirm that selected base ref is valid, then try again.
Loading
base: 962b65daeea2
Choose a base ref
...
head repository: ngscopeclient/scopehal
Failed to load repositories. Confirm that selected head ref is valid, then try again.
Loading
compare: cb4eb5bee7e4
Choose a head ref
  • 1 commit
  • 6 files changed
  • 1 contributor

Commits on Dec 19, 2020

  1. Initial OpenCL kernel creation logic. Seems to work. Test kernel does…

    …n't actually do anything yet and is never called.
    azonenberg committed Dec 19, 2020
    Copy the full SHA
    cb4eb5b View commit details
Showing with 144 additions and 3 deletions.
  1. +52 −1 scopehal/Filter.cpp
  2. +14 −1 scopehal/Filter.h
  3. +37 −0 scopehal/scopehal.cpp
  4. +3 −0 scopehal/scopehal.h
  5. +6 −1 scopeprotocols/FIRFilter.cpp
  6. +32 −0 scopeprotocols/kernels/FIRFilter.cl
53 changes: 52 additions & 1 deletion scopehal/Filter.cpp
Original file line number Diff line number Diff line change
@@ -62,18 +62,69 @@ Gdk::Color Filter::m_standardColors[STANDARD_COLOR_COUNT] =
Filter::Filter(
OscilloscopeChannel::ChannelType type,
const string& color,
Category cat)
Category cat,
const string& kernelPath,
const string& kernelName)
: OscilloscopeChannel(NULL, "", type, color, 1) //TODO: handle this better?
, m_category(cat)
, m_dirty(true)
, m_usingDefault(true)
{
m_physical = false;
m_filters.emplace(this);

//Load our OpenCL kernel, if we have one
#ifdef HAVE_OPENCL

m_kernel = NULL;
m_program = NULL;

if(kernelPath != "")
{
//Build the program
string kernelSource = ReadFile(kernelPath);
cl::Program::Sources sources(1, make_pair(&kernelSource[0], kernelSource.length()));
m_program = new cl::Program(*g_clContext, sources);
cl_int err = m_program->build(g_contextDevices);
if(err != CL_SUCCESS)
{
LogError("Failed to build OpenCL program from %s (code %d)\n", kernelPath.c_str(), err);
string log;
m_program->getBuildInfo<string>(g_contextDevices[0], CL_PROGRAM_BUILD_LOG, &log);
LogDebug("Build log:\n");
LogDebug("%s\n", log.c_str());

delete m_program;
m_program = NULL;
return;
}

//Make the kernel
m_kernel = new cl::Kernel(*m_program, kernelName.c_str(), &err);
if(err != CL_SUCCESS)
{
LogError("Failed to create OpenCL kernel %s in %s (code %d)\n",
kernelName.c_str(), kernelPath.c_str(), err);

delete m_program;
delete m_kernel;
m_program = NULL;
m_kernel = NULL;
}
}

#endif
}

Filter::~Filter()
{
#ifdef HAVE_OPENCL
delete m_kernel;
delete m_program;
m_kernel = NULL;
m_program = NULL;
#endif

m_filters.erase(this);

for(auto c : m_inputs)
15 changes: 14 additions & 1 deletion scopehal/Filter.h
Original file line number Diff line number Diff line change
@@ -62,7 +62,12 @@ class Filter : public OscilloscopeChannel
CAT_RF //Frequency domain analysis (FFT etc) and other RF stuff
};

Filter(OscilloscopeChannel::ChannelType type, const std::string& color, Category cat);
Filter(
OscilloscopeChannel::ChannelType type,
const std::string& color,
Category cat,
const std::string& kernelPath = "",
const std::string& kernelName = "");
virtual ~Filter();

/**
@@ -238,6 +243,14 @@ class Filter : public OscilloscopeChannel
//Common text formatting
virtual std::string GetTextForAsciiChannel(int i, size_t stream);

#ifdef HAVE_OPENCL

//OpenCL state
cl::Program* m_program;
cl::Kernel* m_kernel;

#endif

public:
typedef Filter* (*CreateProcType)(const std::string&);
static void DoAddDecoderClass(const std::string& name, CreateProcType proc);
37 changes: 37 additions & 0 deletions scopehal/scopehal.cpp
Original file line number Diff line number Diff line change
@@ -70,7 +70,10 @@ bool g_hasAvx512DQ = false;
bool g_hasAvx512VL = false;
bool g_hasAvx2 = false;

#ifdef HAVE_OPENCL
cl::Context* g_clContext = NULL;
vector<cl::Device> g_contextDevices;
#endif

AlignedAllocator<float, 32> g_floatVectorAllocator;

@@ -198,6 +201,8 @@ void DetectGPUFeatures()
g_clContext = NULL;
return;
}

g_contextDevices = g_clContext->getInfo<CL_CONTEXT_DEVICES>();
}
}
}
@@ -532,3 +537,35 @@ uint64_t next_pow2(uint64_t v)
return v;
#endif
}

/**
@brief Returns the contents of a file
*/
string ReadFile(const string& path)
{
//Read the file
FILE* fp = fopen(path.c_str(), "rb");
if(!fp)
{
LogWarning("ReadFile: Could not open file \"%s\"\n", path.c_str());
return "";
}
fseek(fp, 0, SEEK_END);
size_t fsize = ftell(fp);
fseek(fp, 0, SEEK_SET);
char* buf = new char[fsize + 1];
if(fsize != fread(buf, 1, fsize, fp))
{
LogWarning("ReadFile: Could not read file \"%s\"\n", path.c_str());
delete[] buf;
fclose(fp);
return "";
}
buf[fsize] = 0;
fclose(fp);

string ret(buf);
delete[] buf;

return ret;
}
3 changes: 3 additions & 0 deletions scopehal/scopehal.h
Original file line number Diff line number Diff line change
@@ -101,6 +101,8 @@ std::string Trim(const std::string& str);
std::string TrimQuotes(const std::string& str);
std::string BaseName(const std::string& path);

std::string ReadFile(const std::string& path);

std::string to_string_sci(double d);

void TransportStaticInit();
@@ -131,6 +133,7 @@ extern bool g_hasAvx2;

#ifdef HAVE_OPENCL
extern cl::Context* g_clContext;
extern std::vector<cl::Device> g_contextDevices;
#endif

#endif
7 changes: 6 additions & 1 deletion scopeprotocols/FIRFilter.cpp
Original file line number Diff line number Diff line change
@@ -37,7 +37,12 @@ using namespace std;
// Construction / destruction

FIRFilter::FIRFilter(const string& color)
: Filter(OscilloscopeChannel::CHANNEL_TYPE_ANALOG, color, CAT_MATH)
: Filter(
OscilloscopeChannel::CHANNEL_TYPE_ANALOG,
color,
CAT_MATH,
"kernels/FIRFilter.cl",
"FIRFilter")
, m_filterTypeName("Filter Type")
, m_filterLengthName("Length")
, m_stopbandAttenName("Stopband Attenuation")
32 changes: 32 additions & 0 deletions scopeprotocols/kernels/FIRFilter.cl
Original file line number Diff line number Diff line change
@@ -0,0 +1,32 @@
/***********************************************************************************************************************
* *
* LIBSCOPEHAL v0.1 *
* *
* Copyright (c) 2012-2020 Andrew D. Zonenberg *
* All rights reserved. *
* *
* Redistribution and use in source and binary forms, with or without modification, are permitted provided that the *
* following conditions are met: *
* *
* * Redistributions of source code must retain the above copyright notice, this list of conditions, and the *
* following disclaimer. *
* *
* * Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the *
* following disclaimer in the documentation and/or other materials provided with the distribution. *
* *
* * Neither the name of the author nor the names of any contributors may be used to endorse or promote products *
* derived from this software without specific prior written permission. *
* *
* THIS SOFTWARE IS PROVIDED BY THE AUTHORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED *
* TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL *
* THE AUTHORS BE HELD LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES *
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR *
* BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT *
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE *
* POSSIBILITY OF SUCH DAMAGE. *
* *
***********************************************************************************************************************/

__kernel void FIRFilter()
{
}