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: f5187e9efc18
Choose a base ref
...
head repository: ngscopeclient/scopehal
Failed to load repositories. Confirm that selected head ref is valid, then try again.
Loading
compare: e7d8e19726bf
Choose a head ref
  • 1 commit
  • 6 files changed
  • 1 contributor

Commits on Dec 19, 2020

  1. Verified

    This commit was created on GitHub.com and signed with GitHub’s verified signature. The key has expired.
    Copy the full SHA
    e7d8e19 View commit details
Showing with 189 additions and 99 deletions.
  1. +19 −21 scopehal/Filter.cpp
  2. +74 −73 scopehal/scopehal.cpp
  3. +1 −0 scopehal/scopehal.h
  4. +64 −0 scopeprotocols/FIRFilter.cpp
  5. +13 −4 scopeprotocols/FIRFilter.h
  6. +18 −1 scopeprotocols/kernels/FIRFilter.cl
40 changes: 19 additions & 21 deletions scopehal/Filter.cpp
Original file line number Diff line number Diff line change
@@ -81,36 +81,34 @@ Filter::Filter(

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)
try
{
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;
string kernelSource = ReadFile(kernelPath);
cl::Program::Sources sources(1, make_pair(&kernelSource[0], kernelSource.length()));
m_program = new cl::Program(*g_clContext, sources);
m_program->build(g_contextDevices);
m_kernel = new cl::Kernel(*m_program, kernelName.c_str());
}

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

if(e.err() == CL_BUILD_PROGRAM_FAILURE)
{
LogError("Failed to build OpenCL program from %s\n", kernelPath.c_str());
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;
delete m_kernel;
m_program = NULL;
m_kernel = NULL;
return;
}

}

