[Top][All Lists]
[Date Prev][Date Next][Thread Prev][Thread Next][Date Index][Thread Index]
[Commit-gnuradio] r9818 - in gnuradio/branches/developers/nldudok1/gpgpu
|
From: |
nldudok1 |
|
Subject: |
[Commit-gnuradio] r9818 - in gnuradio/branches/developers/nldudok1/gpgpu-wip: gr-cuda/src/lib testbed testbed/wfm |
|
Date: |
Wed, 22 Oct 2008 08:35:58 -0600 (MDT) |
Author: nldudok1
Date: 2008-10-22 08:35:54 -0600 (Wed, 22 Oct 2008)
New Revision: 9818
Added:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda_iir_filter2_fff.cc
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda_iir_filter2_fff.h
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda_iir_filter2_fff.i
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda_iir_filter_fff.cc
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda_iir_filter_fff.h
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda_iir_filter_fff.i
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cudai_iir.cu
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cudai_iir.h
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cudai_iir_kernel.cu
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cudai_iir_kernel.h
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/run_13.2.iir.py
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/wfm/check_wfm2.py
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/wfm/fm_emph_cuda.py
Removed:
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/core
Modified:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/Makefile.am
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/Makefile.in
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda.i
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cudai_general_kernel.cu
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/check_13.1.qdemod.py
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/check_13.py
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/wfm/check_wfm.py
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/wfm/cuda_wfm_rcv.py
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/wfm/cuda_wfm_rcv.pyc
Log:
added iir filter
Modified:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/Makefile.am
===================================================================
--- gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/Makefile.am
2008-10-22 08:55:52 UTC (rev 9817)
+++ gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/Makefile.am
2008-10-22 14:35:54 UTC (rev 9818)
@@ -121,7 +121,11 @@
cudai_fft.cc \
cudai_general_kernel.cu \
cuda_fft_filter_ccc.cc \
- cuda_add_cc.cc
+ cuda_add_cc.cc \
+ cudai_iir.cu \
+ cuda_iir_filter_fff.cc \
+ cudai_iir_kernel.cu \
+ cuda_iir_filter2_fff.cc
#cuda_fir_fff_cuda.cc cuda_fir_fff.h cuda_fir_fff_kernel_opt.cu
@@ -187,7 +191,11 @@
cudai_general_kernel.h \
cuda_fft_filter_ccc.h \
gr_cuda.h \
- cuda_add_cc.h
+ cuda_add_cc.h \
+ cudai_iir.h \
+ cuda_iir_filter_fff.h \
+ cudai_iir_kernel.h \
+ cuda_iir_filter2_fff.h
# These swig headers get installed in ${prefix}/include/gnuradio/swig
@@ -200,10 +208,11 @@
cuda_multiply_const_ff.i \
cuda_nop.i \
cuda_fft_filter_ccc.i \
- cuda_add_cc.i
+ cuda_add_cc.i \
+ cuda_iir_filter_fff.i \
+ cuda_iir_filter2_fff.i
-
#noinst_HEADERS = \
# cuda_quadrature_demod_cf_kernel.cu
Modified:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/Makefile.in
===================================================================
--- gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/Makefile.in
2008-10-22 08:55:52 UTC (rev 9817)
+++ gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/Makefile.in
2008-10-22 14:35:54 UTC (rev 9818)
@@ -124,7 +124,9 @@
cuda_null_sink.lo cuda_head.lo cuda_test.lo \
cuda_multiply_const_ff.lo cuda_multiply_const_ff_kernel.lo \
cuda_nop.lo cudai_fft.lo cudai_general_kernel.lo \
- cuda_fft_filter_ccc.lo cuda_add_cc.lo
+ cuda_fft_filter_ccc.lo cuda_add_cc.lo cudai_iir.lo \
+ cuda_iir_filter_fff.lo cudai_iir_kernel.lo \
+ cuda_iir_filter2_fff.lo
_cuda_la_OBJECTS = $(am__cuda_la_OBJECTS)
_cuda_la_LINK = $(LIBTOOL) --tag=CXX $(AM_LIBTOOLFLAGS) \
$(LIBTOOLFLAGS) --mode=link $(CXXLD) $(AM_CXXFLAGS) \
@@ -406,7 +408,11 @@
cudai_fft.cc \
cudai_general_kernel.cu \
cuda_fft_filter_ccc.cc \
- cuda_add_cc.cc
+ cuda_add_cc.cc \
+ cudai_iir.cu \
+ cuda_iir_filter_fff.cc \
+ cudai_iir_kernel.cu \
+ cuda_iir_filter2_fff.cc
#cuda_fir_fff_cuda.cc cuda_fir_fff.h cuda_fir_fff_kernel_opt.cu
@@ -465,7 +471,11 @@
cudai_general_kernel.h \
cuda_fft_filter_ccc.h \
gr_cuda.h \
- cuda_add_cc.h
+ cuda_add_cc.h \
+ cudai_iir.h \
+ cuda_iir_filter_fff.h \
+ cudai_iir_kernel.h \
+ cuda_iir_filter2_fff.h
# These swig headers get installed in ${prefix}/include/gnuradio/swig
@@ -478,7 +488,9 @@
cuda_multiply_const_ff.i \
cuda_nop.i \
cuda_fft_filter_ccc.i \
- cuda_add_cc.i
+ cuda_add_cc.i \
+ cuda_iir_filter_fff.i \
+ cuda_iir_filter2_fff.i
#noinst_HEADERS = \
@@ -566,6 +578,8 @@
@AMDEP_TRUE@@am__include@ @address@hidden/$(DEPDIR)/address@hidden@
@AMDEP_TRUE@@am__include@ @address@hidden/$(DEPDIR)/address@hidden@
@AMDEP_TRUE@@am__include@ @address@hidden/$(DEPDIR)/address@hidden@
address@hidden@@am__include@ @address@hidden/$(DEPDIR)/address@hidden@
address@hidden@@am__include@ @address@hidden/$(DEPDIR)/address@hidden@
@AMDEP_TRUE@@am__include@ @address@hidden/$(DEPDIR)/address@hidden@
@AMDEP_TRUE@@am__include@ @address@hidden/$(DEPDIR)/address@hidden@
@AMDEP_TRUE@@am__include@ @address@hidden/$(DEPDIR)/address@hidden@
Modified: gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda.i
===================================================================
--- gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda.i
2008-10-22 08:55:52 UTC (rev 9817)
+++ gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda.i
2008-10-22 14:35:54 UTC (rev 9818)
@@ -21,6 +21,8 @@
#include "cuda_nop.h"
#include "cuda_fft_filter_ccc.h"
#include "cuda_add_cc.h"
+#include "cuda_iir_filter_fff.h"
+#include "cuda_iir_filter2_fff.h"
#include <stdexcept>
%}
@@ -32,6 +34,8 @@
%include "cuda_nop.i"
%include "cuda_add_cc.i"
%include "cuda_fft_filter_ccc.i"
+%include "cuda_iir_filter_fff.i"
+%include "cuda_iir_filter2_fff.i"
GR_SWIG_BLOCK_MAGIC(cuda,cuda_to_host)
Added:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda_iir_filter2_fff.cc
===================================================================
---
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda_iir_filter2_fff.cc
(rev 0)
+++
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda_iir_filter2_fff.cc
2008-10-22 14:35:54 UTC (rev 9818)
@@ -0,0 +1,194 @@
+/* -*- c++ -*- */
+/*
+ * Copyright 2004,2008 Free Software Foundation, Inc.
+ *
+ * This file is part of GNU Radio
+ *
+ * GNU Radio is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU General Public License as published by
+ * the Free Software Foundation; either version 3, or (at your option)
+ * any later version.
+ *
+ * GNU Radio is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with GNU Radio; see the file COPYING. If not, write to
+ * the Free Software Foundation, Inc., 51 Franklin Street,
+ * Boston, MA 02110-1301, USA.
+ */
+
+#ifdef HAVE_CONFIG_H
+#include "config.h"
+#endif
+
+#include <stdexcept>
+
+#include <cuda_iir_filter2_fff.h>
+#include <gr_io_signature.h>
+#include <stdio.h>
+
+#include <cuda_runtime.h>
+#include <cutil.h>
+
+#include "gr_cuda.h"
+
+#include <cudai_iir_kernel.h>
+
+#include <gri_memcpy_cuda_kernel.h>
+
+cuda_iir_filter2_fff_sptr
+cuda_make_iir_filter2_fff (const std::vector<float> &fftaps,
+ const std::vector<float> &fbtaps) throw
(std::invalid_argument)
+{
+ return cuda_iir_filter2_fff_sptr (new cuda_iir_filter2_fff (fftaps, fbtaps));
+}
+
+cuda_iir_filter2_fff::cuda_iir_filter2_fff (const std::vector<float> &fftaps,
+ const std::vector<float> &fbtaps) throw
(std::invalid_argument)
+
+ : gr_sync_block ("iir_filter2_fff",
+ gr_make_io_signature (1, 1, sizeof (float),GR_BUFFER_CUDA),
+ gr_make_io_signature (1, 1, sizeof (float),GR_BUFFER_CUDA)),
+ d_iir (fftaps, fbtaps),
+ d_updated (false),
+ dd_m(0),
+ dd_n(0),
+ dd_latest_m(0),
+ dd_latest_n(0),
+ dd_fftaps(0),
+ dd_fbtaps(0),
+ dd_prev_input(0),
+ dd_prev_output(0)
+
+{
+ d_new_fftaps = fftaps;
+ d_new_fbtaps = fbtaps;
+ // fprintf (stderr, "cuda_iir_filter2_fff::ctor\n");
+}
+
+cuda_iir_filter2_fff::~cuda_iir_filter2_fff ()
+{
+ //free(d_params);
+ GR_CUDA_SAFE_CALL(cudaFree(dd_latest_n)); dd_latest_n=NULL;
+ GR_CUDA_SAFE_CALL(cudaFree(dd_latest_m)); dd_latest_m=NULL;
+ GR_CUDA_SAFE_CALL(cudaFree(dd_fftaps)); dd_fftaps=NULL;
+ GR_CUDA_SAFE_CALL(cudaFree(dd_fbtaps)); dd_fbtaps=NULL;
+ GR_CUDA_SAFE_CALL(cudaFree(dd_prev_output)); dd_prev_output=NULL;
+ GR_CUDA_SAFE_CALL(cudaFree(dd_prev_input)); dd_prev_input=NULL;
+ GR_CUDA_SAFE_CALL(cudaFree(dd_n)); dd_n=NULL;
+ GR_CUDA_SAFE_CALL(cudaFree(dd_m)); dd_m=NULL;
+}
+
+bool
+cuda_iir_filter2_fff::start()
+{
+ fprintf(stderr,"cuda_iir_filter2_fff::start()\n");
+
+
//d_params=(cudai_general_kernel_params*)malloc(sizeof(cudai_general_kernel_params));
+ //cudai_get_general_kernel_params ( d_params,1);
+ if(dd_latest_n) GR_CUDA_SAFE_CALL(cudaFree(dd_latest_n));
+ if(dd_latest_m) GR_CUDA_SAFE_CALL(cudaFree(dd_latest_m));
+ if(dd_fftaps) GR_CUDA_SAFE_CALL(cudaFree(dd_fftaps));
+ if(dd_fbtaps) GR_CUDA_SAFE_CALL(cudaFree(dd_fbtaps));
+ if(dd_prev_output) GR_CUDA_SAFE_CALL(cudaFree(dd_prev_output));
+ if(dd_prev_input) GR_CUDA_SAFE_CALL(cudaFree(dd_prev_input));
+ if(dd_n) GR_CUDA_SAFE_CALL(cudaFree(dd_n));
+ if(dd_m) GR_CUDA_SAFE_CALL(cudaFree(dd_m));
+
+ GR_CUDA_SAFE_CALL(cudaMalloc((void**) &(dd_latest_n),sizeof(int)));
+ GR_CUDA_SAFE_CALL(cudaMalloc((void**) &(dd_latest_m),sizeof(int)));
+ GR_CUDA_SAFE_CALL(cudaMalloc((void**) &(dd_n),sizeof(unsigned int)));
+ GR_CUDA_SAFE_CALL(cudaMalloc((void**) &(dd_m),sizeof(unsigned int)));
+ dd_fftaps=NULL;
+ dd_fbtaps=NULL;
+ dd_prev_output=NULL;
+ dd_prev_input=NULL;
+ set_kernel_taps( d_new_fftaps, d_new_fbtaps);
+ return true;
+}
+
+bool
+cuda_iir_filter2_fff::stop()
+{
+ fprintf(stderr,"cuda_iir_filter2_fff::stop\n");
+ //free(d_params);
+ GR_CUDA_SAFE_CALL(cudaFree(dd_latest_n)); dd_latest_n=NULL;
+ GR_CUDA_SAFE_CALL(cudaFree(dd_latest_m)); dd_latest_m=NULL;
+ GR_CUDA_SAFE_CALL(cudaFree(dd_fftaps)); dd_fftaps=NULL;
+ GR_CUDA_SAFE_CALL(cudaFree(dd_fbtaps)); dd_fbtaps=NULL;
+ GR_CUDA_SAFE_CALL(cudaFree(dd_prev_output)); dd_prev_output=NULL;
+ GR_CUDA_SAFE_CALL(cudaFree(dd_prev_input)); dd_prev_input=NULL;
+ GR_CUDA_SAFE_CALL(cudaFree(dd_n)); dd_n=NULL;
+ GR_CUDA_SAFE_CALL(cudaFree(dd_m)); dd_m=NULL;
+ return true;
+}
+
+void
+cuda_iir_filter2_fff::set_taps (const std::vector<float> &fftaps,
+ const std::vector<float> &fbtaps) throw
(std::invalid_argument)
+{
+
+ d_new_fftaps = fftaps;
+ d_new_fbtaps = fbtaps;
+ d_updated = true;
+}
+
+ /*!
+ * \brief copy new taps to cuda device.
+ */
+ void
+cuda_iir_filter2_fff::set_kernel_taps (const std::vector<float> &fftaps,
+ const std::vector<float> &fbtaps) throw (std::invalid_argument)
+ {
+ fprintf(stderr,"cuda_iir_filter2_fff::set_kernel_taps(..)\n");
+
+ GR_CUDA_SAFE_CALL(cudaMemset(dd_latest_n,0,sizeof(int)));
+ GR_CUDA_SAFE_CALL(cudaMemset(dd_latest_m,0,sizeof(int)));
+
+ int n = fftaps.size ();
+ GR_CUDA_SAFE_CALL(cudaMemcpy(dd_n,&n,sizeof(n),cudaMemcpyHostToDevice));
+ int m = fbtaps.size ();
+ GR_CUDA_SAFE_CALL(cudaMemcpy(dd_m,&m,sizeof(m),cudaMemcpyHostToDevice));
+ if(dd_fftaps) GR_CUDA_SAFE_CALL(cudaFree(dd_fftaps));
+ GR_CUDA_SAFE_CALL(cudaMalloc((void**) &(dd_fftaps),sizeof(float)*n));
+
GR_CUDA_SAFE_CALL(cudaMemcpy(dd_fftaps,(float*)&(fftaps[0]),sizeof(float)*n,cudaMemcpyHostToDevice));
+ if(dd_fbtaps) GR_CUDA_SAFE_CALL(cudaFree(dd_fbtaps));
+ GR_CUDA_SAFE_CALL(cudaMalloc((void**) &(dd_fbtaps),sizeof(float)*m));
+
GR_CUDA_SAFE_CALL(cudaMemcpy(dd_fbtaps,(float*)&(fbtaps[0]),sizeof(float)*m,cudaMemcpyHostToDevice));
+ if(dd_prev_input) GR_CUDA_SAFE_CALL(cudaFree(dd_prev_input));
+ GR_CUDA_SAFE_CALL(cudaMalloc((void**) &(dd_prev_input),sizeof(float)*2*n));
+ if(dd_prev_output) GR_CUDA_SAFE_CALL(cudaFree(dd_prev_output));
+ GR_CUDA_SAFE_CALL(cudaMalloc((void**)
&(dd_prev_output),sizeof(float)*2*m));
+ GR_CUDA_SAFE_CALL(cudaMemset(dd_prev_input,0,sizeof(float)*2*n));
+ GR_CUDA_SAFE_CALL(cudaMemset(dd_prev_output,0,sizeof(float)*2*m));
+ }
+
+int
+cuda_iir_filter2_fff::work (int noutput_items,
+ gr_vector_const_void_star &input_items,
+ gr_vector_void_star &output_items)
+{
+ const float *iptr = (const float *) input_items[0];
+ float *optr = (float *) output_items[0];
+
+
+ if (d_updated){
+ //d_iir.set_taps (d_new_fftaps, d_new_fbtaps);
+ set_kernel_taps (d_new_fftaps, d_new_fbtaps);
+ d_updated = false;
+ }
+
+ //d_iir.filter_n (out, in, noutput_items);
+ cudai_iir_fff_filter2_n_kernel_params kernel_params;
+ cudai_get_iir_fff_filter2_n_kernel_params(&kernel_params,noutput_items);
+
+
+ cudai_iir_fff_filter2_n(optr,iptr, (const unsigned int) noutput_items,
&kernel_params,
+ dd_m, dd_n, dd_latest_m,
dd_latest_n,
+ dd_fftaps, dd_fbtaps,
dd_prev_input, dd_prev_output);
+
+ return noutput_items;
+};
Added:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda_iir_filter2_fff.h
===================================================================
---
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda_iir_filter2_fff.h
(rev 0)
+++
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda_iir_filter2_fff.h
2008-10-22 14:35:54 UTC (rev 9818)
@@ -0,0 +1,126 @@
+/*
+ * Copyright 2004,2008 Free Software Foundation, Inc.
+ *
+ * This file is part of GNU Radio
+ *
+ * GNU Radio is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU General Public License as published by
+ * the Free Software Foundation; either version 3, or (at your option)
+ * any later version.
+ *
+ * GNU Radio is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with GNU Radio; see the file COPYING. If not, write to
+ * the Free Software Foundation, Inc., 51 Franklin Street,
+ * Boston, MA 02110-1301, USA.
+ */
+
+#ifndef INCLUDED_GR_IIR_FILTER2_FFF_H
+#define INCLUDED_GR_IIR_FILTER2_FFF_H
+
+#include <gr_sync_block.h>
+#include <gri_iir.h>
+#include <stdexcept>
+
+class cuda_iir_filter2_fff;
+typedef boost::shared_ptr<cuda_iir_filter2_fff> cuda_iir_filter2_fff_sptr;
+cuda_iir_filter2_fff_sptr
+cuda_make_iir_filter2_fff (const std::vector<float> &fftaps,
+ const std::vector<float> &fbtaps) throw
(std::invalid_argument);
+
+/*!
+ * \brief cuda IIR filter with float input, float output and float taps
+ * \ingroup filter
+ *
+ * This filter uses the Direct Form I implementation, where
+ * \p fftaps contains the feed-forward taps, and \p fbtaps the feedback ones.
+ *
+ *
+ * The input and output satisfy a difference equation of the form
+
+ \f[
+ y[n] - \sum_{k=1}^{M} a_k y[n-k] = \sum_{k=0}^{N} b_k x[n-k]
+ \f]
+
+ * with the corresponding rational system function
+
+ \f[
+ H(z) = \frac{\sum_{k=0}^{M} b_k z^{-k}}{1 - \sum_{k=1}^{N} a_k z^{-k}}
+ \f]
+
+ * Note that some texts define the system function with a + in the denominator.
+ * If you're using that convention, you'll need to negate the feedback taps.
+ */
+class cuda_iir_filter2_fff : public gr_sync_block
+{
+ private:
+ friend cuda_iir_filter2_fff_sptr
+ cuda_make_iir_filter2_fff (const std::vector<float> &fftaps,
+ const std::vector<float> &fbtaps) throw
(std::invalid_argument);
+
+ gri_iir<float,float,float> d_iir;
+ std::vector<float> d_new_fftaps;
+ std::vector<float> d_new_fbtaps;
+ bool d_updated;
+
+ //following members live in cuda device memory
+ unsigned int *dd_m;
+ unsigned int *dd_n;
+ int *dd_latest_m;
+ int * dd_latest_n;
+ float * dd_fftaps;
+ float * dd_fbtaps;
+ float * dd_prev_input;
+ float * dd_prev_output;
+
+ /*!
+ * Construct an IIR filter with the given taps
+ */
+ cuda_iir_filter2_fff (const std::vector<float> &fftaps,
+ const std::vector<float> &fbtaps) throw
(std::invalid_argument);
+
+ public:
+ ~cuda_iir_filter2_fff ();
+ /*!
+ * \brief Called to enable drivers, etc for i/o devices.
+ *
+ * This allows a block to enable an associated driver to begin
+ * transfering data just before we start to execute the scheduler.
+ * The end result is that this reduces latency in the pipeline when
+ * dealing with audio devices, usrps, etc.
+ *
+ * In this case initialize the cuda memory for the iir taps
+ * In cuda all device memory has to be created in the same thread
+ * as where it is used. So we have to initialize and destroy it with
+ * every start and stop.
+ */
+ virtual bool start();
+
+ /*!
+ * \brief Called to disable drivers, etc for i/o devices.
+ * In this case destroy the cuda memory for the iir taps
+ */
+ virtual bool stop();
+ /*!
+ * \brief set new taps.
+ */
+
+ void set_taps (const std::vector<float> &fftaps,
+ const std::vector<float> &fbtaps) throw
(std::invalid_argument);
+
+ /*!
+ * \brief copy new taps to cuda device.
+ */
+ void set_kernel_taps (const std::vector<float> &fftaps,
+ const std::vector<float> &fbtaps) throw
(std::invalid_argument);
+
+ int work (int noutput_items,
+ gr_vector_const_void_star &input_items,
+ gr_vector_void_star &output_items);
+};
+
+#endif
Added:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda_iir_filter2_fff.i
===================================================================
---
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda_iir_filter2_fff.i
(rev 0)
+++
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda_iir_filter2_fff.i
2008-10-22 14:35:54 UTC (rev 9818)
@@ -0,0 +1,40 @@
+/* -*- c++ -*- */
+/*
+ * Copyright 2004,2008 Free Software Foundation, Inc.
+ *
+ * This file is part of GNU Radio
+ *
+ * GNU Radio is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU General Public License as published by
+ * the Free Software Foundation; either version 3, or (at your option)
+ * any later version.
+ *
+ * GNU Radio is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with GNU Radio; see the file COPYING. If not, write to
+ * the Free Software Foundation, Inc., 51 Franklin Street,
+ * Boston, MA 02110-1301, USA.
+ */
+
+GR_SWIG_BLOCK_MAGIC(cuda,iir_filter2_fff);
+
+cuda_iir_filter2_fff_sptr
+cuda_make_iir_filter2_fff (const std::vector<float> &fftaps,
+ const std::vector<float> &fbtaps) throw
(std::invalid_argument);
+
+class cuda_iir_filter2_fff : public gr_sync_block
+{
+ private:
+ cuda_iir_filter2_fff (const std::vector<float> &fftaps,
+ const std::vector<float> &fbtaps) throw
(std::invalid_argument);
+
+ public:
+ ~cuda_iir_filter2_fff ();
+
+ void set_taps (const std::vector<float> &fftaps,
+ const std::vector<float> &fbtaps) throw
(std::invalid_argument);
+};
Added:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda_iir_filter_fff.cc
===================================================================
---
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda_iir_filter_fff.cc
(rev 0)
+++
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda_iir_filter_fff.cc
2008-10-22 14:35:54 UTC (rev 9818)
@@ -0,0 +1,95 @@
+/* -*- c++ -*- */
+/*
+ * Copyright 2004,2008 Free Software Foundation, Inc.
+ *
+ * This file is part of GNU Radio
+ *
+ * GNU Radio is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU General Public License as published by
+ * the Free Software Foundation; either version 3, or (at your option)
+ * any later version.
+ *
+ * GNU Radio is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with GNU Radio; see the file COPYING. If not, write to
+ * the Free Software Foundation, Inc., 51 Franklin Street,
+ * Boston, MA 02110-1301, USA.
+ */
+
+#ifdef HAVE_CONFIG_H
+#include "config.h"
+#endif
+
+#include <cuda_iir_filter_fff.h>
+#include <gr_io_signature.h>
+#include <stdio.h>
+
+
+cuda_iir_filter_fff_sptr
+cuda_make_iir_filter_fff (const std::vector<float> &fftaps,
+ const std::vector<float> &fbtaps) throw
(std::invalid_argument)
+{
+ fprintf(stderr,"cuda_make_iir_filter_fff (..)\n");
+ return cuda_iir_filter_fff_sptr (new cuda_iir_filter_fff (fftaps, fbtaps));
+}
+
+cuda_iir_filter_fff::cuda_iir_filter_fff (const std::vector<float> &fftaps,
+ const std::vector<float> &fbtaps) throw
(std::invalid_argument)
+
+ : gr_sync_block ("iir_filter_fff",
+ gr_make_io_signature (1, 1, sizeof (float),GR_BUFFER_CUDA),
+ gr_make_io_signature (1, 1, sizeof (float),GR_BUFFER_CUDA)),
+// d_iir (fftaps, fbtaps),
+ d_new_fftaps(fftaps),
+ d_new_fbtaps(fbtaps),
+ d_updated (true),
+ d_do_cuda_init(true)
+{
+ fprintf(stderr,"cuda_iir_filter_fff::cuda_iir_filter_fff (..)\n");
+ d_iir=NULL;
+ // fprintf (stderr, "cuda_iir_filter_fff::ctor\n");
+}
+
+cuda_iir_filter_fff::~cuda_iir_filter_fff ()
+{
+ fprintf(stderr,"cuda_iir_filter_fff::~cuda_iir_filter_fff ()\n");
+ delete d_iir;
+}
+
+void
+cuda_iir_filter_fff::set_taps (const std::vector<float> &fftaps,
+ const std::vector<float> &fbtaps) throw
(std::invalid_argument)
+{
+ fprintf(stderr,"cuda_iir_filter_fff::set_taps (..)\n");
+ d_new_fftaps = fftaps;
+ d_new_fbtaps = fbtaps;
+ d_updated = true;
+}
+
+int
+cuda_iir_filter_fff::work (int noutput_items,
+ gr_vector_const_void_star &input_items,
+ gr_vector_void_star &output_items)
+{
+ const float *in = (const float *) input_items[0];
+ float *out = (float *) output_items[0];
+ fprintf(stderr,"cuda_iir_filter_fff::work(..)
noutput_items=%i\n",noutput_items);
+
+ if ((d_do_cuda_init) || (NULL==d_iir))
+ {
+ d_iir = new cudai_iir_fff (d_new_fftaps, d_new_fbtaps);
+ d_do_cuda_init=false;
+ d_updated = false;
+ } else if (d_updated)
+ {
+ d_iir->set_taps (d_new_fftaps, d_new_fbtaps);
+ d_updated = false;
+ }
+
+ d_iir->filter_n (out, in, noutput_items);
+ return noutput_items;
+};
Added:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda_iir_filter_fff.h
===================================================================
---
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda_iir_filter_fff.h
(rev 0)
+++
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda_iir_filter_fff.h
2008-10-22 14:35:54 UTC (rev 9818)
@@ -0,0 +1,89 @@
+/* -*- c++ -*- */
+/*
+ * Copyright 2004,2008 Free Software Foundation, Inc.
+ *
+ * This file is part of GNU Radio
+ *
+ * GNU Radio is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU General Public License as published by
+ * the Free Software Foundation; either version 3, or (at your option)
+ * any later version.
+ *
+ * GNU Radio is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with GNU Radio; see the file COPYING. If not, write to
+ * the Free Software Foundation, Inc., 51 Franklin Street,
+ * Boston, MA 02110-1301, USA.
+ */
+
+#ifndef INCLUDED_GR_IIR_FILTER_FFF_H
+#define INCLUDED_GR_IIR_FILTER_FFF_H
+
+#include <gr_sync_block.h>
+#include <cudai_iir.h>
+#include <stdexcept>
+
+class cuda_iir_filter_fff;
+typedef boost::shared_ptr<cuda_iir_filter_fff> cuda_iir_filter_fff_sptr;
+cuda_iir_filter_fff_sptr
+cuda_make_iir_filter_fff (const std::vector<float> &fftaps,
+ const std::vector<float> &fbtaps) throw
(std::invalid_argument);
+
+/*!
+ * \brief IIR filter with float input, float output and float taps
+ * \ingroup filter
+ *
+ * This filter uses the Direct Form I implementation, where
+ * \p fftaps contains the feed-forward taps, and \p fbtaps the feedback ones.
+ *
+ *
+ * The input and output satisfy a difference equation of the form
+
+ \f[
+ y[n] - \sum_{k=1}^{M} a_k y[n-k] = \sum_{k=0}^{N} b_k x[n-k]
+ \f]
+
+ * with the corresponding rational system function
+
+ \f[
+ H(z) = \frac{\sum_{k=0}^{M} b_k z^{-k}}{1 - \sum_{k=1}^{N} a_k z^{-k}}
+ \f]
+
+ * Note that some texts define the system function with a + in the denominator.
+ * If you're using that convention, you'll need to negate the feedback taps.
+ */
+class cuda_iir_filter_fff : public gr_sync_block
+{
+ private:
+ friend cuda_iir_filter_fff_sptr
+ cuda_make_iir_filter_fff (const std::vector<float> &fftaps,
+ const std::vector<float> &fbtaps) throw
(std::invalid_argument);
+
+ cudai_iir_fff *d_iir;
+ std::vector<float> d_new_fftaps;
+ std::vector<float> d_new_fbtaps;
+ bool d_updated;
+ bool d_do_cuda_init;
+
+ /*!
+ * Construct an IIR filter with the given taps
+ */
+ cuda_iir_filter_fff (const std::vector<float> &fftaps,
+ const std::vector<float> &fbtaps) throw
(std::invalid_argument);
+
+ public:
+ ~cuda_iir_filter_fff ();
+
+ void set_taps (const std::vector<float> &fftaps,
+ const std::vector<float> &fbtaps) throw
(std::invalid_argument);
+
+ int work (int noutput_items,
+ gr_vector_const_void_star &input_items,
+ gr_vector_void_star &output_items);
+};
+
+#endif
Added:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda_iir_filter_fff.i
===================================================================
---
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda_iir_filter_fff.i
(rev 0)
+++
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cuda_iir_filter_fff.i
2008-10-22 14:35:54 UTC (rev 9818)
@@ -0,0 +1,40 @@
+/* -*- c++ -*- */
+/*
+ * Copyright 2004,2008 Free Software Foundation, Inc.
+ *
+ * This file is part of GNU Radio
+ *
+ * GNU Radio is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU General Public License as published by
+ * the Free Software Foundation; either version 3, or (at your option)
+ * any later version.
+ *
+ * GNU Radio is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with GNU Radio; see the file COPYING. If not, write to
+ * the Free Software Foundation, Inc., 51 Franklin Street,
+ * Boston, MA 02110-1301, USA.
+ */
+
+GR_SWIG_BLOCK_MAGIC(cuda,iir_filter_fff);
+
+cuda_iir_filter_fff_sptr
+cuda_make_iir_filter_fff (const std::vector<float> &fftaps,
+ const std::vector<float> &fbtaps) throw
(std::invalid_argument);
+
+class cuda_iir_filter_fff : public gr_sync_block
+{
+ private:
+ cuda_iir_filter_fff (const std::vector<float> &fftaps,
+ const std::vector<float> &fbtaps) throw
(std::invalid_argument);
+
+ public:
+ ~cuda_iir_filter_fff ();
+
+ void set_taps (const std::vector<float> &fftaps,
+ const std::vector<float> &fbtaps) throw
(std::invalid_argument);
+};
Modified:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cudai_general_kernel.cu
===================================================================
---
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cudai_general_kernel.cu
2008-10-22 08:55:52 UTC (rev 9817)
+++
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cudai_general_kernel.cu
2008-10-22 14:35:54 UTC (rev 9818)
@@ -380,7 +380,7 @@
int
cudai_get_general_kernel_output_multiple(cudai_general_kernel_params *params)
{
- int output_multiple=params->num_outputs_padded;
+ int output_multiple=params->num_outputs_padded;//TODO we want a way bigger
minimal num_outputs //
return output_multiple;
}
Added:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cudai_iir.cu
===================================================================
---
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cudai_iir.cu
(rev 0)
+++
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cudai_iir.cu
2008-10-22 14:35:54 UTC (rev 9818)
@@ -0,0 +1,341 @@
+/* -*- c++ -*- */
+/*
+ * Copyright 2002,2008 Free Software Foundation, Inc.
+ *
+ * This file is part of GNU Radio
+ *
+ * GNU Radio is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU General Public License as published by
+ * the Free Software Foundation; either version 3, or (at your option)
+ * any later version.
+ *
+ * GNU Radio is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with GNU Radio; see the file COPYING. If not, write to
+ * the Free Software Foundation, Inc., 51 Franklin Street,
+ * Boston, MA 02110-1301, USA.
+ */
+
+#ifndef INCLUDED_CUDAI_IIR_CU
+#define INCLUDED_CUDAI_IIR_CU
+
+#include <vector>
+#include <stdexcept>
+
+#include <cutil.h>
+#include "gr_cuda.h"
+typedef float2 gr_complex;
+
+#include "cudai_general_kernel.h"
+#include "cudai_iir.h"
+
+#define LOCAL_CUDA_SYNC( message ) do { \
+ cudaError err = cudaThreadSynchronize(); \
+ if( cudaSuccess != err) { \
+ fprintf(stderr, "%s Cuda error in file '%s' in line %i : %s.\n", \
+ message,__FILE__, __LINE__, cudaGetErrorString( err) ); \
+ exit(EXIT_FAILURE); \
+ } } while (0)
+
+//typedef struct __align__(16){
+// unsigned int r, g, b, a;
+//} RGBA32;
+
+//typedef struct __align__(16){
+//__device__ int latest_n;
+//__device__ int latest_m;
+//__device__ std::vector<float> fftaps;
+//__device__ std::vector<float> fbtaps;
+//
+//__device__ std::vector<float> prev_output;
+//__device__ std::vector<float> prev_input;
+//} status_struct;
+
+#if 1
+__device__ int *dd_latest_n;//[1];
+__device__ int *dd_latest_m;
+__device__ float *dd_fftaps;
+__device__ float *dd_fbtaps;
+
+__device__ float *dd_prev_output;
+__device__ float *dd_prev_input;
+__device__ unsigned int *dd_n;
+__device__ unsigned int *dd_m;
+#else
+ int *dd_latest_n;//[1];
+ int *dd_latest_m;
+ float *dd_fftaps;
+ float *dd_fbtaps;
+
+ float *dd_prev_output;
+ float *dd_prev_input;
+ unsigned int *dd_n;
+ unsigned int *dd_m;
+#endif
+
+cudai_iir_fff::cudai_iir_fff (const std::vector<float>& fftaps,
+ const std::vector<float>& fbtaps) throw (std::invalid_argument)
+ {
+ fprintf(stderr,"cudai_iir_fff::cudai_iir_fff(fftaps,fbtaps)\n");
+ init_cuda();
+ set_taps (fftaps, fbtaps);
+
+ }
+
+cudai_iir_fff::cudai_iir_fff () : d_latest_n(0),d_latest_m(0) {
+ fprintf(stderr,"cudai_iir_fff::cudai_iir_fff()\n");
+ init_cuda();
+ }
+
+cudai_iir_fff::~cudai_iir_fff () {
+ fprintf(stderr,"cudai_iir_fff::~cudai_iir_fff ()\n");
+ free(d_params);
+ GR_CUDA_SAFE_CALL(cudaFree(dd_latest_n));
+ GR_CUDA_SAFE_CALL(cudaFree(dd_latest_m));
+ GR_CUDA_SAFE_CALL(cudaFree(dd_fftaps));
+ GR_CUDA_SAFE_CALL(cudaFree(dd_fbtaps));
+ GR_CUDA_SAFE_CALL(cudaFree(dd_prev_output));
+ GR_CUDA_SAFE_CALL(cudaFree(dd_prev_input));
+ GR_CUDA_SAFE_CALL(cudaFree(dd_n));
+ GR_CUDA_SAFE_CALL(cudaFree(dd_m));
+ }
+
+void
+cudai_iir_fff::init_cuda()
+{
+ fprintf(stderr,"cudai_iir_fff::init_cuda()\n");
+
d_params=(cudai_general_kernel_params*)malloc(sizeof(cudai_general_kernel_params));
+ cudai_get_general_kernel_params ( d_params,1);
+ GR_CUDA_SAFE_CALL(cudaMalloc((void**) &(dd_latest_n),sizeof(int)));
+ GR_CUDA_SAFE_CALL(cudaMalloc((void**) &(dd_latest_m),sizeof(int)));
+ GR_CUDA_SAFE_CALL(cudaMalloc((void**) &(dd_n),sizeof(unsigned int)));
+ GR_CUDA_SAFE_CALL(cudaMalloc((void**) &(dd_m),sizeof(unsigned int)));
+ dd_fftaps=NULL;
+ dd_fbtaps=NULL;
+ dd_prev_output=NULL;
+ dd_prev_input=NULL;
+}
+ /*!
+ * \brief install new taps.
+ */
+ void
+cudai_iir_fff::set_taps (const std::vector<float> &fftaps,
+ const std::vector<float> &fbtaps) throw (std::invalid_argument)
+ {
+ fprintf(stderr,"cudai_iir_fff::set_taps(..)\n");
+
+ d_latest_n = 0;
+ GR_CUDA_SAFE_CALL(cudaMemset(dd_latest_n,0,sizeof(int)));
+ d_latest_m = 0;
+ GR_CUDA_SAFE_CALL(cudaMemset(dd_latest_m,0,sizeof(int)));
+ d_fftaps = fftaps;
+ d_fbtaps = fbtaps;
+
+ int n = fftaps.size ();
+ GR_CUDA_SAFE_CALL(cudaMemcpy(dd_n,&n,sizeof(n),cudaMemcpyHostToDevice));
+ int m = fbtaps.size ();
+ GR_CUDA_SAFE_CALL(cudaMemcpy(dd_m,&m,sizeof(m),cudaMemcpyHostToDevice));
+ if(dd_fftaps) GR_CUDA_SAFE_CALL(cudaFree(dd_fftaps));
+ GR_CUDA_SAFE_CALL(cudaMalloc((void**) &(dd_fftaps),sizeof(fftaps[0])*n));
+
GR_CUDA_SAFE_CALL(cudaMemcpy(dd_fftaps,(float*)&fftaps[0],sizeof(fftaps[0])*n,cudaMemcpyHostToDevice));
+ if(dd_fbtaps) GR_CUDA_SAFE_CALL(cudaFree(dd_fbtaps));
+ GR_CUDA_SAFE_CALL(cudaMalloc((void**) &(dd_fbtaps),sizeof(fbtaps[0])*m));
+
GR_CUDA_SAFE_CALL(cudaMemcpy(dd_fbtaps,(float*)&fbtaps[0],sizeof(fbtaps[0])*m,cudaMemcpyHostToDevice));
+ d_prev_input.resize (2 * n);
+ if(dd_prev_input) GR_CUDA_SAFE_CALL(cudaFree(dd_prev_input));
+ GR_CUDA_SAFE_CALL(cudaMalloc((void**)
&(dd_prev_input),sizeof(d_prev_input[0])*2*n));
+ d_prev_output.resize (2 * m);
+ if(dd_prev_output) GR_CUDA_SAFE_CALL(cudaFree(dd_prev_output));
+ GR_CUDA_SAFE_CALL(cudaMalloc((void**)
&(dd_prev_output),sizeof(d_prev_output[0])*2*m));
+
+
+ for (int i = 0; i < 2 * n; i++){
+ d_prev_input[i] = 0;
+ }
+ GR_CUDA_SAFE_CALL(cudaMemset(dd_prev_input,0,sizeof(d_prev_input[0])*2*n));
+ for (int i = 0; i < 2 * m; i++){
+ d_prev_output[i] = 0;
+ }
+
GR_CUDA_SAFE_CALL(cudaMemset(dd_prev_output,0,sizeof(d_prev_output[0])*2*m));
+ }
+
+//
+// general case. We may want to specialize this
+//
+
+float
+cudai_iir_fff::filter (const float input)
+{
+ fprintf(stderr,"cudai_iir_fff::filter(..)\n");
+ float acc;
+ unsigned i = 0;
+ unsigned n = ntaps_ff ();
+ unsigned m = ntaps_fb ();
+
+ if (n == 0)
+ return (float) 0;
+
+ int latest_n = d_latest_n;
+ int latest_m = d_latest_m;
+
+ acc = d_fftaps[0] * input;
+ for (i = 1; i < n; i ++)
+ acc += (d_fftaps[i] * d_prev_input[latest_n + i]);
+ for (i = 1; i < m; i ++)
+ acc += (d_fbtaps[i] * d_prev_output[latest_m + i]);
+
+ // store the values twice to avoid having to handle wrap-around in the loop
+ d_prev_output[latest_m] = acc;
+ d_prev_output[latest_m+m] = acc;
+ d_prev_input[latest_n] = input;
+ d_prev_input[latest_n+n] = input;
+
+ latest_n--;
+ latest_m--;
+ if (latest_n < 0)
+ latest_n += n;
+ if (latest_m < 0)
+ latest_m += m;
+
+ d_latest_m = latest_m;
+ d_latest_n = latest_n;
+ return (float) acc;
+}
+
+
+//__device__ void
+//cudai_iir_fff_filter (float *output, const float *input, unsigned int n,
unsigned int m)
+
+__device__ float
+cudai_iir_fff_filter (const float input)
+{
+ float acc;
+ unsigned i = 0;
+ unsigned n = dd_n[0];//ntaps_ff ();
+ unsigned m = dd_m[0];//ntaps_fb ();
+
+ if (n == 0)
+ return (float) 0;
+
+ int latest_n = dd_latest_n[0];
+ int latest_m = dd_latest_m[0];
+
+ acc = dd_fftaps[0] * input;
+ for (i = 1; i < n; i ++)
+ acc += (dd_fftaps[i] * dd_prev_input[latest_n + i]);
+ for (i = 1; i < m; i ++)
+ acc += (dd_fbtaps[i] * dd_prev_output[latest_m + i]);
+
+ // store the values twice to avoid having to handle wrap-around in the loop
+ dd_prev_output[latest_m] = acc;
+ dd_prev_output[latest_m+m] = acc;
+ dd_prev_input[latest_n] = input;
+ dd_prev_input[latest_n+n] = input;
+
+ latest_n--;
+ latest_m--;
+ if (latest_n < 0)
+ latest_n += n;
+ if (latest_m < 0)
+ latest_m += m;
+
+ dd_latest_m[0] = latest_m;
+ dd_latest_n[0] = latest_n;
+ return (float) acc;
+}
+
+/*! \brief complex addition with constant
+ * \param g_idata_a first operand input data in global device memory
+ * \param konst second operand
+ * \param g_odata result data in global device memory
+ */
+__global__ void cudai_iir_fff_filter_n_kernel(float* g_odata, const float*
g_idata_a, const unsigned int size)
+{
+ //const unsigned int num_threads = blockDim.x * gridDim.x;
+ const unsigned int thread_id = blockIdx.x * blockDim.x + threadIdx.x;
+ //for (unsigned int i = thread_id; i < size; i += num_threads)
+ //{
+ // g_odata[i]=cudai_iir_fff_filter (g_idata_a[i]);
+ //}
+ //Since every output is dependant on the previous output we only use thread
0 (single-threaded)
+ //This hurts performance in a big way, but there is no way around.
+ if(0==thread_id)
+ for( unsigned int i=0;i<size;i++)
+ g_odata[i]=cudai_iir_fff_filter (g_idata_a[i]);
+}
+
+/*void
+cudai_iir_fff::filter_n (float output[],
+ const float input[],
+ long n)
+{
+ //for (int i = 0; i < n; i++)
+ // output[i] = filter (input[i]);
+ dim3 grid( 256, 1, 1);
+ dim3 threads( 256, 1, 1);
+ cudai_iir_fff_filter_n_kernel<<< grid, threads,0>>>(output,input,(unsigned
int)n);
+ CUT_CHECK_ERROR("cudai_iir_fff_filter_n_kernel");
+}*/
+
+int
+cudai_iir_fff::get_kernel_params ( unsigned int num_outputs )
+{
+
fprintf(stderr,"cudai_iir_fff::get_kernel_params(num_outputs=%i)\n",num_outputs);
+ //return cudai_get_general_kernel_params ( d_params, num_outputs );
+ d_params->dynamic_shared_mem_size=0;
+ dim3 griddim( 1, 1, 1);
+ dim3 threaddim( 1, 1, 1);
+ d_params->griddim=griddim;
+ d_params->threaddim=threaddim;
+ d_params->num_inputs_padded=num_outputs;
+ d_params->num_outputs_padded=num_outputs;
+ d_params->num_inputs=num_outputs;
+ d_params->num_outputs=num_outputs;
+ return 0;
+}
+
+int
+cudai_iir_fff::get_output_multiple()
+{
+ fprintf(stderr,"cudai_iir_fff::get_output_multiple()\n");
+ return 1;//(int) (d_params->num_outputs_padded);
+}
+
+void
+cudai_iir_fff::filter_n_device (float *device_output,
+ const float *device_input_a,
+ long long_n)
+{
+ fprintf(stderr,"cudai_iir_fff::filter_n_device(..)\n");
+ unsigned int n=(unsigned int)long_n;
+ //if(NULL==d_params)
+ //{
+ // cudai_get_general_kernel_params ( d_params,n);
+ //}
+ //cudaThreadSynchronize();
+ LOCAL_CUDA_SYNC( "cudai_iir_fff::filter_n" );
+ cudai_iir_fff_filter_n_kernel<<< d_params->griddim, d_params->threaddim,
d_params->dynamic_shared_mem_size>>>
+ (device_output,device_input_a, n);
+ char errortxt[1024];
+ sprintf(errortxt,"cudai_iir_fff_filter_n_kernel() execution
failed\ngriddim.x %i,%i,%i threaddim.x %i,%i,%i smemsize %i noutputspadded
%i\n",d_params->griddim.x,d_params->griddim.y,d_params->griddim.z,
d_params->threaddim.x,d_params->threaddim.y,d_params->threaddim.z,
d_params->dynamic_shared_mem_size, d_params->num_outputs_padded);
+ CUT_CHECK_ERROR(errortxt);
+}
+
+void
+cudai_iir_fff::filter_n (float *device_output,
+ const float *device_input,
+ long n)
+{
+ fprintf(stderr,"cudai_iir_fff::filter_n (..)\n");
+ filter_n_device (device_output,device_input,n);
+ //if((d_fftaps.size ()<=4) && (d_fbtaps.size ()<=4))
+ // filter_n_maxn4_maxm4 (device_output,device_input,n);
+ //else
+ // throw std::runtime_error ("cudai_iir_fff::filter_n(..) ERROR too many
taps\n");
+}
+#endif /* INCLUDED_CUDAI_IIR_CU */
+
Added:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cudai_iir.h
===================================================================
--- gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cudai_iir.h
(rev 0)
+++ gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cudai_iir.h
2008-10-22 14:35:54 UTC (rev 9818)
@@ -0,0 +1,108 @@
+/* -*- c++ -*- */
+/*
+ * Copyright 2008 Free Software Foundation, Inc.
+ *
+ * This file is part of GNU Radio
+ *
+ * GNU Radio is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU General Public License as published by
+ * the Free Software Foundation; either version 3, or (at your option)
+ * any later version.
+ *
+ * GNU Radio is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with GNU Radio; see the file COPYING. If not, write to
+ * the Free Software Foundation, Inc., 51 Franklin Street,
+ * Boston, MA 02110-1301, USA.
+ */
+
+#ifndef INCLUDED_CUDAI_IIR_H
+#define INCLUDED_CUDAI_IIR_H
+
+#include <vector>
+#include <stdexcept>
+#include <cuda_runtime.h>
+#include "cudai_general_kernel.h"
+/*!
+ * \brief Infinite Impulse Response filter (IIR), float in float out float taps
+ */
+
+ class cudai_iir_fff {
+public:
+ /*!
+ * \brief Construct an IIR with the given taps.
+ *
+ * This filter uses the Direct Form I implementation, where
+ * \p fftaps contains the feed-forward taps, and \p fbtaps the feedback ones.
+ *
+ * \p fftaps and \p fbtaps must have equal numbers of taps
+ *
+ * The input and output satisfy a difference equation of the form
+
+ \f[
+ y[n] - \sum_{k=1}^{M} a_k y[n-k] = \sum_{k=0}^{N} b_k x[n-k]
+ \f]
+
+ * with the corresponding rational system function
+
+ \f[
+ H(z) = \frac{\sum_{k=0}^{N} b_k z^{-k}}{1 - \sum_{k=1}^{M} a_k z^{-k}}
+ \f]
+
+ * Note that some texts define the system function with a + in the
denominator.
+ * If you're using that convention, you'll need to negate the feedback taps.
+ */
+ cudai_iir_fff (const std::vector<float>& fftaps,
+ const std::vector<float>& fbtaps) throw (std::invalid_argument);
+
+ cudai_iir_fff ();
+
+ ~cudai_iir_fff ();
+
+ /*!
+ * \brief compute a single output value.
+ * \returns the filtered input value.
+ */
+ float filter (const float input);
+
+ void filter_n_device (float *device_output,
+ const float *device_input_a,
+ long long_n);
+ /*!
+ * \brief compute an array of N output values.
+ * \p input must have N valid entries.
+ */
+ //void filter_n (float output[], const float input[], long n);
+ void filter_n (float *device_output, const float *device_input, long n);
+
+ /*!
+ * \return number of taps in filter.
+ */
+ unsigned ntaps_ff () const { return d_fftaps.size (); }
+ unsigned ntaps_fb () const { return d_fbtaps.size (); }
+
+ /*!
+ * \brief install new taps.
+ */
+ void set_taps (const std::vector<float> &fftaps,
+ const std::vector<float> &fbtaps) throw
(std::invalid_argument);
+
+protected:
+ void init_cuda();
+ int get_kernel_params ( unsigned int num_outputs );
+ int get_output_multiple();
+
+ std::vector<float> d_fftaps;
+ std::vector<float> d_fbtaps;
+ int d_latest_n;
+ int d_latest_m;
+ std::vector<float> d_prev_output;
+ std::vector<float> d_prev_input;
+ cudai_general_kernel_params *d_params;
+};
+
+#endif // #ifndef CUDAI_GENERAL_KERNEL_H
Added:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cudai_iir_kernel.cu
===================================================================
---
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cudai_iir_kernel.cu
(rev 0)
+++
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cudai_iir_kernel.cu
2008-10-22 14:35:54 UTC (rev 9818)
@@ -0,0 +1,180 @@
+/* -*- c++ -*- */
+/*
+ * Copyright 2008 Free Software Foundation, Inc.
+ *
+ * This file is part of GNU Radio
+ *
+ * GNU Radio is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU General Public License as published by
+ * the Free Software Foundation; either version 3, or (at your option)
+ * any later version.
+ *
+ * GNU Radio is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with GNU Radio; see the file COPYING. If not, write to
+ * the Free Software Foundation, Inc., 51 Franklin Street,
+ * Boston, MA 02110-1301, USA.
+ */
+#ifndef CUDAI_IIR_FFF2_KERNEL_CU
+#define CUDAI_IIR_FFF2_KERNEL_CU
+
+/////
+// system includes
+#include <stdlib.h>
+#include <stdio.h>
+#include <string.h>
+#include <math.h>
+
+#include <cutil.h>
+#include "gr_cuda.h"
+
+#include "cudai_iir_kernel.h"
+
+#define LOCAL_CUDA_SYNC( message ) do { \
+ cudaError err = cudaThreadSynchronize(); \
+ if( cudaSuccess != err) { \
+ fprintf(stderr, "%s Cuda error in file '%s' in line %i : %s.\n", \
+ message,__FILE__, __LINE__, cudaGetErrorString( err) ); \
+ exit(EXIT_FAILURE); \
+ } } while (0)
+
+//define SDATA( index) CUT_BANK_CHECKER(f_shared, index)
+
+
+/*! \brief cuda iir filter float input, float output float taps
+ * \param g_idata_a first operand input data in global device memory
+ * \param g_odata result data in global device memory
+ */
+__global__ void cudai_iir_fff_filter2_n_kernel(float* g_odata, const float*
g_idata_a, const unsigned int size,
+ unsigned int *dd_m, unsigned int
*dd_n, int *dd_latest_m, int * dd_latest_n,
+ float * dd_fftaps, float *
dd_fbtaps, float * dd_prev_input, float * dd_prev_output)
+{
+ //const unsigned int num_threads = blockDim.x * gridDim.x;
+ const unsigned int thread_id = blockIdx.x * blockDim.x + threadIdx.x;
+ //for (unsigned int i = thread_id; i < size; i += num_threads)
+ //{
+ // g_odata[i]=cudai_iir_fff_filter (g_idata_a[i]);
+ //}
+ //Since every output is dependant on the previous output we only use thread
0 (single-threaded)
+ //This hurts performance in a big way, but there is no way around.
+ if(0==thread_id)
+ {
+ unsigned n = dd_n[0];//ntaps_ff ();
+ unsigned m = dd_m[0];//ntaps_fb ();
+
+ if (n == 0)
+ return;// (float) 0;
+
+ int latest_n = dd_latest_n[0];
+ int latest_m = dd_latest_m[0];
+ for( unsigned int j=0;j<size;j++)
+ //g_odata[i]=cudai_iir_fff_filter (g_idata_a[i]);
+ {
+ float acc;
+ unsigned i = 0;
+
+ float input=g_idata_a[j];//TODO use all threads to cache input in
shared mem
+ acc = dd_fftaps[0] * input;
+ for (i = 1; i < n; i ++)
+ acc += (dd_fftaps[i] * dd_prev_input[latest_n + i]);
+ for (i = 1; i < m; i ++)
+ acc += (dd_fbtaps[i] * dd_prev_output[latest_m + i]);
+
+ // store the values twice to avoid having to handle wrap-around in the
loop
+ dd_prev_output[latest_m] = acc;
+ dd_prev_output[latest_m+m] = acc;
+ dd_prev_input[latest_n] = input;
+ dd_prev_input[latest_n+n] = input;
+
+ latest_n--;
+ latest_m--;
+ if (latest_n < 0)
+ latest_n += n;
+ if (latest_m < 0)
+ latest_m += m;
+
+
+ //return (float) acc;
+ g_odata[j]=acc;
+ }
+ dd_latest_m[0] = latest_m;
+ dd_latest_n[0] = latest_n;
+ }
+
+}
+
+
+void
+cudai_iir_fff_filter2_n(float* device_output, const float* device_input_a,
const unsigned int n, cudai_iir_fff_filter2_n_kernel_params *params,
+ unsigned int *dd_m, unsigned int
*dd_n, int *dd_latest_m, int * dd_latest_n,
+ float * dd_fftaps, float *
dd_fbtaps, float * dd_prev_input, float * dd_prev_output)
+{
+ cudai_iir_fff_filter2_n_kernel_params tmp_params;
+ if(NULL==params)
+ {
+ params=&tmp_params;
+ cudai_get_iir_fff_filter2_n_kernel_params ( params, n);
+ }
+ cudaThreadSynchronize();//TODO remove this
+ LOCAL_CUDA_SYNC( "cudai_iir_fff_filter2_n" );
+ cudai_iir_fff_filter2_n_kernel<<< params->griddim, params->threaddim,
params->dynamic_shared_mem_size>>>
+ (device_output,device_input_a, n,
+
dd_m,dd_n,dd_latest_m,dd_latest_n,dd_fftaps,dd_fbtaps,dd_prev_input,dd_prev_output);
+ char errortxt[1024];
+ sprintf(errortxt,"cudai_iir_fff_filter2_n_kernel() execution
failed\ngriddim.x %i,%i,%i threaddim.x %i,%i,%i smemsize %i noutputspadded %i
n %i\n",params->griddim.x,params->griddim.y,params->griddim.z,
params->threaddim.x,params->threaddim.y,params->threaddim.z,
params->dynamic_shared_mem_size, params->num_outputs_padded, n);
+ CUT_CHECK_ERROR(errortxt);
+}
+
+
+int
+cudai_get_iir_fff_filter2_n_kernel_params_fixed (
cudai_iir_fff_filter2_n_kernel_params *params, unsigned int num_outputs )
+{
+ int result=0;
+ //const unsigned int max_num_threads_per_block = MAX_NUM_THREADS_ALL;
//can use the maximum number of threads if wanted
+ //unsigned int max_num_blocks = MAX_NUM_BLOCKS_ALL;
+
+ unsigned int num_blocks=1 ;// = gridDim.x;
//NUM_CUDABLOCKS
+ unsigned int num_threads_per_block=1;// = blockDim.x;
//NUM_THREADS;
+ unsigned int num_outputs_per_block=num_threads_per_block;
+
+ const unsigned int num_outputs_per_grid= num_outputs_per_block*num_blocks;
//(blockDim.x)*gridDim.x
+
+ size_t dynamic_shared_mem_size =
0;//256*sizeof(float);//0;//num_threads_per_block*sizeof(gr_complex);
+ dim3 griddim( num_blocks, 1, 1);
+ dim3 threaddim( num_threads_per_block, 1, 1);
+
+ params->griddim=griddim;
+ params->threaddim=threaddim;
+ params->dynamic_shared_mem_size=dynamic_shared_mem_size;
+ params->num_outputs_padded=num_outputs_per_grid;
+ params->num_inputs_padded=num_outputs_per_grid;
+ params->num_inputs=0;//num_outputs_per_grid;//num_outputs;
+ params->num_outputs=0;//num_outputs_per_grid;//num_outputs;
+
+ //Now you can do the kernel invocation like this:
+ //cudai_XXXX_kernel<<< params->griddim, params->threaddim,
params->dynamic_shared_mem_size >>>(g_odata, g_idata,....,
params->num_outputs_padded*X);
+ return result;
+}
+
+
+
+int
+cudai_get_iir_fff_filter2_n_kernel_params (
cudai_iir_fff_filter2_n_kernel_params *params, unsigned int num_outputs )
+{
+ return cudai_get_iir_fff_filter2_n_kernel_params_fixed ( params, num_outputs
);
+}
+
+int
+cudai_get_iir_fff_filter2_n_output_multiple(cudai_iir_fff_filter2_n_kernel_params
*params)
+{
+ int output_multiple=params->num_outputs_padded;//TODO we want a way bigger
minimal num_outputs //params->num_outputs_padded;
+ return output_multiple;
+}
+
+
+
+#endif // #ifndef CUDAI_IIR_FFF2_KERNEL_CU
Added:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cudai_iir_kernel.h
===================================================================
---
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cudai_iir_kernel.h
(rev 0)
+++
gnuradio/branches/developers/nldudok1/gpgpu-wip/gr-cuda/src/lib/cudai_iir_kernel.h
2008-10-22 14:35:54 UTC (rev 9818)
@@ -0,0 +1,45 @@
+#ifndef CUDAI_IIR_FFF2_KERNEL_H
+#define CUDAI_IIR_FFF2_KERNEL_H
+
+#include "cuda_kernel_common.h"
+
+extern "C" {
+
+struct cudai_iir_fff_filter2_n_kernel_params{
+ size_t dynamic_shared_mem_size;
+ dim3 griddim;
+ dim3 threaddim;
+ unsigned num_inputs_padded;
+ unsigned num_outputs_padded;
+ unsigned num_inputs;
+ unsigned num_outputs;
+} ;
+
+void
+cudai_iir_fff_filter2_n(float* device_output, const float* device_input_a,
const unsigned int n, cudai_iir_fff_filter2_n_kernel_params *params,
+ unsigned int *dd_m, unsigned int
*dd_n, int *dd_latest_m, int * dd_latest_n,
+ float * dd_fftaps, float *
dd_fbtaps, float * dd_prev_input, float * dd_prev_output);
+
+/*! \brief get fixed kernel parameters which are large enough for optimal
calculation
+ * \param params pointer to already existing kernel_params structure
which will be filled
+ * \param num_outputs expected minimal number_of_outputs when kernel is
actually invoked
+ */
+int cudai_get_iir_fff_filter2_n_kernel_params_fixed (
cudai_iir_fff_filter2_n_kernel_params *params, unsigned int num_outputs );
+
+
+/*! \brief get kernel parameters
+ * \param params pointer to already existing kernel_params structure
which will be filled
+ * \param num_outputs expected minimal number_of_outputs when kernel is
actually invoked
+ * This will probably be a redirect to one of the specialized
get_general_kernel_params routines
+ */
+int cudai_get_iir_fff_filter2_n_kernel_params (
cudai_iir_fff_filter2_n_kernel_params *params, unsigned int num_outputs );
+
+/*! \brief the chunk size in which the kernel will calculate outputs
+ * \param params pointer to already filled kernel_params structure
+ * \returns output_multiple
+ * Make num_outputs a multiple of output_multiple when you invoke the kernel
for optimal performance
+ */
+int
cudai_get_iir_fff_filter2_n_output_multiple(cudai_iir_fff_filter2_n_kernel_params
*params);
+
+} //extern "C"
+#endif /*CUDAI_IIR_FFF2_KERNEL_H*/
Modified:
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/check_13.1.qdemod.py
===================================================================
---
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/check_13.1.qdemod.py
2008-10-22 08:55:52 UTC (rev 9817)
+++
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/check_13.1.qdemod.py
2008-10-22 14:35:54 UTC (rev 9818)
@@ -57,7 +57,7 @@
src=gr.noise_source_c ( gr.GR_UNIFORM,1.0, 1)
head = gr.head(itemsize_in,options.num_samples)
#dst=gr.null_sink(itemsize_out)
- check=gr.check_compare_ff(1.0e-6,False,True)
+ check=gr.check_compare_ff(1.0e-5,False,True)
dst=gr.null_sink(gr.sizeof_float)
if True: #options.use_cuda:
#dst_cuda=cuda.null_sink(itemsize_out)
@@ -78,6 +78,8 @@
#testblock4= cuda.quadrature_demod_cuda_cf (1.0 )
conv_in_cuda=cuda.host_to_cuda(itemsize_in)
conv_out_cuda=cuda.cuda_to_host(itemsize_out)
+ out_file_cuda=gr.file_sink(itemsize_out, "out_cuda.raw")
+ self.out_vector_cuda=gr.vector_sink_f()
if True: #else:
#nop_gr=gr.nop(itemsize_in)
@@ -97,12 +99,19 @@
conv_in_gr=gr.kludge_copy(itemsize_in)
conv_out_gr=gr.kludge_copy(itemsize_out)
+ out_file_gr=gr.file_sink(itemsize_out, "out_gr.raw")
+ self.out_vector_gr=gr.vector_sink_f()
+
self.connect (src,head)
self.connect(head,conv_in_cuda,testblock1_cuda,conv_out_cuda)
self.connect(head,conv_in_gr,testblock1_gr,conv_out_gr)
- self.connect(conv_out_gr,test_gr,(check,0))
- self.connect(conv_out_cuda,test_cuda,(check,1))
+ self.connect(conv_out_gr,(check,0)) #test_gr
+ self.connect(conv_out_cuda,(check,1)) #test_cuda,
+ self.connect(conv_out_gr,out_file_gr) #test_gr
+ self.connect(conv_out_gr,self.out_vector_gr)
+ self.connect(conv_out_cuda,out_file_cuda) #test_cuda,
+ self.connect(conv_out_cuda,self.out_vector_cuda) #test_cuda,
self.connect(check,dst)
@@ -115,6 +124,10 @@
#print 'Blocked waiting for GDB attach (pid = %d)' % (os.getpid(),)
#raw_input ('Press Enter to continue: ')
# remainder of your test code follows...
- my_top_block().run()
+ tb=my_top_block()
+ tb.run()
+ print "out gr", tb.out_vector_gr.data()
+ print "out cuda", tb.out_vector_cuda.data()
+
except KeyboardInterrupt:
pass
Modified: gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/check_13.py
===================================================================
--- gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/check_13.py
2008-10-22 08:55:52 UTC (rev 9817)
+++ gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/check_13.py
2008-10-22 14:35:54 UTC (rev 9818)
@@ -84,12 +84,12 @@
1, 1, itemsize_out,
1,options.output_multiple,1.0,
True)
-#,
-# True,gr_consume_type_t cons_type, gr_produce_type_t prod_type
+ #,
+ # True,gr_consume_type_t cons_type, gr_produce_type_t prod_type
#nop_cuda = cuda.nop(itemsize_in)
if(options.fir):
print "fir"
- testblock1_cuda= cuda.fir_filter_ccc(options.decimation,taps)
+ testblock1_cuda= cuda.fir_filter_ccf(options.decimation,taps)
else:
print "fft"
testblock1_cuda= cuda.fft_filter_ccc(options.decimation,taps)
@@ -99,6 +99,7 @@
#testblock4= cuda.quadrature_demod_cuda_cf (1.0 )
conv_in_cuda=cuda.host_to_cuda(itemsize_in)
conv_out_cuda=cuda.cuda_to_host(itemsize_out)
+ self.out_vector_gr=gr.vector_sink_c()
if True: #else:
#nop_gr=gr.nop(itemsize_in)
@@ -108,7 +109,7 @@
1,options.output_multiple,1.0,
True)
if(options.fir):
- testblock1_gr= gr.fir_filter_ccc(options.decimation,taps)
+ testblock1_gr= gr.fir_filter_ccf(options.decimation,taps)
else:
testblock1_gr= gr.fft_filter_ccc(options.decimation,taps)
@@ -120,6 +121,7 @@
conv_in_gr=gr.kludge_copy(itemsize_in)
conv_out_gr=gr.kludge_copy(itemsize_out)
+ self.out_vector_cuda=gr.vector_sink_c()
self.connect (src,head)
self.connect(head,conv_in_cuda,testblock1_cuda,conv_out_cuda)
@@ -128,9 +130,12 @@
self.connect(conv_out_cuda,(check,1))
self.connect(check,dst)
+ self.connect(conv_out_gr,self.out_vector_gr)
+ self.connect(conv_out_cuda,self.out_vector_cuda) #test_cuda,
+
if __name__ == '__main__':
try:
# insert this in your test code to debug with gdb ...
@@ -138,6 +143,10 @@
#print 'Blocked waiting for GDB attach (pid = %d)' % (os.getpid(),)
#raw_input ('Press Enter to continue: ')
# remainder of your test code follows...
- my_top_block().run()
+ #my_top_block().run()
+ tb=my_top_block()
+ tb.run()
+ print "out gr", tb.out_vector_gr.data()
+ print "out cuda", tb.out_vector_cuda.data()
except KeyboardInterrupt:
pass
Deleted: gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/core
Added: gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/run_13.2.iir.py
===================================================================
--- gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/run_13.2.iir.py
(rev 0)
+++ gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/run_13.2.iir.py
2008-10-22 14:35:54 UTC (rev 9818)
@@ -0,0 +1,158 @@
+#!/usr/bin/env python
+#
+# Copyright 2004,2005,2007 Free Software Foundation, Inc.
+#
+# This file is part of GNU Radio
+#
+# GNU Radio is free software; you can redistribute it and/or modify
+# it under the terms of the GNU General Public License as published by
+# the Free Software Foundation; either version 3, or (at your option)
+# any later version.
+#
+# GNU Radio is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+# GNU General Public License for more details.
+#
+# You should have received a copy of the GNU General Public License
+# along with GNU Radio; see the file COPYING. If not, write to
+# the Free Software Foundation, Inc., 51 Franklin Street,
+# Boston, MA 02110-1301, USA.
+#
+
+from gnuradio import gr
+from gnuradio import audio
+from gnuradio.eng_option import eng_option
+from optparse import OptionParser
+
+
+import math
+from gnuradio import cuda
+
+class my_top_block(gr.top_block):
+
+ def __init__(self):
+ gr.top_block.__init__(self)
+
+ parser = OptionParser(option_class=eng_option)
+
+ parser.add_option("-n", "--num-samples", type="int", default=1000000,
+ help="set number of samples to process (1000000)")
+ parser.add_option("-m", "--output-multiple", type="int", default=1,
+ help="set output_multiple (1)")
+
+ #parser.add_option ("-c", "--cuda", action="store_true",
dest="use_cuda",
+ # default=False, help="use cuda")
+ (options, args) = parser.parse_args ()
+ if len(args) != 0:
+ parser.print_help()
+ raise SystemExit, 1
+
+ #
+ # 1
+ # H(s) = -------
+ # 1 + s
+ #
+ # tau is the RC time constant.
+ # critical frequency: w_p = 1/tau
+ #
+ # We prewarp and use the bilinear z-transform to get our IIR
coefficients.
+ # See "Digital Signal Processing: A Practical Approach" by Ifeachor
and Jervis
+ #
+
+ """
+ FM Deemphasis IIR filter.
+ """
+
+ fs=320.0e3 #sampling frequency in Hz
+ tau=75.0e-6 # Time constant in seconds (75us in US, 50us in
EUR)
+ w_p = 1/tau
+ w_pp = math.tan (w_p / (fs * 2)) # prewarped analog freq
+
+ a1 = (w_pp - 1)/(w_pp + 1)
+ b0 = w_pp/(1 + w_pp)
+ b1 = b0
+
+ btaps = [b0, b1]
+ ataps = [1, a1]
+
+ if 0:
+ print "btaps =", btaps
+ print "ataps =", ataps
+ global plot1
+ plot1 = gru.gnuplot_freqz (gru.freqz (btaps, ataps), fs, True)
+
+ #deemph = gr.iir_filter_ffd(btaps, ataps)
+ #src=gr.null_source (gr.sizeof_float)
+ #src=gr.vector_source_f(range(1,32768,1),True)
+
+ #the block to test
+
+ itemsize_in=gr.sizeof_float
+ itemsize_out=gr.sizeof_float
+ src=gr.noise_source_f ( gr.GR_UNIFORM,1.0, 1)
+ head = gr.head(itemsize_in,options.num_samples)
+ #dst=gr.null_sink(itemsize_out)
+ check=gr.check_compare_ff(1.0e-6,False,True)
+ dst=gr.null_sink(gr.sizeof_float)
+ if True: #options.use_cuda:
+ #dst_cuda=cuda.null_sink(itemsize_out)
+ test_cuda= cuda.test("test",
+ 1, 1, itemsize_out,
+ 1, 1, itemsize_out,
+ 1,options.output_multiple,1.0,
+ True)
+ # ,
+ # True,gr_consume_type_t cons_type, gr_produce_type_t
prod_type
+ #nop_cuda = cuda.nop(itemsize_in)
+ print "cuda.iir_filter_fff"
+ testblock1_cuda= cuda.iir_filter2_fff(btaps, ataps)
+
+ #testblock4= cuda.multiply_const_cc(1.0)
+ #testblock5= cuda.multiply_const_cc(1.0)
+ #testblock6= cuda.multiply_const_cc(1.0)
+ #testblock4= cuda.quadrature_demod_cuda_cf (1.0 )
+ conv_in_cuda=cuda.host_to_cuda(itemsize_in)
+ conv_out_cuda=cuda.cuda_to_host(itemsize_out)
+
+ if False: #else:
+ #nop_gr=gr.nop(itemsize_in)
+ test_gr= gr.test("test",
+ 1, 1, itemsize_out,
+ 1, 1, itemsize_out,
+ 1,options.output_multiple,1.0,
+ True)
+ print "gr.iir_filter_ffd"
+ testblock1_gr= gr.iir_filter_ffd(btaps, ataps)
+
+ #testblock1_gr= gr.add_const_cc(complex(1,0))
+ #testblock4= gr.multiply_const_cc(1.0)
+ #testblock5= gr.multiply_const_cc(1.0)
+ #testblock6= gr.multiply_const_cc(1.0)
+ #testblock4= gr.quadrature_demod_cf(1.0)
+
+ conv_in_gr=gr.kludge_copy(itemsize_in)
+ conv_out_gr=gr.kludge_copy(itemsize_out)
+
+ self.connect (src,head)
+ self.connect(head,conv_in_cuda,testblock1_cuda,conv_out_cuda)
+ #self.connect(head,conv_in_gr,testblock1_gr,conv_out_gr)
+ #self.connect(conv_out_gr,test_gr,(check,0))
+ #self.connect(conv_out_cuda,test_cuda,(check,1))
+ #self.connect(check,dst)
+ self.connect(conv_out_cuda,dst)
+
+
+
+
+if __name__ == '__main__':
+ try:
+ # insert this in your test code to debug with gdb ...
+ #import os
+ #print 'Blocked waiting for GDB attach (pid = %d)' % (os.getpid(),)
+ #raw_input ('Press Enter to continue: ')
+ # remainder of your test code follows...
+ my_top_block().run()
+ except KeyboardInterrupt:
+ pass
+
Property changes on:
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/run_13.2.iir.py
___________________________________________________________________
Name: svn:executable
+ *
Modified:
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/wfm/check_wfm.py
===================================================================
--- gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/wfm/check_wfm.py
2008-10-22 08:55:52 UTC (rev 9817)
+++ gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/wfm/check_wfm.py
2008-10-22 14:35:54 UTC (rev 9818)
@@ -85,7 +85,7 @@
demod_rate=1.0
#dst=gr.null_sink(itemsize_out)
- check=gr.check_compare_ff(1.0e-6,False,True)
+ check=gr.check_compare_ff(1.0e-5,False,True)
dst=gr.null_sink(gr.sizeof_float)
if True: #options.use_cuda:
#dst_cuda=cuda.null_sink(itemsize_out)
@@ -94,10 +94,11 @@
# 1, 1, itemsize_out,
# 1,options.output_multiple,1.0,
# True)
-#,
-# True,gr_consume_type_t cons_type, gr_produce_type_t prod_type
+ #,
+ # True,gr_consume_type_t cons_type, gr_produce_type_t
prod_type
#nop_cuda = cuda.nop(itemsize_in)
testblock1_cuda= cuda_wfm_rcv.wfm_rcv (demod_rate,
options.decimation)
+ testblock2_cuda= cuda_wfm_rcv.wfm_rcv (demod_rate,
options.decimation)
#conv_in_cuda=cuda.host_to_cuda(itemsize_in)
#conv_out_cuda=cuda.cuda_to_host(itemsize_out)
@@ -109,12 +110,15 @@
1,options.output_multiple,1.0,
True)
testblock1_gr= blks2.wfm_rcv (demod_rate, options.decimation)
+ testblock2_gr= blks2.wfm_rcv (demod_rate, options.decimation)
#conv_in_gr=gr.kludge_copy(itemsize_in)
#conv_out_gr=gr.kludge_copy(itemsize_out)
self.connect(src,head)
self.connect(head,testblock1_cuda,(check,0))
- self.connect(head,testblock1_gr,(check,1))
+ #self.connect(head,testblock2_gr,(check,0))
+ #self.connect(head,testblock1_gr,(check,1))
+ self.connect(head,testblock2_cuda,(check,1))
self.connect(check,test_gr,dst)
Added: gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/wfm/check_wfm2.py
===================================================================
--- gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/wfm/check_wfm2.py
(rev 0)
+++ gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/wfm/check_wfm2.py
2008-10-22 14:35:54 UTC (rev 9818)
@@ -0,0 +1,136 @@
+#!/usr/bin/env python
+#
+# Copyright 2004,2005,2007 Free Software Foundation, Inc.
+#
+# This file is part of GNU Radio
+#
+# GNU Radio is free software; you can redistribute it and/or modify
+# it under the terms of the GNU General Public License as published by
+# the Free Software Foundation; either version 3, or (at your option)
+# any later version.
+#
+# GNU Radio is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+# GNU General Public License for more details.
+#
+# You should have received a copy of the GNU General Public License
+# along with GNU Radio; see the file COPYING. If not, write to
+# the Free Software Foundation, Inc., 51 Franklin Street,
+# Boston, MA 02110-1301, USA.
+#
+
+from gnuradio import gr, gru, eng_notation, optfir
+
+import cuda_wfm_rcv
+
+from gnuradio import audio
+from gnuradio import blks2
+from gnuradio.eng_option import eng_option
+from optparse import OptionParser
+import sys
+import math
+
+
+from gnuradio import cuda
+
+class my_top_block(gr.top_block):
+
+ def __init__(self):
+ gr.top_block.__init__(self)
+
+ parser = OptionParser(option_class=eng_option)
+
+ parser.add_option("-n", "--num-samples", type="int", default=1000000,
+ help="set number of samples to process (1000000)")
+ parser.add_option("-m", "--output-multiple", type="int", default=1,
+ help="set output_multiple (1)")
+ parser.add_option("-d", "--decimation", type="int", default=1,
+ help="set decimation (1)")
+ parser.add_option("-t", "--num_taps", type="int", default=64,
+ help="set number of taps (64)")
+ parser.add_option ("-f", "--fir", action="store_true",
+ default=False, help="use fir filter in stead of fft
filter")
+ #parser.add_option ("-c", "--cuda", action="store_true",
dest="use_cuda",
+ # default=False, help="use cuda")
+ (options, args) = parser.parse_args ()
+ if len(args) != 0:
+ parser.print_help()
+ raise SystemExit, 1
+
+
+ #src=gr.null_source (gr.sizeof_float)
+ #src=gr.vector_source_f(range(1,32768,1),True)
+
+ #the block to test
+ taps=range(1,options.num_taps+1,1)
+ tapsum=0.0
+ for i in range(0,options.num_taps,1):
+ tapsum=tapsum+1.0*taps[i]
+
+ tapavg=tapsum/(1.0*options.num_taps)
+ for i in range(0,options.num_taps,1):
+ taps[i]=(1.0*taps[i]-tapavg)/tapsum
+ print taps
+ #taps2=range(2,options.num_taps+2,1)
+ #decimation1=1
+ #decimation2=decimation1
+ #decimation3=decimation1
+ #decimation4=1
+ itemsize_in=gr.sizeof_gr_complex
+ itemsize_out=gr.sizeof_float
+ src=gr.noise_source_c ( gr.GR_UNIFORM,1.0, 1)
+ head = gr.head(itemsize_in,options.num_samples)
+
+ demod_rate=1.0
+
+ #dst=gr.null_sink(itemsize_out)
+ check=gr.check_compare_ff(1.0e-5,False,True)
+ dst=gr.null_sink(gr.sizeof_float)
+ if True: #options.use_cuda:
+ #dst_cuda=cuda.null_sink(itemsize_out)
+ #test_cuda= cuda.test("test",
+ # 1, 1, itemsize_in,
+ # 1, 1, itemsize_out,
+ # 1,options.output_multiple,1.0,
+ # True)
+ #,
+ # True,gr_consume_type_t cons_type, gr_produce_type_t
prod_type
+ #nop_cuda = cuda.nop(itemsize_in)
+ testblock1_cuda= cuda_wfm_rcv.wfm_rcv (demod_rate,
options.decimation)
+ testblock2_cuda= cuda_wfm_rcv.wfm_rcv (demod_rate,
options.decimation)
+ #conv_in_cuda=cuda.host_to_cuda(itemsize_in)
+ #conv_out_cuda=cuda.cuda_to_host(itemsize_out)
+
+ if True: #else:
+ #nop_gr=gr.nop(itemsize_in)
+ test_gr= gr.test("test",
+ 1, 1, itemsize_out,
+ 1, 1, itemsize_out,
+ 1,options.output_multiple,1.0,
+ True)
+ testblock1_gr= blks2.wfm_rcv (demod_rate, options.decimation)
+ testblock2_gr= blks2.wfm_rcv (demod_rate, options.decimation)
+ #conv_in_gr=gr.kludge_copy(itemsize_in)
+ #conv_out_gr=gr.kludge_copy(itemsize_out)
+
+ self.connect(src,head)
+ #self.connect(head,testblock1_cuda,(check,0))
+ self.connect(head,testblock2_gr,(check,0))
+ self.connect(head,testblock1_gr,(check,1))
+ #self.connect(head,testblock2_cuda,(check,1))
+ self.connect(check,test_gr,dst)
+
+
+
+
+if __name__ == '__main__':
+ try:
+ # insert this in your test code to debug with gdb ...
+ #import os
+ #print 'Blocked waiting for GDB attach (pid = %d)' % (os.getpid(),)
+ #raw_input ('Press Enter to continue: ')
+ # remainder of your test code follows...
+ my_top_block().run()
+ except KeyboardInterrupt:
+ pass
Property changes on:
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/wfm/check_wfm2.py
___________________________________________________________________
Name: svn:executable
+ *
Modified:
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/wfm/cuda_wfm_rcv.py
===================================================================
--- gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/wfm/cuda_wfm_rcv.py
2008-10-22 08:55:52 UTC (rev 9817)
+++ gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/wfm/cuda_wfm_rcv.py
2008-10-22 14:35:54 UTC (rev 9818)
@@ -43,9 +43,9 @@
volume = 20.
- max_dev = 75e3
- fm_demod_gain = quad_rate/(2*math.pi*max_dev)
- audio_rate = quad_rate / audio_decimation
+ max_dev = 75.0e3
+ fm_demod_gain = float(quad_rate)/(2*math.pi*max_dev)
+ audio_rate = float(quad_rate) / audio_decimation
# We assign to self so that outsiders can grab the demodulator
@@ -57,8 +57,25 @@
self.fm_demod = cuda.quadrature_demod_cuda_cf (fm_demod_gain)
# input: float; output: float
- self.deemph = fm_deemph (audio_rate)
+ #self.deemph = fm_deemph (audio_rate)
+
+ """
+ FM Deemphasis IIR filter.
+ """
+ fs=audio_rate #sampling frequency in Hz
+ tau=75.0e-6 # Time constant in seconds (75us in US, 50us in
EUR)
+ w_p = 1/tau
+ w_pp = math.tan (w_p / (fs * 2)) # prewarped analog freq
+
+ a1 = (w_pp - 1)/(w_pp + 1)
+ b0 = w_pp/(1 + w_pp)
+ b1 = b0
+
+ btaps = [b0, b1]
+ ataps = [1, a1]
+ self.deemph=cuda.iir_filter2_fff(btaps, ataps)
+
# compute FIR filter taps for audio filter
width_of_transition_band = audio_rate / 32
audio_coeffs = gr.firdes.low_pass (1.0, # gain
@@ -66,9 +83,10 @@
audio_rate/2 -
width_of_transition_band,
width_of_transition_band,
gr.firdes.WIN_HAMMING)
+ print "num_audio_coeffs",len(audio_coeffs)
# input: float; output: float
#self.audio_filter = gr.fir_filter_fff (audio_decimation, audio_coeffs)
self.audio_filter = cuda.fir_filter_fff (audio_decimation,
audio_coeffs)
self.c2h=cuda.cuda_to_host(gr.sizeof_float)
- self.connect (self, self.h2c,self.fm_demod, self.audio_filter,
self.c2h,self.deemph, self)
+ self.connect (self, self.h2c,self.fm_demod, self.audio_filter,
self.c2h,self) #self.deemph,
Modified:
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/wfm/cuda_wfm_rcv.pyc
===================================================================
(Binary files differ)
Added:
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/wfm/fm_emph_cuda.py
===================================================================
--- gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/wfm/fm_emph_cuda.py
(rev 0)
+++ gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/wfm/fm_emph_cuda.py
2008-10-22 14:35:54 UTC (rev 9818)
@@ -0,0 +1,151 @@
+#
+# Copyright 2005,2007 Free Software Foundation, Inc.
+#
+# This file is part of GNU Radio
+#
+# GNU Radio is free software; you can redistribute it and/or modify
+# it under the terms of the GNU General Public License as published by
+# the Free Software Foundation; either version 3, or (at your option)
+# any later version.
+#
+# GNU Radio is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+# GNU General Public License for more details.
+#
+# You should have received a copy of the GNU General Public License
+# along with GNU Radio; see the file COPYING. If not, write to
+# the Free Software Foundation, Inc., 51 Franklin Street,
+# Boston, MA 02110-1301, USA.
+#
+
+from gnuradio import gr
+import math
+from gnuradio import cuda
+
+#
+# 1
+# H(s) = -------
+# 1 + s
+#
+# tau is the RC time constant.
+# critical frequency: w_p = 1/tau
+#
+# We prewarp and use the bilinear z-transform to get our IIR coefficients.
+# See "Digital Signal Processing: A Practical Approach" by Ifeachor and Jervis
+#
+
+class fm_deemph(gr.hier_block2):
+ """
+ FM Deemphasis IIR filter.
+ """
+
+
+ def __init__(self, fs, tau=75e-6):
+ """
+ @param fs: sampling frequency in Hz
+ @type fs: float
+ @param tau: Time constant in seconds (75us in US, 50us in EUR)
+ @type tau: float
+ """
+ gr.hier_block2.__init__(self, "fm_deemph",
+ gr.io_signature(1, 1, gr.sizeof_float), # Input
signature
+ gr.io_signature(1, 1, gr.sizeof_float)) #
Output signature
+
+ w_p = 1/tau
+ w_pp = math.tan (w_p / (fs * 2)) # prewarped analog freq
+
+ a1 = (w_pp - 1)/(w_pp + 1)
+ b0 = w_pp/(1 + w_pp)
+ b1 = b0
+
+ btaps = [b0, b1]
+ ataps = [1, a1]
+
+ if 0:
+ print "btaps =", btaps
+ print "ataps =", ataps
+ global plot1
+ plot1 = gru.gnuplot_freqz (gru.freqz (btaps, ataps), fs, True)
+
+ deemph = gr.iir_filter_ffd(btaps, ataps)
+ self.connect(self, deemph, self)
+
+#
+# 1 + s*t1
+# H(s) = ----------
+# 1 + s*t2
+#
+# I think this is the right transfer function.
+#
+#
+# This fine ASCII rendition is based on Figure 5-15
+# in "Digital and Analog Communication Systems", Leon W. Couch II
+#
+#
+# R1
+# +-----||------+
+# | |
+# o------+ +-----+--------o
+# | C1 | |
+# +----/\/\/\/--+ \
+# /
+# \ R2
+# /
+# \
+# |
+# o--------------------------+--------o
+#
+# f1 = 1/(2*pi*t1) = 1/(2*pi*R1*C)
+#
+# 1 R1 + R2
+# f2 = ------- = ------------
+# 2*pi*t2 2*pi*R1*R2*C
+#
+# t1 is 75us in US, 50us in EUR
+# f2 should be higher than our audio bandwidth.
+#
+#
+# The Bode plot looks like this:
+#
+#
+# /----------------
+# /
+# / <-- slope = 20dB/decade
+# /
+# -------------/
+# f1 f2
+#
+# We prewarp and use the bilinear z-transform to get our IIR coefficients.
+# See "Digital Signal Processing: A Practical Approach" by Ifeachor and Jervis
+#
+
+class fm_preemph(gr.hier_block2):
+ """
+ FM Preemphasis IIR filter.
+ """
+ def __init__(self, fs, tau=75e-6):
+ """
+ @param fs: sampling frequency in Hz
+ @type fs: float
+ @param tau: Time constant in seconds (75us in US, 50us in EUR)
+ @type tau: float
+ """
+
+ gr.hier_block2.__init__(self, "fm_deemph",
+ gr.io_signature(1, 1, gr.sizeof_float), # Input
signature
+ gr.io_signature(1, 1, gr.sizeof_float)) #
Output signature
+
+ # FIXME make this compute the right answer
+
+ btaps = [1]
+ ataps = [1]
+
+ if 0:
+ print "btaps =", btaps
+ print "ataps =", ataps
+ global plot2
+ plot2 = gru.gnuplot_freqz (gru.freqz (btaps, ataps), fs, True)
+
+ preemph = gr.iir_filter_ffd(btaps, ataps)
+ self.connect(self, preemph, self)
| [Prev in Thread] |
Current Thread |
[Next in Thread] |
- [Commit-gnuradio] r9818 - in gnuradio/branches/developers/nldudok1/gpgpu-wip: gr-cuda/src/lib testbed testbed/wfm,
nldudok1 <=