[Top][All Lists]
[Date Prev][Date Next][Thread Prev][Thread Next][Date Index][Thread Index]
[Commit-gnuradio] r9522 - in gnuradio/branches/developers/nldudok1/gpgpu
|
From: |
nldudok1 |
|
Subject: |
[Commit-gnuradio] r9522 - in gnuradio/branches/developers/nldudok1/gpgpu-wip: . gnuradio-core/src/lib/runtime testbed |
|
Date: |
Sun, 7 Sep 2008 13:41:44 -0600 (MDT) |
Author: nldudok1
Date: 2008-09-07 13:41:42 -0600 (Sun, 07 Sep 2008)
New Revision: 9522
Added:
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/cuda_test_dial_tone.py
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/cuda_test_dial_tone_wav.py
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/cuda_test_vector.py
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/cuda_test_vector_int.py
Modified:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_buffer.cc
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_buffer.h
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_buffer.i
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_flat_flowgraph.cc
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_flat_flowgraph.h
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_io_signature.h
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_runtime_types.h
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_single_threaded_scheduler.cc
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_vmcircbuf.cc
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_vmcircbuf.h
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_vmcircbuf_cuda.cc
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_vmcircbuf_cuda.h
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/qa_gr_buffer.cc
Log:
allocate cuda buffers in work thread of block
Modified:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_buffer.cc
===================================================================
---
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_buffer.cc
2008-09-07 19:35:50 UTC (rev 9521)
+++
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_buffer.cc
2008-09-07 19:41:42 UTC (rev 9522)
@@ -32,6 +32,9 @@
#include <assert.h>
#include <algorithm>
+#include <gr_block.h>
+#include <gr_io_signature.h>
+
static long s_buffer_count = 0; // counts for debugging storage
mgmt
static long s_buffer_reader_count = 0;
@@ -77,21 +80,21 @@
}
-gr_buffer::gr_buffer (int nitems, size_t sizeof_item)
+gr_buffer::gr_buffer (int nitems, size_t sizeof_item, gr_block_sptr link)
: d_base (0), d_bufsize (0), d_vmcircbuf (0),
- d_sizeof_item (sizeof_item), d_write_index (0),
- d_done (false)
+ d_sizeof_item (sizeof_item), d_link(link),
+ d_write_index (0), d_done (false)
{
- if (!allocate_buffer (nitems, sizeof_item))
+ if (!allocate_buffer (nitems,
sizeof_item,link->output_signature()->buffer_type()))
throw std::bad_alloc ();
s_buffer_count++;
}
gr_buffer_sptr
-gr_make_buffer (int nitems, size_t sizeof_item)
+gr_make_buffer (int nitems, size_t sizeof_item, gr_block_sptr link)
{
- return gr_buffer_sptr (new gr_buffer (nitems, sizeof_item));
+ return gr_buffer_sptr (new gr_buffer (nitems, sizeof_item, link));
}
gr_buffer::~gr_buffer ()
@@ -101,18 +104,32 @@
s_buffer_count--;
}
+bool
+gr_buffer::start()
+{
+ bool retval = d_vmcircbuf->start();
+ d_base = (char *) d_vmcircbuf->pointer_to_first_copy ();
+ return retval;
+}
+
+bool
+gr_buffer::stop()
+{
+ bool retval = d_vmcircbuf->stop();
+ return retval;
+}
/*!
* sets d_vmcircbuf, d_base, d_bufsize.
* returns true iff successful.
*/
bool
-gr_buffer::allocate_buffer (int nitems, size_t sizeof_item)
+gr_buffer::allocate_buffer (int nitems, size_t sizeof_item,gr_buffer_type
buffer_type)
{
int orig_nitems = nitems;
// Any buffersize we come up with must be a multiple of min_nitems.
- int granularity = gr_vmcircbuf_sysconfig::granularity ();
+ int granularity = gr_vmcircbuf_sysconfig::granularity (buffer_type);
int min_nitems = minimum_buffer_items (sizeof_item, granularity);
// Round-up nitems to a multiple of min_nitems.
@@ -133,7 +150,7 @@
}
d_bufsize = nitems;
- d_vmcircbuf = gr_vmcircbuf_sysconfig::make (d_bufsize * d_sizeof_item);
+ d_vmcircbuf = gr_vmcircbuf_sysconfig::make (d_bufsize *
d_sizeof_item,buffer_type);
if (d_vmcircbuf == 0){
std::cerr << "gr_buffer::allocate_buffer: failed to allocate buffer of
size "
<< d_bufsize * d_sizeof_item / 1024 << " KB\n";
@@ -146,7 +163,7 @@
int
-gr_buffer::space_available () const
+gr_buffer::space_available ()
{
if (d_readers.empty ())
return d_bufsize - 1; // See comment below
@@ -175,19 +192,28 @@
void
gr_buffer::update_write_pointer (int nitems)
{
- d_vmcircbuf->update_circular(d_write_index,nitems);
+ d_vmcircbuf->update_circular(d_write_index* d_sizeof_item,nitems*
d_sizeof_item);
+ //scoped_lock guard(*mutex());//MDVH temporary disable mutex untill
full merge with mb-sched cude
d_write_index = index_add (d_write_index, nitems);
}
+void
+gr_buffer::set_done (bool done)
+{
+ //scoped_lock guard(*mutex());//MDVH temporary disable mutex untill
full merge with mb-sched cude
+ d_done = done;
+}
+
gr_buffer_reader_sptr
-gr_buffer_add_reader (gr_buffer_sptr buf, int nzero_preload)
+gr_buffer_add_reader (gr_buffer_sptr buf, int nzero_preload, gr_block_sptr
link)
{
if (nzero_preload < 0)
throw std::invalid_argument("gr_buffer_add_reader: nzero_preload must be
>= 0");
gr_buffer_reader_sptr r (new gr_buffer_reader (buf,
buf->index_sub(buf->d_write_index,
-
nzero_preload)));
+ nzero_preload),
+ link));
buf->d_readers.push_back (r.get ());
return r;
@@ -215,8 +241,9 @@
// ----------------------------------------------------------------------------
-gr_buffer_reader::gr_buffer_reader (gr_buffer_sptr buffer, unsigned int
read_index)
- : d_buffer (buffer), d_read_index (read_index)
+gr_buffer_reader::gr_buffer_reader(gr_buffer_sptr buffer, unsigned int
read_index,
+ gr_block_sptr link)
+ : d_buffer(buffer), d_read_index(read_index), d_link(link)
{
s_buffer_reader_count++;
}
@@ -242,6 +269,7 @@
void
gr_buffer_reader::update_read_pointer (int nitems)
{
+ //scoped_lock guard(*mutex());//MDVH temporary disable mutex untill
full merge with mb-sched cude
d_read_index = d_buffer->index_add (d_read_index, nitems);
}
Modified:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_buffer.h
===================================================================
---
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_buffer.h
2008-09-07 19:35:50 UTC (rev 9521)
+++
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_buffer.h
2008-09-07 19:41:42 UTC (rev 9522)
@@ -24,6 +24,8 @@
#define INCLUDED_GR_BUFFER_H
#include <gr_runtime_types.h>
+#include <boost/weak_ptr.hpp>
+//include <boost/thread.hpp> //MDVH temporary disabled untill full merge
with mb-sched code
class gr_vmcircbuf;
@@ -33,8 +35,12 @@
* The total size of the buffer will be rounded up to a system
* dependent boundary. This is typically the system page size, but
* under MS windows is 64KB.
+ *
+ * \param nitems is the minimum number of items the buffer will hold.
+ * \param sizeof_item is the size of an item in bytes.
+ * \param link is the buffer that writes to this buffer.
*/
-gr_buffer_sptr gr_make_buffer (int nitems, size_t sizeof_item);
+gr_buffer_sptr gr_make_buffer (int nitems, size_t sizeof_item, gr_block_sptr
link=gr_block_sptr());
/*!
@@ -43,14 +49,38 @@
*/
class gr_buffer {
public:
+
+ //typedef boost::unique_lock<boost::mutex> scoped_lock;//MDVH temporary
disable mutex untill full merge with mp-sched code
+
virtual ~gr_buffer ();
/*!
+ * \brief Called to enable drivers, etc for external memory.
+ *
+ * This allows a buffer to enable an associated driver and memory
+ * just before we start to execute the scheduler.
+ * The actual allocation and de-allocation of memory can also be delayed
+ * untill start and stop are called.
+ * This way the allocation will be done by the same thread that calls
+ * work on the blocks using this buffer.
+ */
+ virtual bool start();
+
+ /*!
+ * \brief Called to disable drivers, etc for external memory.
+ */
+ virtual bool stop();
+ /*!
* \brief return number of items worth of space available for writing
*/
- int space_available () const;
+ int space_available ();
/*!
+ * \brief return size of this buffer in items
+ */
+ int bufsize() const { return d_bufsize; }
+
+ /*!
* \brief return pointer to write buffer.
*
* The return value points at space that can hold at least
@@ -63,17 +93,26 @@
*/
void update_write_pointer (int nitems);
-
- void set_done (bool done) { d_done = done; }
+ void set_done (bool done);
bool done () const { return d_done; }
+ /*!
+ * \brief Return the block that writes this buffer.
+ */
+ gr_block_sptr link() { return gr_block_sptr(d_link); }
+
+ size_t nreaders() const { return d_readers.size(); }
+ gr_buffer_reader* reader(size_t index) { return d_readers[index]; }
+
+ //boost::mutex *mutex() { return &d_mutex; }//MDVH temporary disable mutex
untill full merge with mp-sched code
+
// -------------------------------------------------------------------------
private:
friend class gr_buffer_reader;
- friend gr_buffer_sptr gr_make_buffer (int nitems, size_t sizeof_item);
- friend gr_buffer_reader_sptr gr_buffer_add_reader (gr_buffer_sptr buf, int
nzero_preload);
+ friend gr_buffer_sptr gr_make_buffer (int nitems, size_t sizeof_item,
gr_block_sptr link);
+ friend gr_buffer_reader_sptr gr_buffer_add_reader (gr_buffer_sptr buf, int
nzero_preload, gr_block_sptr link);
protected:
char *d_base; // base address of
buffer
@@ -81,8 +120,14 @@
private:
gr_vmcircbuf *d_vmcircbuf;
size_t d_sizeof_item; // in bytes
+ std::vector<gr_buffer_reader *> d_readers;
+ boost::weak_ptr<gr_block> d_link; // block that writes
this buffer
+
+ //
+ // The mutex protects d_write_index, d_done and the d_read_index's in the
buffer readers.
+ //
+ //boost::mutex d_mutex;//MDVH temporary
disable mutex untill full merge with mp-sched code
unsigned int d_write_index; // in items
[0,d_bufsize)
- std::vector<gr_buffer_reader *> d_readers;
bool d_done;
unsigned
@@ -109,18 +154,22 @@
return s;
}
- virtual bool allocate_buffer (int nitems, size_t sizeof_item);
+ virtual bool allocate_buffer (int nitems, size_t sizeof_item, gr_buffer_type
buffer_type=GR_BUFFER_DEFAULT);
/*!
* \brief constructor is private. Use gr_make_buffer to create instances.
*
* Allocate a buffer that holds at least \p nitems of size \p sizeof_item.
*
+ * \param nitems is the minimum number of items the buffer will hold.
+ * \param sizeof_item is the size of an item in bytes.
+ * \param link is the buffer that writes to this buffer.
+ *
* The total size of the buffer will be rounded up to a system
* dependent boundary. This is typically the system page size, but
* under MS windows is 64KB.
*/
- gr_buffer (int nitems, size_t sizeof_item);
+ gr_buffer (int nitems, size_t sizeof_item, gr_block_sptr link);
/*!
* \brief disassociate \p reader from this buffer
@@ -132,8 +181,10 @@
/*!
* \brief create a new gr_buffer_reader and attach it to buffer \p buf
* \param nzero_preload -- number of zero items to "preload" into buffer.
+ * \param link is the buffer that reads using this gr_buffer_reader.
*/
-gr_buffer_reader_sptr gr_buffer_add_reader (gr_buffer_sptr buf, int
nzero_preload);
+gr_buffer_reader_sptr
+gr_buffer_add_reader (gr_buffer_sptr buf, int nzero_preload, gr_block_sptr
link=gr_block_sptr());
//! returns # of gr_buffers currently allocated
long gr_buffer_ncurrently_allocated ();
@@ -147,8 +198,10 @@
*/
class gr_buffer_reader {
+ public:
- public:
+ //typedef gr_buffer::scoped_lock scoped_lock;//MDVH temporary disable mutex
untill full merge with mp-sched code
+
~gr_buffer_reader ();
/*!
@@ -183,19 +236,29 @@
void set_done (bool done) { d_buffer->set_done (done); }
bool done () const { return d_buffer->done (); }
+ //boost::mutex *mutex() { return d_buffer->mutex(); }//MDVH temporary
disable mutex untill full merge with mp-sched code
+
+
+ /*!
+ * \brief Return the block that reads via this reader.
+ */
+ gr_block_sptr link() { return gr_block_sptr(d_link); }
+
// -------------------------------------------------------------------------
private:
friend class gr_buffer;
- friend gr_buffer_reader_sptr gr_buffer_add_reader (gr_buffer_sptr buf, int
nzero_preload);
+ friend gr_buffer_reader_sptr
+ gr_buffer_add_reader (gr_buffer_sptr buf, int nzero_preload, gr_block_sptr
link);
gr_buffer_sptr d_buffer;
unsigned int d_read_index; // in items
[0,d->buffer.d_bufsize)
+ boost::weak_ptr<gr_block> d_link; // block that reads via this
buffer reader
//! constructor is private. Use gr_buffer::add_reader to create instances
- gr_buffer_reader (gr_buffer_sptr buffer, unsigned int read_index);
+ gr_buffer_reader (gr_buffer_sptr buffer, unsigned int read_index,
gr_block_sptr link);
};
//! returns # of gr_buffer_readers currently allocated
Modified:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_buffer.i
===================================================================
---
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_buffer.i
2008-09-07 19:35:50 UTC (rev 9521)
+++
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_buffer.i
2008-09-07 19:41:42 UTC (rev 9522)
@@ -26,14 +26,14 @@
%rename(buffer) gr_make_buffer;
%ignore gr_buffer;
-gr_buffer_sptr gr_make_buffer (int nitems, size_t sizeof_item);
+gr_buffer_sptr gr_make_buffer (int nitems, size_t sizeof_item, gr_block_sptr
link);
class gr_buffer {
public:
~gr_buffer ();
private:
- gr_buffer (int nitems, size_t sizeof_item);
+ gr_buffer (int nitems, size_t sizeof_item, gr_block_sptr link);
};
@@ -43,7 +43,7 @@
%ignore gr_buffer_reader;
%rename(buffer_add_reader) gr_buffer_add_reader;
-gr_buffer_reader_sptr gr_buffer_add_reader (gr_buffer_sptr buf, int
nzero_preload);
+gr_buffer_reader_sptr gr_buffer_add_reader (gr_buffer_sptr buf, int
nzero_preload, gr_block_sptr link);
class gr_buffer_reader {
public:
@@ -51,7 +51,7 @@
private:
friend class gr_buffer;
- gr_buffer_reader (gr_buffer_sptr buffer, unsigned int read_index);
+ gr_buffer_reader (gr_buffer_sptr buffer, unsigned int read_index,
gr_block_sptr link);
};
Modified:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_flat_flowgraph.cc
===================================================================
---
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_flat_flowgraph.cc
2008-09-07 19:35:50 UTC (rev 9521)
+++
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_flat_flowgraph.cc
2008-09-07 19:41:42 UTC (rev 9522)
@@ -109,9 +109,9 @@
nitems = std::max(nitems,
static_cast<int>(2*(decimation*multiple+history)));
}
- if( GR_BUFFER_DEFAULT==block->output_signature()->buffer_type())
- return gr_make_buffer(nitems, item_size);
- else if(
GR_BUFFER_EXTERNAL_BLOCK_HAS_ITS_OWN_FACTORY==block->output_signature()->buffer_type())
+ //if( GR_BUFFER_DEFAULT==block->output_signature()->buffer_type())
+
+ if(
GR_BUFFER_EXTERNAL_BLOCK_HAS_ITS_OWN_FACTORY==block->output_signature()->buffer_type())
{
gr_buffer_sptr ext_buffer;
//gr_buffer_sptr ext_buffer=grblock->make_buffer(nitems, item_size);//TODO
implement grblock->make_buffer
@@ -120,9 +120,13 @@
return ext_buffer;
} else
{
- throw std::runtime_error("allocate_buffer block->output_signature() has
illegal buffer_type");
- //return 0;
+ return gr_make_buffer(nitems, item_size, grblock);
}
+ //else
+ //{
+ // throw std::runtime_error("allocate_buffer block->output_signature() has
illegal buffer_type");
+ // //return 0;
+ //}
}
void
@@ -151,7 +155,7 @@
if (GR_FLAT_FLOWGRAPH_DEBUG)
std::cout << "Setting input " << dst_port << " from edge " << (*e) <<
std::endl;
- detail->set_input(dst_port, gr_buffer_add_reader(src_buffer,
grblock->history()-1));
+ detail->set_input(dst_port, gr_buffer_add_reader(src_buffer,
grblock->history()-1, grblock));
}
}
@@ -238,7 +242,7 @@
std::cout << "needs a new reader" << std::endl;
// Create new buffer reader and assign
- detail->set_input(i, gr_buffer_add_reader(src_buffer,
block->history()-1));
+ detail->set_input(i, gr_buffer_add_reader(src_buffer,
block->history()-1, block));
}
}
}
@@ -282,3 +286,14 @@
}
}
+
+gr_block_vector_t
+gr_flat_flowgraph::make_block_vector(gr_basic_block_vector_t &blocks)
+{
+ gr_block_vector_t result;
+ for (gr_basic_block_viter_t p = blocks.begin(); p != blocks.end(); p++) {
+ result.push_back(make_gr_block_sptr(*p));
+ }
+
+ return result;
+}
Modified:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_flat_flowgraph.h
===================================================================
---
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_flat_flowgraph.h
2008-09-07 19:35:50 UTC (rev 9521)
+++
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_flat_flowgraph.h
2008-09-07 19:41:42 UTC (rev 9522)
@@ -55,6 +55,11 @@
void dump();
+ /*!
+ * Make a vector of gr_block from a vector of gr_basic_block
+ */
+ static gr_block_vector_t make_block_vector(gr_basic_block_vector_t &blocks);
+
private:
gr_flat_flowgraph();
Modified:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_io_signature.h
===================================================================
---
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_io_signature.h
2008-09-07 19:35:50 UTC (rev 9521)
+++
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_io_signature.h
2008-09-07 19:41:42 UTC (rev 9522)
@@ -25,13 +25,6 @@
#include <gr_runtime_types.h>
- enum gr_buffer_type {
- GR_BUFFER_DEFAULT = 0, // default buffer = circular buffer
in memory of host processor
- // created by
gr_vmcircbuf_sysconfig->get_default_factory()
- GR_BUFFER_EXTERNAL_BLOCK_HAS_ITS_OWN_FACTORY = 1, // block creates its
own buffer by block_instance->get_buffer_factory()
- // This memory might not be
in host main memory, so no direct copies
- // or pointer dereferencing
allowed.
- };
/*!
* \brief Create an i/o signature
*
Modified:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_runtime_types.h
===================================================================
---
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_runtime_types.h
2008-09-07 19:35:50 UTC (rev 9521)
+++
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_runtime_types.h
2008-09-07 19:41:42 UTC (rev 9522)
@@ -25,6 +25,18 @@
#include <gr_types.h>
+ enum gr_buffer_type {
+ GR_BUFFER_DEFAULT = 0, // default buffer = circular buffer
in memory of host processor
+ // created by
gr_vmcircbuf_sysconfig->get_default_factory()
+ GR_BUFFER_EXTERNAL_BLOCK_HAS_ITS_OWN_FACTORY = 1, // block creates its
own buffer by block_instance->get_buffer_factory()
+ // This memory might not be
in host main memory, so no direct copies
+ // or pointer dereferencing
allowed.
+ GR_BUFFER_CUDA = 2, //CUDA buffer lives on device memory
of videocard
+ //The circular buffer behind this
will be created with "gr_vmcircbuf_cuda_factory"
+ // This memory is not in host main
memory, so no direct copies
+ // or pointer dereferencing allowed.
+ };
+
/*
* typedefs for smart pointers we use throughout the runtime system
*/
Modified:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_single_threaded_scheduler.cc
===================================================================
---
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_single_threaded_scheduler.cc
2008-09-07 19:35:50 UTC (rev 9521)
+++
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_single_threaded_scheduler.cc
2008-09-07 19:41:42 UTC (rev 9522)
@@ -150,6 +150,10 @@
d_blocks[i]->detail()->set_done (false); // reset any done flags
for (unsigned i = 0; i < d_blocks.size (); i++) // enable any drivers,
etc.
+ for (unsigned j = 0 ; j< d_blocks[i]->detail()->noutputs ();j++)
+ d_blocks[i]->detail()->output(j)->start();
+
+ for (unsigned i = 0; i < d_blocks.size (); i++) // enable any drivers,
etc.
d_blocks[i]->start();
@@ -365,4 +369,8 @@
for (unsigned i = 0; i < d_blocks.size (); i++) // disable any drivers,
etc.
d_blocks[i]->stop();
+
+ for (unsigned i = 0; i < d_blocks.size (); i++) // disble any buffer
drivers, etc.
+ for (unsigned j = 0 ; j< d_blocks[i]->detail()->noutputs ();j++)
+ d_blocks[i]->detail()->output(j)->stop();
}
Modified:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_vmcircbuf.cc
===================================================================
---
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_vmcircbuf.cc
2008-09-07 19:35:50 UTC (rev 9521)
+++
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_vmcircbuf.cc
2008-09-07 19:41:42 UTC (rev 9522)
@@ -207,19 +207,44 @@
throw std::runtime_error ("gr_vmcircbuf_sysconfig");
}
+gr_vmcircbuf_factory *
+gr_vmcircbuf_sysconfig::get_factory (gr_buffer_type buffer_type)
+{
+ gr_vmcircbuf_factory *factory = 0;
+ switch(buffer_type)
+ {
+ case GR_BUFFER_DEFAULT:
+ factory=get_default_factory ();
+ break;
+ case GR_BUFFER_EXTERNAL_BLOCK_HAS_ITS_OWN_FACTORY:
+ throw std::runtime_error("get_factory is not able to make a
GR_BUFFER_EXTERNAL_BLOCK_HAS_ITS_OWN_FACTORY buffer.");
+ factory=0;
+ break;
+ case GR_BUFFER_CUDA:
+ factory=gr_vmcircbuf_cuda_factory::singleton ();
+ break;
+ default:
+ factory=get_default_factory ();
+ }
+ bool verbose = false;
+ if (verbose)
+ fprintf (stderr, "gr_vmcircbuf_sysconfig: using %s\n",
+ factory->name ());
+ return factory;
+}
+
std::vector<gr_vmcircbuf_factory *>
gr_vmcircbuf_sysconfig::all_factories ()
{
std::vector<gr_vmcircbuf_factory *> result;
- result.push_back (gr_vmcircbuf_cuda_factory::singleton ());
result.push_back (gr_vmcircbuf_createfilemapping_factory::singleton ());
result.push_back (gr_vmcircbuf_sysv_shm_factory::singleton ());
result.push_back (gr_vmcircbuf_mmap_shm_open_factory::singleton ());
result.push_back (gr_vmcircbuf_mmap_tmpfile_factory::singleton ());
result.push_back (gr_vmcircbuf_noncircular_factory::singleton ());
+ result.push_back (gr_vmcircbuf_cuda_factory::singleton ());
-
return result;
}
Modified:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_vmcircbuf.h
===================================================================
---
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_vmcircbuf.h
2008-09-07 19:35:50 UTC (rev 9521)
+++
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_vmcircbuf.h
2008-09-07 19:41:42 UTC (rev 9522)
@@ -25,6 +25,7 @@
#include <vector>
+#include <gr_runtime_types.h>
/*!
* \brief abstract class to implement doubly mapped virtual memory circular
buffers
* \ingroup base
@@ -42,6 +43,24 @@
public:
virtual ~gr_vmcircbuf ();
+ /*!
+ * \brief Called to enable drivers, etc for external memory.
+ *
+ * This allows a gr_vmcricbuf to enable an associated driver and memory
+ * just before we start to execute the scheduler.
+ * The actual allocation and de-allocation of memory can also be delayed
+ * untill start and stop are called.
+ * This way the allocation will be done by the same thread that calls
+ * work on the blocks using this buffer.
+ * This default implementation does nothing.
+ */
+ virtual bool start() { return true; }
+
+ /*!
+ * \brief Called to disable drivers, etc for external memory.
+ * This default implementation does nothing.
+ */
+ virtual bool stop() { return true; }
// ACCESSORS
void *pointer_to_first_copy () const { return d_base; }
void *pointer_to_second_copy () const { return d_base + d_size; }
@@ -251,11 +270,12 @@
* else find the first working factory and use it.
*/
static gr_vmcircbuf_factory *get_default_factory ();
-
+ static gr_vmcircbuf_factory *get_factory (gr_buffer_type buffer_type);
static int granularity () { return
get_default_factory()->granularity(); }
+ static int granularity (gr_buffer_type buffer_type) { return
get_factory(buffer_type)->granularity(); }
static gr_vmcircbuf *make (int size) { return
get_default_factory()->make(size); }
-
+ static gr_vmcircbuf *make (int size, gr_buffer_type buffer_type) { return
get_factory(buffer_type)->make(size); }
// N.B. not all factories are guaranteed to work.
// It's too hard to check everything at config time, so we check at runtime
Modified:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_vmcircbuf_cuda.cc
===================================================================
---
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_vmcircbuf_cuda.cc
2008-09-07 19:35:50 UTC (rev 9521)
+++
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_vmcircbuf_cuda.cc
2008-09-07 19:41:42 UTC (rev 9522)
@@ -51,29 +51,86 @@
#else
d_is_circular=false;
d_no_direct_pointer_access=true;
- cudaError_t res=cudaMalloc( (void**) &(d_base),size*2 );
- //CUDA_SAFE_CALL( cudaMalloc( (void**) &(d_base),size ));
+ d_base=NULL;
+ /*cudaError_t res=cudaMalloc( (void**) &(d_base),size*2 );
+ ////CUDA_SAFE_CALL( cudaMalloc( (void**) &(d_base),size ));
if (res != cudaSuccess){
d_base=NULL; // cleanup
perror ("gr_vmcircbuf_cuda: cudaMalloc");
throw std::runtime_error ("gr_vmcircbuf_cuda");
- }
-
+ }*/
// Now remember the important stuff
d_size = size;
+ fprintf (stderr, "gr_vmcircbuf_cuda::gr_vmcircbuf_cuda d_base=0x%x size=%i
*2 \n",d_base,size);//DEBUG
#endif
}
+bool
+gr_vmcircbuf_cuda::start()
+{
+bool retval=false;
+#if defined(HAVE_CUDA_RUNTIME_H)
+ if(d_base !=NULL)
+ fprintf (stderr, "gr_vmcircbuf_cuda::start() ERROR d_base != NULL. Buffer
already created? d_base=0x%x size=%i *2 \n",d_base,d_size);//DEBUG
+ //throw std::runtime_error ("gr_vmcircbuf_cuda:start() d_base != NULL.
Buffer already created?");
+ cudaError_t res=cudaMalloc( (void**) &(d_base),d_size*2 );
+ fprintf (stderr, "gr_vmcircbuf_cuda::start() d_base=0x%x size=%i *2
\n",d_base,d_size);//DEBUG
+ //CUDA_SAFE_CALL( cudaMalloc( (void**) &(d_base),size ));
+ if (res != cudaSuccess){
+ d_base=NULL; // cleanup
+ perror ("gr_vmcircbuf_cuda: cudaMalloc");
+ throw std::runtime_error ("gr_vmcircbuf_cuda");
+ } else retval=true;
+#endif /*HAVE_CUDA_RUNTIME_H*/
+return retval;
+}
+
+bool
+gr_vmcircbuf_cuda::stop()
+{
+bool retval=false;
+#if defined(HAVE_CUDA_RUNTIME_H)
+ fprintf (stderr, "gr_vmcircbuf_cuda::stop() d_base=0x%x size=%i *2
\n",d_base,d_size);//DEBUG
+ cudaError_t res=cudaFree(d_base);
+ if (res !=cudaSuccess){
+ fprintf(stderr,"ERROR: gr_vmcircbuf_cuda::stop() cudaFree didn't return
cudaSuccess\n");
+ perror ("gr_vmcircbuf_cuda::stop() cudaFree");
+ }else
+ {
+ d_base=NULL;
+ retval=true;
+ //fprintf(stderr,"OK: gr_vmcircbuf_cuda: cudaFree returned cudaSuccess\n");
+ }
+#endif /*HAVE_CUDA_RUNTIME_H*/
+ return retval;
+}
void
gr_vmcircbuf_cuda::copy_buf_to_buf(void * dst,void * src,unsigned int len)
{
#if defined(HAVE_CUDA_RUNTIME_H)
+ //fprintf (stderr, "gr_vmcircbuf_cuda::copy_buf_to_buf dst=0x%x src=0x%x
len=%i=0x%x \n",dst,src,len,len);//DEBUG
assert(len<=d_size);
cudaError_t res=cudaMemcpy(dst, src, len, cudaMemcpyDeviceToDevice);
if (res != cudaSuccess){
- perror ("gr_vmcircbuf_cuda:copy_buf_to_buf cudaMemcpy");
- throw std::runtime_error ("gr_vmcircbuf_cuda");
+ //try recreating the buffer only one time:
+ static bool first=true;
+ if(first)
+ {
+ cudaError_t res=cudaMalloc( (void**) &(d_base),d_size*2 );
+ fprintf (stderr, "gr_vmcircbuf_cuda::copy_buf_to_buf recreating buffer
d_base=0x%x size=%i *2 \n",d_base,d_size);//DEBUG
+ //CUDA_SAFE_CALL( cudaMalloc( (void**) &(d_base),size ));
+ if (res != cudaSuccess){
+ d_base=NULL; // cleanup
+ perror ("gr_vmcircbuf_cuda: cudaMalloc");
+ throw std::runtime_error ("gr_vmcircbuf_cuda");
+ }
+ first=false;
+ } else
+ {
+ perror ("gr_vmcircbuf_cuda:copy_buf_to_buf cudaMemcpy");
+ throw std::runtime_error ("gr_vmcircbuf_cuda");
+ }
}
#endif /*HAVE_CUDA_RUNTIME_H*/
}
@@ -108,7 +165,8 @@
gr_vmcircbuf_cuda::~gr_vmcircbuf_cuda ()
{
-#if defined(HAVE_CUDA_RUNTIME_H)
+#if defined(HAVE_CUDA_RUNTIME_H)
+ fprintf (stderr, "gr_vmcircbuf_cuda::~gr_vmcircbuf_cuda() d_base=0x%x
size=%i *2 \n",d_base,d_size);//DEBUG
cudaError_t res=cudaFree(d_base);
if (res !=cudaSuccess){
fprintf(stderr,"ERROR: gr_vmcircbuf_cuda: cudaFree didn't return
cudaSuccess\n");
Modified:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_vmcircbuf_cuda.h
===================================================================
---
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_vmcircbuf_cuda.h
2008-09-07 19:35:50 UTC (rev 9521)
+++
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/gr_vmcircbuf_cuda.h
2008-09-07 19:41:42 UTC (rev 9522)
@@ -36,6 +36,27 @@
gr_vmcircbuf_cuda (int size);
virtual ~gr_vmcircbuf_cuda ();
+
+ /*!
+ * \brief Called to enable drivers, etc for external memory.
+ *
+ * This allows a gr_vmcricbuf to enable an associated driver and memory
+ * just before we start to execute the scheduler.
+ * The actual allocation of this implementation
+ * will be delayed untill start is called.
+ * This way the allocation will be done by the same thread that calls
+ * work on the blocks using this buffer.
+ */
+ virtual bool start();
+
+ /*!
+ * \brief Called to disable drivers, etc for external memory.
+ * The actual de-allocation of this implementation
+ * will be delayed untill stop is called.
+ * This way the de-allocation will be done by the same thread that calls
+ * work on the blocks using this buffer.
+ */
+ virtual bool stop();
// MODIFIERS
/*!
* \brief copy buffer memory into another place in this buffer or from one
buffer to a similar buffer.
Modified:
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/qa_gr_buffer.cc
===================================================================
---
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/qa_gr_buffer.cc
2008-09-07 19:35:50 UTC (rev 9521)
+++
gnuradio/branches/developers/nldudok1/gpgpu-wip/gnuradio-core/src/lib/runtime/qa_gr_buffer.cc
2008-09-07 19:41:42 UTC (rev 9522)
@@ -29,6 +29,8 @@
#include <stdlib.h>
#include <gr_random.h>
+#include <gr_nop.h>
+
static void
leak_check (void f ())
{
@@ -52,7 +54,8 @@
int nitems = 4000 / sizeof (int);
int counter = 0;
- gr_buffer_sptr buf (gr_make_buffer (nitems, sizeof (int)));
+ //gr_buffer_sptr buf(gr_make_buffer(nitems, sizeof (int), gr_block_sptr()));
+ gr_buffer_sptr buf(gr_make_buffer(nitems, sizeof (int), gr_make_nop(1)));
int last_sa;
int sa;
@@ -87,8 +90,8 @@
int write_counter = 0;
int read_counter = 0;
- gr_buffer_sptr buf (gr_make_buffer (nitems, sizeof (int)));
- gr_buffer_reader_sptr r1 (gr_buffer_add_reader (buf, 0));
+ gr_buffer_sptr buf(gr_make_buffer(nitems, sizeof (int), gr_make_nop(1)));
+ gr_buffer_reader_sptr r1 (gr_buffer_add_reader (buf, 0, gr_make_nop(1)));
int sa;
@@ -162,8 +165,8 @@
int nitems = (64 * (1L << 10)) / sizeof (int); // 64K worth of ints
- gr_buffer_sptr buf (gr_make_buffer (nitems, sizeof (int)));
- gr_buffer_reader_sptr r1 (gr_buffer_add_reader (buf, 0));
+ gr_buffer_sptr buf(gr_make_buffer (nitems, sizeof (int), gr_make_nop(1)));
+ gr_buffer_reader_sptr r1 (gr_buffer_add_reader (buf, 0, gr_make_nop(1)));
int read_counter = 0;
int write_counter = 0;
@@ -229,7 +232,7 @@
int nitems = (64 * (1L << 10)) / sizeof (int);
static const int N = 5;
- gr_buffer_sptr buf (gr_make_buffer (nitems, sizeof (int)));
+ gr_buffer_sptr buf(gr_make_buffer(nitems, sizeof (int), gr_make_nop(1)));
gr_buffer_reader_sptr reader[N];
int read_counter[N];
int write_counter = 0;
@@ -237,7 +240,7 @@
for (int i = 0; i < N; i++){
read_counter[i] = 0;
- reader[i] = gr_buffer_add_reader (buf, 0);
+ reader[i] = gr_buffer_add_reader (buf, 0, gr_make_nop(1));
}
for (int lc = 0; lc < 1000; lc++){
Added:
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/cuda_test_dial_tone.py
===================================================================
---
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/cuda_test_dial_tone.py
(rev 0)
+++
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/cuda_test_dial_tone.py
2008-09-07 19:41:42 UTC (rev 9522)
@@ -0,0 +1,66 @@
+#!/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
+
+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("-O", "--audio-output", type="string", default="",
+ help="pcm output device name. E.g., hw:0,0 or
/dev/dsp")
+ parser.add_option("-r", "--sample-rate", type="eng_float",
default=48000,
+ help="set sample rate to RATE (48000)")
+ (options, args) = parser.parse_args ()
+ if len(args) != 0:
+ parser.print_help()
+ raise SystemExit, 1
+
+ sample_rate = int(options.sample_rate)
+ ampl = 0.1
+
+ src0 = gr.sig_source_f (sample_rate, gr.GR_SIN_WAVE, 350, ampl)
+ src1 = gr.sig_source_f (sample_rate, gr.GR_SIN_WAVE, 440, ampl)
+ h2c=cuda.host_to_cuda(gr.sizeof_float)
+ c2h=cuda.cuda_to_host(gr.sizeof_float)
+ dst = audio.sink (sample_rate, options.audio_output)
+ self.connect (src0, (dst, 0))
+ self.connect (src1, h2c,c2h,(dst, 1))
+
+
+if __name__ == '__main__':
+ try:
+ # insert this in your test code...
+ 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/cuda_test_dial_tone.py
___________________________________________________________________
Name: svn:executable
+ *
Added:
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/cuda_test_dial_tone_wav.py
===================================================================
---
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/cuda_test_dial_tone_wav.py
(rev 0)
+++
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/cuda_test_dial_tone_wav.py
2008-09-07 19:41:42 UTC (rev 9522)
@@ -0,0 +1,65 @@
+#!/usr/bin/env python
+#
+# Copyright 2004,2005,2007,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.
+#
+
+# GNU Radio example program to record a dial tone to a WAV file
+
+from gnuradio import gr
+from gnuradio.eng_option import eng_option
+from optparse import OptionParser
+
+from gnuradio import cuda
+
+class my_top_block(gr.top_block):
+
+ def __init__(self):
+ gr.top_block.__init__(self)
+
+ usage = "%prog: [options] filename"
+ parser = OptionParser(option_class=eng_option, usage=usage)
+ parser.add_option("-r", "--sample-rate", type="eng_float",
default=48000,
+ help="set sample rate to RATE (48000)")
+ parser.add_option("-N", "--samples", type="eng_float", default=None,
+ help="number of samples to record")
+ (options, args) = parser.parse_args ()
+ if len(args) != 1 or options.samples is None:
+ parser.print_help()
+ raise SystemExit, 1
+
+ sample_rate = int(options.sample_rate)
+ ampl = 0.1
+
+ src0 = gr.sig_source_f (sample_rate, gr.GR_SIN_WAVE, 350, ampl)
+ src1 = gr.sig_source_f (sample_rate, gr.GR_SIN_WAVE, 440, ampl)
+ head0 = gr.head(gr.sizeof_float, int(options.samples))
+ head1 = gr.head(gr.sizeof_float, int(options.samples))
+ h2c=cuda.host_to_cuda(gr.sizeof_float)
+ c2h=cuda.cuda_to_host(gr.sizeof_float)
+ dst = gr.wavfile_sink(args[0], 2, int(options.sample_rate), 16)
+
+ self.connect(src0, head0, (dst, 0))
+ self.connect(src1, head1, h2c , c2h, (dst, 1))
+
+if __name__ == '__main__':
+ try:
+ my_top_block().run()
+ except KeyboardInterrupt:
+ pass
Property changes on:
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/cuda_test_dial_tone_wav.py
___________________________________________________________________
Name: svn:executable
+ *
Added:
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/cuda_test_vector.py
===================================================================
--- gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/cuda_test_vector.py
(rev 0)
+++ gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/cuda_test_vector.py
2008-09-07 19:41:42 UTC (rev 9522)
@@ -0,0 +1,86 @@
+#!/usr/bin/env python
+#
+# Copyright 2004,2005,2007,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.
+#
+
+# GNU Radio example program to record a dial tone to a WAV file
+
+from gnuradio import gr
+from gnuradio.eng_option import eng_option
+from optparse import OptionParser
+
+from gnuradio import cuda
+def ramp_source_s ():
+ period = 2**16
+ src = gr.vector_source_s (range (-period/2, period/2, 1), True)
+ return src
+
+
+class my_top_block(gr.top_block):
+
+ def __init__(self):
+ gr.top_block.__init__(self)
+
+ usage = "%prog: [options]"
+ parser = OptionParser(option_class=eng_option, usage=usage)
+ parser.add_option("-N", "--samples", type="eng_float", default=None,
+ help="number of samples to process")
+ (options, args) = parser.parse_args ()
+ if len(args) != 0 or options.samples is None:
+ parser.print_help()
+ raise SystemExit, 1
+
+ self.nsamples=int(options.samples)
+ period = 2**16
+ self.data_in=range (-period/2, period/2, 1)
+ src = gr.vector_source_s (self.data_in, True)
+ head = gr.head(gr.sizeof_short, self.nsamples)
+ h2c=cuda.host_to_cuda(gr.sizeof_short)
+ c2h=cuda.cuda_to_host(gr.sizeof_short)
+ self.dst = gr.vector_sink_s()
+ self.dst_cuda= gr.vector_sink_s()
+
+ self.connect(src, head, self.dst)
+ self.connect(head, h2c,c2h,self.dst_cuda)
+
+ def check(self):
+ self.data_out=self.dst.data()
+ self.data_out_cuda=self.dst_cuda.data()
+ len1=self.nsamples
+ len2=len(self.data_in)
+ len3=len(self.data_out)
+ #print self.data_out
+ print "len1,len2,len3=",len1,len2,len3
+ for i in range(self.nsamples):
+ if self.data_out[i]!=self.data_out_cuda[i]:
+ print "FAULT at i=",i," self.data_out[i]=",self.data_out[i],"
self.data_out_cuda[i]=",self.data_out_cuda[i]
+ elif i>24570:
+ print "OK at i=",i," self.data_out[i]=",self.data_out[i],"
self.data_out_cuda[i]=",self.data_out_cuda[i]
+
+
+if __name__ == '__main__':
+ try:
+ tb=my_top_block()
+ tb.run()
+ tb.check()
+ except KeyboardInterrupt:
+ pass
+
+
Property changes on:
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/cuda_test_vector.py
___________________________________________________________________
Name: svn:executable
+ *
Added:
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/cuda_test_vector_int.py
===================================================================
---
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/cuda_test_vector_int.py
(rev 0)
+++
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/cuda_test_vector_int.py
2008-09-07 19:41:42 UTC (rev 9522)
@@ -0,0 +1,99 @@
+#!/usr/bin/env python
+#
+# Copyright 2004,2005,2007,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.
+#
+
+# GNU Radio example program to record a dial tone to a WAV file
+
+from gnuradio import gr
+from gnuradio.eng_option import eng_option
+from optparse import OptionParser
+
+from gnuradio import cuda
+def ramp_source_s ():
+ period = 2**16
+ src = gr.vector_source_s (range (-period/2, period/2, 1), True)
+ return src
+
+
+class my_top_block(gr.top_block):
+
+ def __init__(self):
+ gr.top_block.__init__(self)
+
+ usage = "%prog: [options]"
+ parser = OptionParser(option_class=eng_option, usage=usage)
+ parser.add_option("-N", "--samples", type="eng_float", default=None,
+ help="number of samples to process")
+ (options, args) = parser.parse_args ()
+ if len(args) != 0 or options.samples is None:
+ parser.print_help()
+ raise SystemExit, 1
+
+ self.nsamples=int(options.samples)
+ period = self.nsamples #2**16
+ self.data_in=range (0x00ffffff, period+0x00ffffff, 1) #range
(-period/2, period/2, 1)
+ #self.data_in=range (0,2,1)
+ src = gr.vector_source_i (self.data_in, True)
+ head = gr.head(gr.sizeof_int, self.nsamples)
+ h2c=cuda.host_to_cuda(gr.sizeof_int)
+ c2h=cuda.cuda_to_host(gr.sizeof_int)
+ self.dst = gr.vector_sink_i()
+ self.dst_cuda= gr.vector_sink_i()
+
+ self.connect(src, head, self.dst)
+ self.connect(head, h2c,c2h,self.dst_cuda)
+
+ def check(self):
+ self.data_out=self.dst.data()
+ self.data_out_cuda=self.dst_cuda.data()
+ len1=self.nsamples
+ len2=len(self.data_in)
+ len3=len(self.data_out)
+ status=(self.data_out[0]==self.data_out_cuda[0])
+ last_status=status
+ #print self.data_out
+ counter=0
+ print "len1,len2,len3=",len1,len2,len3
+ for i in range(self.nsamples):
+ last_status=status
+ counter=counter+1
+ if self.data_out[i]!=self.data_out_cuda[i]:
+ status=False
+ #print "FAULT at i=",i," self.data_out[i]=",self.data_out[i],"
self.data_out_cuda[i]=",self.data_out_cuda[i]
+ else:
+ status=True
+ #print "OK at i=",i," self.data_out[i]=",self.data_out[i],"
self.data_out_cuda[i]=",self.data_out_cuda[i]
+ if(status!=last_status):
+ print last_status, " at i=",i-1," data_out=",self.data_out[i-1],"
data_out_cuda=",self.data_out_cuda[i-1]
+ print "xor at",i-1, self.data_out[i-1] ^ self.data_out_cuda[i-1]
+ print status, " at i=",i," data_out=",self.data_out[i],"
data_out_cuda=",self.data_out_cuda[i]
+ print "xor at",i, self.data_out[i] ^ self.data_out_cuda[i]
+ print "counter=",counter," ..."
+ counter=0
+if __name__ == '__main__':
+ try:
+ tb=my_top_block()
+ tb.run()
+ tb.check()
+ except KeyboardInterrupt:
+ pass
+
+
Property changes on:
gnuradio/branches/developers/nldudok1/gpgpu-wip/testbed/cuda_test_vector_int.py
___________________________________________________________________
Name: svn:executable
+ *
| [Prev in Thread] |
Current Thread |
[Next in Thread] |
- [Commit-gnuradio] r9522 - in gnuradio/branches/developers/nldudok1/gpgpu-wip: . gnuradio-core/src/lib/runtime testbed,
nldudok1 <=