uhd: removed convert types, replaced by convert

4 files changed, 17 insertions, 631 deletions

diff --git a/host/lib/convert/gen_convert_impl.py b/host/lib/convert/gen_convert_impl.py
index 71095ab97..70d437ba2 100644
--- a/host/lib/convert/gen_convert_impl.py
+++ b/host/lib/convert/gen_convert_impl.py
@@ -26,6 +26,7 @@ typedef size_t pred_type;
 \#include <uhd/utils/algorithm.hpp>
 \#include <boost/lexical_cast.hpp>
 \#include <boost/detail/endian.hpp>
+\#include <boost/cstdint.hpp>
 \#include <stdexcept>
 
 enum dir_type{
@@ -33,6 +34,13 @@ enum dir_type{
     DIR_CPU_TO_OTW = 1
 };
 
+struct pred_error : std::runtime_error{
+    pred_error(const std::string &what)
+    :std::runtime_error("convert::make_pred: " + what){
+        /* NOP */
+    }
+};
+
 pred_type make_pred(const std::string &markup, dir_type &dir){
     pred_type pred = 0;
 
@@ -60,10 +68,10 @@ pred_type make_pred(const std::string &markup, dir_type &dir){
 
         if      (cpu_type == "fc32") pred |= $ph.fc32_p;
         else if (cpu_type == "sc16") pred |= $ph.sc16_p;
-        else throw std::runtime_error("unhandled io type " + cpu_type);
+        else throw pred_error("unhandled io type " + cpu_type);
 
         if (otw_type == "item32") pred |= $ph.item32_p;
-        else throw std::runtime_error("unhandled otw type " + otw_type);
+        else throw pred_error("unhandled otw type " + otw_type);
 
         int num_inputs = boost::lexical_cast<int>(num_inps);
         int num_outputs = boost::lexical_cast<int>(num_outs);
@@ -73,16 +81,16 @@ pred_type make_pred(const std::string &markup, dir_type &dir){
         case 2: pred |= $ph.chan2_p; break;
         case 3: pred |= $ph.chan3_p; break;
         case 4: pred |= $ph.chan4_p; break;
-        default: throw std::runtime_error("unhandled number of channels");
+        default: throw pred_error("unhandled number of channels");
         }
 
         if      (swap_type == "bswap") pred |= $ph.bswap_p;
         else if (swap_type == "nswap") pred |= $ph.nswap_p;
-        else throw std::runtime_error("unhandled swap type");
+        else throw pred_error("unhandled swap type");
 
     }
     catch(...){
-        throw std::runtime_error("convert::make_pred: could not parse markup: " + markup);
+        throw pred_error("could not parse markup: " + markup);
     }
 
     return pred;
@@ -105,18 +113,18 @@ UHD_INLINE pred_type make_pred(
     case otw_type_t::BO_LITTLE_ENDIAN: pred |= $ph.nswap_p; break;
     \#endif
     case otw_type_t::BO_NATIVE:        pred |= $ph.nswap_p; break;
-    default: throw std::runtime_error("unhandled otw byteorder type");
+    default: throw pred_error("unhandled otw byteorder type");
     }
 
     switch(otw_type.get_sample_size()){
     case sizeof(boost::uint32_t): pred |= $ph.item32_p; break;
-    default: throw std::runtime_error("unhandled otw sample size");
+    default: throw pred_error("unhandled otw sample size");
     }
 
     switch(io_type.tid){
     case io_type_t::COMPLEX_FLOAT32: pred |= $ph.fc32_p; break;
     case io_type_t::COMPLEX_INT16:   pred |= $ph.sc16_p; break;
-    default: throw std::runtime_error("unhandled io type id");
+    default: throw pred_error("unhandled io type id");
     }
 
     switch(num_inputs*num_outputs){ //FIXME treated as one value
@@ -124,7 +132,7 @@ UHD_INLINE pred_type make_pred(
     case 2: pred |= $ph.chan2_p; break;
     case 3: pred |= $ph.chan3_p; break;
     case 4: pred |= $ph.chan4_p; break;
-    default: throw std::runtime_error("unhandled number of channels");
+    default: throw pred_error("unhandled number of channels");
     }
 
     return pred;
@@ -148,38 +156,6 @@ class ph:
 
     nbits = 4 #see above
 
-    @staticmethod
-    def has(pred, mask, flag): return (pred & mask) == flag
-
-    @staticmethod
-    def get_swap_type(pred):
-        mask = 0b1
-        if ph.has(pred, mask, ph.bswap_p): return 'bswap'
-        if ph.has(pred, mask, ph.nswap_p): return 'nswap'
-        raise NotImplementedError
-
-    @staticmethod
-    def get_dev_type(pred):
-        mask = 0b0
-        if ph.has(pred, mask, ph.item32_p): return 'item32'
-        raise NotImplementedError
-
-    @staticmethod
-    def get_host_type(pred):
-        mask = 0b10
-        if ph.has(pred, mask, ph.sc16_p): return 'sc16'
-        if ph.has(pred, mask, ph.fc32_p): return 'fc32'
-        raise NotImplementedError
-
-    @staticmethod
-    def get_num_chans(pred):
-        mask = 0b1100
-        if ph.has(pred, mask, ph.chan1_p): return 1
-        if ph.has(pred, mask, ph.chan2_p): return 2
-        if ph.has(pred, mask, ph.chan3_p): return 3
-        if ph.has(pred, mask, ph.chan4_p): return 4
-        raise NotImplementedError
-
 if __name__ == '__main__':
     import sys, os
     file = os.path.basename(__file__)
diff --git a/host/lib/transport/CMakeLists.txt b/host/lib/transport/CMakeLists.txt
index ac92ffe6d..914ba3ba9 100644
--- a/host/lib/transport/CMakeLists.txt
+++ b/host/lib/transport/CMakeLists.txt
@@ -47,25 +47,6 @@ ELSE(ENABLE_USB)
 ENDIF(ENABLE_USB)
 
 ########################################################################
-# Check for SIMD headers
-########################################################################
-MESSAGE(STATUS "")
-
-INCLUDE(CheckIncludeFileCXX)
-CHECK_INCLUDE_FILE_CXX(emmintrin.h HAVE_EMMINTRIN_H)
-
-IF(HAVE_EMMINTRIN_H)
-    ADD_DEFINITIONS(-DHAVE_EMMINTRIN_H)
-ENDIF(HAVE_EMMINTRIN_H)
-
-INCLUDE(CheckIncludeFileCXX)
-CHECK_INCLUDE_FILE_CXX(arm_neon.h HAVE_ARM_NEON_H)
-
-IF(HAVE_ARM_NEON_H)
-    ADD_DEFINITIONS(-DHAVE_ARM_NEON_H)
-ENDIF(HAVE_ARM_NEON_H)
-
-########################################################################
 # Setup defines for interface address discovery
 ########################################################################
 MESSAGE(STATUS "")
@@ -93,21 +74,6 @@ LIBUHD_PYTHON_GEN_SOURCE(
     ${CMAKE_CURRENT_BINARY_DIR}/vrt_if_packet.cpp
 )
 
-LIBUHD_PYTHON_GEN_SOURCE(
-    ${CMAKE_CURRENT_SOURCE_DIR}/gen_convert_types.py
-    ${CMAKE_CURRENT_BINARY_DIR}/convert_types.cpp
-)
-
-# append this directory to the include path so the generated convert types
-# can include the implementation convert types file in the source directory
-INCLUDE_DIRECTORIES(${CMAKE_CURRENT_SOURCE_DIR})
-
-# make the generated convert types depend on the implementation header
-SET_SOURCE_FILES_PROPERTIES(
-    ${CMAKE_CURRENT_BINARY_DIR}/convert_types.cpp PROPERTIES
-    OBJECT_DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/convert_types_impl.hpp
-)
-
 LIBUHD_APPEND_SOURCES(
     ${CMAKE_CURRENT_SOURCE_DIR}/buffer_pool.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/if_addrs.cpp
diff --git a/host/lib/transport/convert_types_impl.hpp b/host/lib/transport/convert_types_impl.hpp
deleted file mode 100644
index 48ff99725..000000000
--- a/host/lib/transport/convert_types_impl.hpp
+++ /dev/null
@@ -1,345 +0,0 @@
-//
-// Copyright 2010 Ettus Research LLC
-//
-// This program 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 of the License, or
-// (at your option) any later version.
-//
-// This program 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 this program.  If not, see <http://www.gnu.org/licenses/>.
-//
-
-#ifndef INCLUDED_LIBUHD_TRANSPORT_CONVERT_TYPES_IMPL_HPP
-#define INCLUDED_LIBUHD_TRANSPORT_CONVERT_TYPES_IMPL_HPP
-
-#include <uhd/config.hpp>
-#include <uhd/utils/byteswap.hpp>
-#include <boost/cstdint.hpp>
-#include <cstring>
-#include <complex>
-
-#ifdef HAVE_EMMINTRIN_H
-    #define USE_EMMINTRIN_H //use sse2 intrinsics
-#endif
-
-#if defined(USE_EMMINTRIN_H)
-    #include <emmintrin.h>
-#endif
-
-#ifdef HAVE_ARM_NEON_H
-    #define USE_ARM_NEON_H
-#endif
-
-#if defined(USE_ARM_NEON_H)
-    #include <arm_neon.h>
-#endif
-
-/***********************************************************************
- * Typedefs
- **********************************************************************/
-typedef std::complex<float>          fc32_t;
-typedef std::complex<boost::int16_t> sc16_t;
-typedef boost::uint32_t              item32_t;
-
-/***********************************************************************
- * Convert complex short buffer to items32
- **********************************************************************/
-static UHD_INLINE item32_t sc16_to_item32(sc16_t num){
-    boost::uint16_t real = num.real();
-    boost::uint16_t imag = num.imag();
-    return (item32_t(real) << 16) | (item32_t(imag) << 0);
-}
-
-static UHD_INLINE void sc16_to_item32_nswap(
-    const sc16_t *input, item32_t *output, size_t nsamps
-){
-    for (size_t i = 0; i < nsamps; i++){
-        output[i] = sc16_to_item32(input[i]);
-    }
-}
-
-static UHD_INLINE void sc16_to_item32_bswap(
-    const sc16_t *input, item32_t *output, size_t nsamps
-){
-    for (size_t i = 0; i < nsamps; i++){
-        output[i] = uhd::byteswap(sc16_to_item32(input[i]));
-    }
-}
-
-/***********************************************************************
- * Convert items32 buffer to complex short
- **********************************************************************/
-static UHD_INLINE sc16_t item32_to_sc16(item32_t item){
-    return sc16_t(
-        boost::int16_t(item >> 16),
-        boost::int16_t(item >> 0)
-    );
-}
-
-static UHD_INLINE void item32_to_sc16_nswap(
-    const item32_t *input, sc16_t *output, size_t nsamps
-){
-    for (size_t i = 0; i < nsamps; i++){
-        output[i] = item32_to_sc16(input[i]);
-    }
-}
-
-static UHD_INLINE void item32_to_sc16_bswap(
-    const item32_t *input, sc16_t *output, size_t nsamps
-){
-    for (size_t i = 0; i < nsamps; i++){
-        output[i] = item32_to_sc16(uhd::byteswap(input[i]));
-    }
-}
-
-/***********************************************************************
- * Convert complex float buffer to items32 (no swap)
- **********************************************************************/
-static const float shorts_per_float = float(32767);
-
-static UHD_INLINE item32_t fc32_to_item32(fc32_t num){
-    boost::uint16_t real = boost::int16_t(num.real()*shorts_per_float);
-    boost::uint16_t imag = boost::int16_t(num.imag()*shorts_per_float);
-    return (item32_t(real) << 16) | (item32_t(imag) << 0);
-}
-
-////////////////////////////////////
-// none-swap
-////////////////////////////////////
-#if defined(USE_EMMINTRIN_H)
-static UHD_INLINE void fc32_to_item32_nswap(
-    const fc32_t *input, item32_t *output, size_t nsamps
-){
-    __m128 scalar = _mm_set_ps1(shorts_per_float);
-
-    //convert blocks of samples with intrinsics
-    size_t i = 0; for (; i < (nsamps & ~0x3); i+=4){
-        //load from input
-        __m128 tmplo = _mm_loadu_ps(reinterpret_cast<const float *>(input+i+0));
-        __m128 tmphi = _mm_loadu_ps(reinterpret_cast<const float *>(input+i+2));
-
-        //convert and scale
-        __m128i tmpilo = _mm_cvtps_epi32(_mm_mul_ps(tmplo, scalar));
-        __m128i tmpihi = _mm_cvtps_epi32(_mm_mul_ps(tmphi, scalar));
-
-        //pack + swap 16-bit pairs
-        __m128i tmpi = _mm_packs_epi32(tmpilo, tmpihi);
-        tmpi = _mm_shufflelo_epi16(tmpi, _MM_SHUFFLE(2, 3, 0, 1));
-        tmpi = _mm_shufflehi_epi16(tmpi, _MM_SHUFFLE(2, 3, 0, 1));
-
-        //store to output
-        _mm_storeu_si128(reinterpret_cast<__m128i *>(output+i), tmpi);
-    }
-
-    //convert remainder
-    for (; i < nsamps; i++){
-        output[i] = fc32_to_item32(input[i]);
-    }
-}
-
-#elif defined(USE_ARM_NEON_H)
-static UHD_INLINE void fc32_to_item32_nswap(
-    const fc32_t *input, item32_t *output, size_t nsamps)
-{
-    size_t i;
-
-    float32x4_t Q0 = vdupq_n_f32(shorts_per_float);
-    for (i=0; i < (nsamps & ~0x03); i+=2) {
-        float32x4_t Q1 = vld1q_f32(reinterpret_cast<const float *>(&input[i]));
-        float32x4_t Q2 = vmulq_f32(Q1, Q0);
-        int32x4_t Q3 = vcvtq_s32_f32(Q2);
-        int16x4_t D8 = vmovn_s32(Q3);
-        int16x4_t D9 = vrev32_s16(D8);
-        vst1_s16((reinterpret_cast<int16_t *>(&output[i])), D9);
-    }
-
-    for (; i < nsamps; i++)
-        output[i] = fc32_to_item32(input[i]);
-}
-
-#else
-static UHD_INLINE void fc32_to_item32_nswap(
-    const fc32_t *input, item32_t *output, size_t nsamps
-){
-    for (size_t i = 0; i < nsamps; i++){
-        output[i] = fc32_to_item32(input[i]);
-    }
-}
-
-#endif
-
-////////////////////////////////////
-// byte-swap
-////////////////////////////////////
-#if defined(USE_EMMINTRIN_H)
-static UHD_INLINE void fc32_to_item32_bswap(
-    const fc32_t *input, item32_t *output, size_t nsamps
-){
-    __m128 scalar = _mm_set_ps1(shorts_per_float);
-
-    //convert blocks of samples with intrinsics
-    size_t i = 0; for (; i < (nsamps & ~0x3); i+=4){
-        //load from input
-        __m128 tmplo = _mm_loadu_ps(reinterpret_cast<const float *>(input+i+0));
-        __m128 tmphi = _mm_loadu_ps(reinterpret_cast<const float *>(input+i+2));
-
-        //convert and scale
-        __m128i tmpilo = _mm_cvtps_epi32(_mm_mul_ps(tmplo, scalar));
-        __m128i tmpihi = _mm_cvtps_epi32(_mm_mul_ps(tmphi, scalar));
-
-        //pack + byteswap -> byteswap 16 bit words
-        __m128i tmpi = _mm_packs_epi32(tmpilo, tmpihi);
-        tmpi = _mm_or_si128(_mm_srli_epi16(tmpi, 8), _mm_slli_epi16(tmpi, 8));
-
-        //store to output
-        _mm_storeu_si128(reinterpret_cast<__m128i *>(output+i), tmpi);
-    }
-
-    //convert remainder
-    for (; i < nsamps; i++){
-        output[i] = uhd::byteswap(fc32_to_item32(input[i]));
-    }
-}
-
-#else
-static UHD_INLINE void fc32_to_item32_bswap(
-    const fc32_t *input, item32_t *output, size_t nsamps
-){
-    for (size_t i = 0; i < nsamps; i++){
-        output[i] = uhd::byteswap(fc32_to_item32(input[i]));
-    }
-}
-
-#endif
-
-/***********************************************************************
- * Convert items32 buffer to complex float
- **********************************************************************/
-static const float floats_per_short = float(1.0/shorts_per_float);
-
-static UHD_INLINE fc32_t item32_to_fc32(item32_t item){
-    return fc32_t(
-        float(boost::int16_t(item >> 16)*floats_per_short),
-        float(boost::int16_t(item >> 0)*floats_per_short)
-    );
-}
-
-////////////////////////////////////
-// none-swap
-////////////////////////////////////
-#if defined(USE_EMMINTRIN_H)
-static UHD_INLINE void item32_to_fc32_nswap(
-    const item32_t *input, fc32_t *output, size_t nsamps
-){
-    __m128 scalar = _mm_set_ps1(floats_per_short/(1 << 16));
-    __m128i zeroi = _mm_setzero_si128();
-
-    //convert blocks of samples with intrinsics
-    size_t i = 0; for (; i < (nsamps & ~0x3); i+=4){
-        //load from input
-        __m128i tmpi = _mm_loadu_si128(reinterpret_cast<const __m128i *>(input+i));
-
-        //unpack + swap 16-bit pairs
-        tmpi = _mm_shufflelo_epi16(tmpi, _MM_SHUFFLE(2, 3, 0, 1));
-        tmpi = _mm_shufflehi_epi16(tmpi, _MM_SHUFFLE(2, 3, 0, 1));
-        __m128i tmpilo = _mm_unpacklo_epi16(zeroi, tmpi); //value in upper 16 bits
-        __m128i tmpihi = _mm_unpackhi_epi16(zeroi, tmpi);
-
-        //convert and scale
-        __m128 tmplo = _mm_mul_ps(_mm_cvtepi32_ps(tmpilo), scalar);
-        __m128 tmphi = _mm_mul_ps(_mm_cvtepi32_ps(tmpihi), scalar);
-
-        //store to output
-        _mm_storeu_ps(reinterpret_cast<float *>(output+i+0), tmplo);
-        _mm_storeu_ps(reinterpret_cast<float *>(output+i+2), tmphi);
-    }
-
-    //convert remainder
-    for (; i < nsamps; i++){
-        output[i] = item32_to_fc32(input[i]);
-    }
-}
-
-#elif defined(USE_ARM_NEON_H)
-static UHD_INLINE void item32_to_fc32_nswap(
-    const item32_t *input, fc32_t *output, size_t nsamps)
-{
-    size_t i;
-
-    float32x4_t Q1 = vdupq_n_f32(floats_per_short);
-    for (i=0; i < (nsamps & ~0x03); i+=2) {
-        int16x4_t D0 = vld1_s16(reinterpret_cast<const int16_t *>(&input[i]));
-        int16x4_t D1 = vrev32_s16(D0);
-        int32x4_t Q2 = vmovl_s16(D1);
-        float32x4_t Q3 = vcvtq_f32_s32(Q2);
-        float32x4_t Q4 = vmulq_f32(Q3, Q1);
-        vst1q_f32((reinterpret_cast<float *>(&output[i])), Q4);
-    }
-
-    for (; i < nsamps; i++)
-        output[i] = item32_to_fc32(input[i]);
-}
-
-#else
-static UHD_INLINE void item32_to_fc32_nswap(
-    const item32_t *input, fc32_t *output, size_t nsamps
-){
-    for (size_t i = 0; i < nsamps; i++){
-        output[i] = item32_to_fc32(input[i]);
-    }
-}
-#endif
-
-////////////////////////////////////
-// byte-swap
-////////////////////////////////////
-#if defined(USE_EMMINTRIN_H)
-static UHD_INLINE void item32_to_fc32_bswap(
-    const item32_t *input, fc32_t *output, size_t nsamps
-){
-    __m128 scalar = _mm_set_ps1(floats_per_short/(1 << 16));
-    __m128i zeroi = _mm_setzero_si128();
-
-    //convert blocks of samples with intrinsics
-    size_t i = 0; for (; i < (nsamps & ~0x3); i+=4){
-        //load from input
-        __m128i tmpi = _mm_loadu_si128(reinterpret_cast<const __m128i *>(input+i));
-
-        //byteswap + unpack -> byteswap 16 bit words
-        tmpi = _mm_or_si128(_mm_srli_epi16(tmpi, 8), _mm_slli_epi16(tmpi, 8));
-        __m128i tmpilo = _mm_unpacklo_epi16(zeroi, tmpi); //value in upper 16 bits
-        __m128i tmpihi = _mm_unpackhi_epi16(zeroi, tmpi);
-
-        //convert and scale
-        __m128 tmplo = _mm_mul_ps(_mm_cvtepi32_ps(tmpilo), scalar);
-        __m128 tmphi = _mm_mul_ps(_mm_cvtepi32_ps(tmpihi), scalar);
-
-        //store to output
-        _mm_storeu_ps(reinterpret_cast<float *>(output+i+0), tmplo);
-        _mm_storeu_ps(reinterpret_cast<float *>(output+i+2), tmphi);
-    }
-
-    //convert remainder
-    for (; i < nsamps; i++){
-        output[i] = item32_to_fc32(uhd::byteswap(input[i]));
-    }
-}
-
-#else
-static UHD_INLINE void item32_to_fc32_bswap(
-    const item32_t *input, fc32_t *output, size_t nsamps
-){
-    for (size_t i = 0; i < nsamps; i++){
-        output[i] = item32_to_fc32(uhd::byteswap(input[i]));
-    }
-}
-
-#endif
-
-#endif /* INCLUDED_LIBUHD_TRANSPORT_CONVERT_TYPES_IMPL_HPP */
diff --git a/host/lib/transport/gen_convert_types.py b/host/lib/transport/gen_convert_types.py
deleted file mode 100755
index f9509c81d..000000000
--- a/host/lib/transport/gen_convert_types.py
+++ /dev/null
@@ -1,211 +0,0 @@
-#!/usr/bin/env python
-#
-# Copyright 2010 Ettus Research LLC
-#
-# This program 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 of the License, or
-# (at your option) any later version.
-#
-# This program 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 this program.  If not, see <http://www.gnu.org/licenses/>.
-#
-
-TMPL_TEXT = """
-#import time
-/***********************************************************************
- * This file was generated by $file on $time.strftime("%c")
- **********************************************************************/
-
-\#include <uhd/config.hpp>
-\#include <uhd/transport/convert_types.hpp>
-\#include <boost/cstdint.hpp>
-\#include <boost/detail/endian.hpp>
-\#include <stdexcept>
-\#include "convert_types_impl.hpp"
-
-using namespace uhd;
-
-/***********************************************************************
- * Generate predicate for jump table
- **********************************************************************/
-UHD_INLINE boost::uint8_t get_pred(
-    const io_type_t &io_type,
-    const otw_type_t &otw_type,
-    size_t num_chans
-){
-    boost::uint8_t pred = 0;
-
-    switch(otw_type.byteorder){
-    \#ifdef BOOST_BIG_ENDIAN
-    case otw_type_t::BO_BIG_ENDIAN:    pred |= $ph.nswap_p; break;
-    case otw_type_t::BO_LITTLE_ENDIAN: pred |= $ph.bswap_p; break;
-    \#else
-    case otw_type_t::BO_BIG_ENDIAN:    pred |= $ph.bswap_p; break;
-    case otw_type_t::BO_LITTLE_ENDIAN: pred |= $ph.nswap_p; break;
-    \#endif
-    case otw_type_t::BO_NATIVE:        pred |= $ph.nswap_p; break;
-    default: throw std::runtime_error("unhandled otw byteorder type");
-    }
-
-    switch(otw_type.get_sample_size()){
-    case sizeof(boost::uint32_t): pred |= $ph.item32_p; break;
-    default: throw std::runtime_error("unhandled otw sample size");
-    }
-
-    switch(io_type.tid){
-    case io_type_t::COMPLEX_FLOAT32: pred |= $ph.fc32_p; break;
-    case io_type_t::COMPLEX_INT16:   pred |= $ph.sc16_p; break;
-    default: throw std::runtime_error("unhandled io type id");
-    }
-
-    switch(num_chans){
-    case 1: pred |= $ph.chan1_p; break;
-    case 2: pred |= $ph.chan2_p; break;
-    case 3: pred |= $ph.chan3_p; break;
-    case 4: pred |= $ph.chan4_p; break;
-    default: throw std::runtime_error("unhandled number of channels");
-    }
-
-    return pred;
-}
-
-/***********************************************************************
- * Convert host type to device type
- **********************************************************************/
-void transport::convert_io_type_to_otw_type(
-    const std::vector<const void *> &io_buffs,
-    const io_type_t &io_type,
-    void *otw_buff,
-    const otw_type_t &otw_type,
-    size_t nsamps_per_io_buff
-){
-    switch(get_pred(io_type, otw_type, io_buffs.size())){
-    #for $pred in range(2**$ph.nbits)
-    case $pred:
-        #set $out_type = $ph.get_dev_type($pred)
-        #set $in_type = $ph.get_host_type($pred)
-        #set $num_chans = $ph.get_num_chans($pred)
-        #set $converter = '_'.join([$in_type, 'to', $out_type])
-        #if $num_chans == 1
-        $(converter)_$ph.get_swap_type($pred)(
-            reinterpret_cast<const $(in_type)_t *>(io_buffs.front()),
-            reinterpret_cast<$(out_type)_t *>(otw_buff),
-            nsamps_per_io_buff
-        );
-        #else
-        for (size_t i = 0, j = 0; i < nsamps_per_io_buff; i++){
-            #for $j in range($num_chans)
-            reinterpret_cast<$(out_type)_t *>(otw_buff)[j++] =
-                #if $ph.get_swap_type($pred) == 'bswap'
-                uhd::byteswap($(converter)(reinterpret_cast<const $(in_type)_t *>(io_buffs[$j])[i]));
-                #else
-                $(converter)(reinterpret_cast<const $(in_type)_t *>(io_buffs[$j])[i]);
-                #end if
-            #end for
-        }
-        #end if
-        break;
-    #end for
-    }
-}
-
-/***********************************************************************
- * Convert device type to host type
- **********************************************************************/
-void transport::convert_otw_type_to_io_type(
-    const void *otw_buff,
-    const otw_type_t &otw_type,
-    std::vector<void *> &io_buffs,
-    const io_type_t &io_type,
-    size_t nsamps_per_io_buff
-){
-    switch(get_pred(io_type, otw_type, io_buffs.size())){
-    #for $pred in range(2**$ph.nbits)
-    case $pred:
-        #set $out_type = $ph.get_host_type($pred)
-        #set $in_type = $ph.get_dev_type($pred)
-        #set $num_chans = $ph.get_num_chans($pred)
-        #set $converter = '_'.join([$in_type, 'to', $out_type])
-        #if $num_chans == 1
-        $(converter)_$ph.get_swap_type($pred)(
-            reinterpret_cast<const $(in_type)_t *>(otw_buff),
-            reinterpret_cast<$(out_type)_t *>(io_buffs.front()),
-            nsamps_per_io_buff
-        );
-        #else
-        for (size_t i = 0, j = 0; i < nsamps_per_io_buff; i++){
-            #for $j in range($num_chans)
-            reinterpret_cast<$(out_type)_t *>(io_buffs[$j])[i] =
-                #if $ph.get_swap_type($pred) == 'bswap'
-                $(converter)(uhd::byteswap(reinterpret_cast<const $(in_type)_t *>(otw_buff)[j++]));
-                #else
-                $(converter)(reinterpret_cast<const $(in_type)_t *>(otw_buff)[j++]);
-                #end if
-            #end for
-        }
-        #end if
-        break;
-    #end for
-    }
-}
-
-"""
-
-def parse_tmpl(_tmpl_text, **kwargs):
-    from Cheetah.Template import Template
-    return str(Template(_tmpl_text, kwargs))
-
-class ph:
-    bswap_p  = 0b00001
-    nswap_p  = 0b00000
-    item32_p = 0b00000
-    sc16_p   = 0b00010
-    fc32_p   = 0b00000
-    chan1_p  = 0b00000
-    chan2_p  = 0b00100
-    chan3_p  = 0b01000
-    chan4_p  = 0b01100
-
-    nbits = 4 #see above
-
-    @staticmethod
-    def has(pred, mask, flag): return (pred & mask) == flag
-
-    @staticmethod
-    def get_swap_type(pred):
-        mask = 0b1
-        if ph.has(pred, mask, ph.bswap_p): return 'bswap'
-        if ph.has(pred, mask, ph.nswap_p): return 'nswap'
-        raise NotImplementedError
-
-    @staticmethod
-    def get_dev_type(pred):
-        mask = 0b0
-        if ph.has(pred, mask, ph.item32_p): return 'item32'
-        raise NotImplementedError
-
-    @staticmethod
-    def get_host_type(pred):
-        mask = 0b10
-        if ph.has(pred, mask, ph.sc16_p): return 'sc16'
-        if ph.has(pred, mask, ph.fc32_p): return 'fc32'
-        raise NotImplementedError
-
-    @staticmethod
-    def get_num_chans(pred):
-        mask = 0b1100
-        if ph.has(pred, mask, ph.chan1_p): return 1
-        if ph.has(pred, mask, ph.chan2_p): return 2
-        if ph.has(pred, mask, ph.chan3_p): return 3
-        if ph.has(pred, mask, ph.chan4_p): return 4
-        raise NotImplementedError
-
-if __name__ == '__main__':
-    import sys
-    open(sys.argv[1], 'w').write(parse_tmpl(TMPL_TEXT, file=__file__, ph=ph))