#endif
147 changes: 74 additions & 73 deletions scopehal/scopehal.cpp
Original file line number Diff line number Diff line change
@@ -124,88 +124,89 @@ void DetectCPUFeatures()
void DetectGPUFeatures()
{
#ifdef HAVE_OPENCL
LogDebug("Detecting OpenCL devices...\n");
LogIndenter li;

//Find platforms and print info
vector<cl::Platform> platforms;
cl::Platform::get(&platforms);
if(platforms.empty())
{
LogNotice("No platforms found, disabling OpenCL\n");
return;
}
else
try
{
for(size_t i=0; i<platforms.size(); i++)
{
LogDebug("Platform %zu\n", i);
LogIndenter li2;

string name;
string profile;
string vendor;
string version;
platforms[i].getInfo(CL_PLATFORM_NAME, &name);
platforms[i].getInfo(CL_PLATFORM_PROFILE, &profile);
platforms[i].getInfo(CL_PLATFORM_VENDOR, &vendor);
platforms[i].getInfo(CL_PLATFORM_VERSION, &version);
LogDebug("CL_PLATFORM_NAME = %s\n", name.c_str());
LogDebug("CL_PLATFORM_PROFILE = %s\n", profile.c_str());
LogDebug("CL_PLATFORM_VENDOR = %s\n", vendor.c_str());
LogDebug("CL_PLATFORM_VERSION = %s\n", version.c_str());

vector<cl::Device> devices;
platforms[i].getDevices(CL_DEVICE_TYPE_GPU, &devices);
if(devices.empty())
LogDebug("No GPUs found\n");
for(size_t j=0; j<devices.size(); j++)
{
LogDebug("Device %zu\n", j);
LogIndenter li3;

string dname;
string dcvers;
string dprof;
string dvendor;
string dversion;
string ddversion;
devices[j].getInfo(CL_DEVICE_NAME, &dname);
devices[j].getInfo(CL_DEVICE_OPENCL_C_VERSION, &dcvers);
devices[j].getInfo(CL_DEVICE_PROFILE, &dprof);
devices[j].getInfo(CL_DEVICE_VENDOR, &dvendor);
devices[j].getInfo(CL_DEVICE_VERSION, &dversion);
devices[j].getInfo(CL_DRIVER_VERSION, &ddversion);

LogDebug("CL_DEVICE_NAME = %s\n", dname.c_str());
LogDebug("CL_DEVICE_OPENCL_C_VERSION = %s\n", dcvers.c_str());
LogDebug("CL_DEVICE_PROFILE = %s\n", dprof.c_str());
LogDebug("CL_DEVICE_VENDOR = %s\n", dvendor.c_str());
LogDebug("CL_DEVICE_VERSION = %s\n", dversion.c_str());
LogDebug("CL_DRIVER_VERSION = %s\n", ddversion.c_str());
}
LogDebug("Detecting OpenCL devices...\n");
LogIndenter li;

//For now, create a context on the first device of the first detected platform and hope for the best
if(!g_clContext)
//Find platforms and print info
vector<cl::Platform> platforms;
cl::Platform::get(&platforms);
if(platforms.empty())
{
LogNotice("No platforms found, disabling OpenCL\n");
return;
}
else
{
for(size_t i=0; i<platforms.size(); i++)
{
vector<cl::Device> devs;
devs.push_back(devices[0]);

//Passing CL_CONTEXT_PLATFORM as parameters seems to make context creation fail. Weird.
cl_int err;
g_clContext = new cl::Context(devs, NULL, NULL, NULL, &err);
if(err != CL_SUCCESS)
LogDebug("Platform %zu\n", i);
LogIndenter li2;

string name;
string profile;
string vendor;
string version;
platforms[i].getInfo(CL_PLATFORM_NAME, &name);
platforms[i].getInfo(CL_PLATFORM_PROFILE, &profile);
platforms[i].getInfo(CL_PLATFORM_VENDOR, &vendor);
platforms[i].getInfo(CL_PLATFORM_VERSION, &version);
LogDebug("CL_PLATFORM_NAME = %s\n", name.c_str());
LogDebug("CL_PLATFORM_PROFILE = %s\n", profile.c_str());
LogDebug("CL_PLATFORM_VENDOR = %s\n", vendor.c_str());
LogDebug("CL_PLATFORM_VERSION = %s\n", version.c_str());

vector<cl::Device> devices;
platforms[i].getDevices(CL_DEVICE_TYPE_GPU, &devices);
if(devices.empty())
LogDebug("No GPUs found\n");
for(size_t j=0; j<devices.size(); j++)
{
LogNotice("OpenCL context creation failed (code %d), disabling OpenCL\n", err);
delete g_clContext;
g_clContext = NULL;
return;
LogDebug("Device %zu\n", j);
LogIndenter li3;

string dname;
string dcvers;
string dprof;
string dvendor;
string dversion;
string ddversion;
devices[j].getInfo(CL_DEVICE_NAME, &dname);
devices[j].getInfo(CL_DEVICE_OPENCL_C_VERSION, &dcvers);
devices[j].getInfo(CL_DEVICE_PROFILE, &dprof);
devices[j].getInfo(CL_DEVICE_VENDOR, &dvendor);
devices[j].getInfo(CL_DEVICE_VERSION, &dversion);
devices[j].getInfo(CL_DRIVER_VERSION, &ddversion);

LogDebug("CL_DEVICE_NAME = %s\n", dname.c_str());
LogDebug("CL_DEVICE_OPENCL_C_VERSION = %s\n", dcvers.c_str());
LogDebug("CL_DEVICE_PROFILE = %s\n", dprof.c_str());
LogDebug("CL_DEVICE_VENDOR = %s\n", dvendor.c_str());
LogDebug("CL_DEVICE_VERSION = %s\n", dversion.c_str());
LogDebug("CL_DRIVER_VERSION = %s\n", ddversion.c_str());
}

g_contextDevices = g_clContext->getInfo<CL_CONTEXT_DEVICES>();
//For now, create a context on the first device of the first detected platform and hope for the best
if(!g_clContext)
{
vector<cl::Device> devs;
devs.push_back(devices[0]);

//Passing CL_CONTEXT_PLATFORM as parameters seems to make context creation fail. Weird.
g_clContext = new cl::Context(devs, NULL, NULL, NULL);
g_contextDevices = g_clContext->getInfo<CL_CONTEXT_DEVICES>();
}
}
}
}
catch(const cl::Error& e)
{
LogError("OpenCL error: %s (%d)\n", e.what(), e.err() );
delete g_clContext;
g_clContext = NULL;
return;
}

#else
LogNotice("OpenCL support not present at compile time. GPU acceleration disabled.\n");
1 change: 1 addition & 0 deletions scopehal/scopehal.h
Original file line number Diff line number Diff line change
@@ -56,6 +56,7 @@
#include "config.h"
#ifdef HAVE_OPENCL
#define CL_TARGET_OPENCL_VERSION 120
#define __CL_ENABLE_EXCEPTIONS
#include <CL/cl.hpp>
#endif

64 changes: 64 additions & 0 deletions scopeprotocols/FIRFilter.cpp
Original file line number Diff line number Diff line change
@@ -251,6 +251,12 @@ void FIRFilter::DoFilterKernel(
float& vmin,
float& vmax)
{
#ifdef HAVE_OPENCL
if(g_clContext && m_kernel)
DoFilterKernelOpenCL(coefficients, din, cap, vmin, vmax);
else
#endif

if(g_hasAvx512F)
DoFilterKernelAVX512F(coefficients, din, cap, vmin, vmax);
else if(g_hasAvx2)
@@ -259,6 +265,64 @@ void FIRFilter::DoFilterKernel(
DoFilterKernelGeneric(coefficients, din, cap, vmin, vmax);
}

#ifdef HAVE_OPENCL
void FIRFilter::DoFilterKernelOpenCL(
std::vector<float>& coefficients,
AnalogWaveform* din,
AnalogWaveform* cap,
float& vmin,
float& vmax)
{
//Setup
size_t len = din->m_samples.size();
size_t filterlen = coefficients.size();
size_t end = len - filterlen;

try
{
//Allocate memory and copy to the GPU
cl::Buffer inbuf(*g_clContext, din->m_samples.begin(), din->m_samples.end(), true, true, NULL);
cl::Buffer coeffbuf(*g_clContext, coefficients.begin(), coefficients.end(), true, true, NULL);
cl::Buffer outbuf(*g_clContext, CL_MEM_WRITE_ONLY, end * sizeof(float));

//Run the filter
cl::Event event;
cl::CommandQueue queue(*g_clContext, g_contextDevices[0], 0);
m_kernel->setArg(0, inbuf);
m_kernel->setArg(1, coeffbuf);
m_kernel->setArg(2, outbuf);
m_kernel->setArg(3, filterlen);
m_kernel->setArg(4, end);
queue.enqueueNDRangeKernel(
*m_kernel,
cl::NullRange,
cl::NDRange(end, 1),
cl::NullRange,
NULL,
&event);

//Done, copy output
event.wait();
cl::copy(queue, outbuf, cap->m_samples.begin(), cap->m_samples.end() );
}
catch(const cl::Error& e)
{
LogError("OpenCL error: %s (%d)\n", e.what(), e.err() );
}

//Final reduction CPU-side for now
vmin = FLT_MAX;
vmax = -FLT_MAX;
for(size_t i=0; i<end; i++)
{
float v = cap->m_samples[i];
vmin = min(vmin, v);
vmax = max(vmax, v);
}
}
#endif


/**
@brief Performs a FIR filter (does not assume symmetric)
*/
17 changes: 13 additions & 4 deletions scopeprotocols/FIRFilter.h
Original file line number Diff line number Diff line change
@@ -60,7 +60,7 @@ class FIRFilter : public Filter

PROTOCOL_DECODER_INITPROC(FIRFilter)

static void DoFilterKernel(
void DoFilterKernel(
std::vector<float>& coefficients,
AnalogWaveform* din,
AnalogWaveform* cap,
@@ -86,21 +86,30 @@ class FIRFilter : public Filter

static float Bessel(float x);

static void DoFilterKernelGeneric(
void DoFilterKernelGeneric(
std::vector<float>& coefficients,
AnalogWaveform* din,
AnalogWaveform* cap,
float& vmin,
float& vmax);

static void DoFilterKernelAVX2(
#ifdef HAVE_OPENCL
void DoFilterKernelOpenCL(
std::vector<float>& coefficients,
AnalogWaveform* din,
AnalogWaveform* cap,
float& vmin,
float& vmax);
#endif

void DoFilterKernelAVX2(
std::vector<float>& coefficients,
AnalogWaveform* din,
AnalogWaveform* cap,
float& vmin,
float& vmax);

static void DoFilterKernelAVX512F(
void DoFilterKernelAVX512F(
std::vector<float>& coefficients,
AnalogWaveform* din,
AnalogWaveform* cap,
19 changes: 18 additions & 1 deletion scopeprotocols/kernels/FIRFilter.cl
Original file line number Diff line number Diff line change
@@ -27,6 +27,23 @@
* *
***********************************************************************************************************************/

__kernel void FIRFilter()
__kernel void FIRFilter(
__global const float* din,
__constant const float* coefficients,
__global float* dout,
unsigned long filterlen,
unsigned long end
)
{
//Make sure we're actually in the block
unsigned long i = get_global_id(0);
if(i > end)
return;

//FIR reduction
float v = 0;
for(unsigned long j=0; j<filterlen; j++)
v += din[i+j] * coefficients[j];

dout[i] = v;
}