diff --git a/CHANGELOG.md b/CHANGELOG.md
index db7f529e..cd5f1a13 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -9,7 +9,7 @@ CHANGELOG
 * Pdext: migrate to `pd-lib-builder` and newer `m_pd.h` and add some initial documentation.
 * deprecated: Fabric generator - no longer supported
 * docs: updates on missing objects and limitations
-* bugfix: issues #24, #50, #106
+* bugfix: issues #24, #50, #100, #106
 
 0.7.0
 -----
diff --git a/docs/09.supported_vanilla_objects.md b/docs/09.supported_vanilla_objects.md
index e6281a57..e3f09221 100644
--- a/docs/09.supported_vanilla_objects.md
+++ b/docs/09.supported_vanilla_objects.md
@@ -171,6 +171,7 @@ snapshot~
 sqrt~
 s~
 tabosc4~
+tabplay~
 tabread4~
 tabread~
 tabwrite~
diff --git a/docs/10.unsupported_vanilla_objects.md b/docs/10.unsupported_vanilla_objects.md
index f3d06ca6..393b5b39 100644
--- a/docs/10.unsupported_vanilla_objects.md
+++ b/docs/10.unsupported_vanilla_objects.md
@@ -92,7 +92,6 @@ rifft~
 scope~
 sigmund~
 slop~
-tabplay~
 tabreceive~
 tabsend~
 threshold~
diff --git a/hvcc/core/hv2ir/HIrTabread.py b/hvcc/core/hv2ir/HIrTabread.py
index 93a229ab..2969d90f 100644
--- a/hvcc/core/hv2ir/HIrTabread.py
+++ b/hvcc/core/hv2ir/HIrTabread.py
@@ -31,7 +31,7 @@ def __init__(
         graph: Optional[HeavyGraph] = None,
         annotations: Optional[Dict] = None
     ) -> None:
-        assert obj_type in {"__tabread~if", "__tabread~f", "__tabreadu~f", "__tabread"}
+        assert obj_type in {"__tabread~if", "__tabread~f", "__tabread_stoppable~f", "__tabreadu~f", "__tabread"}
         super().__init__(obj_type, args=args, graph=graph, annotations=annotations)
 
     def reduce(self) -> Optional[tuple]:
diff --git a/hvcc/core/hv2ir/HeavyParser.py b/hvcc/core/hv2ir/HeavyParser.py
index 2600cad5..180de62d 100644
--- a/hvcc/core/hv2ir/HeavyParser.py
+++ b/hvcc/core/hv2ir/HeavyParser.py
@@ -298,6 +298,7 @@ def reduce(self) -> tuple:
     "slice": HLangSlice,
     "__tabread~if": HIrTabread,
     "__tabread~f": HIrTabread,
+    "__tabread_stoppable~f": HIrTabread,
     "__tabreadu~f": HIrTabread,
     "__tabread": HIrTabread,
     "__tabhead~f": HIrTabhead,
diff --git a/hvcc/core/json/heavy.ir.json b/hvcc/core/json/heavy.ir.json
index cdd11435..ff36a44d 100644
--- a/hvcc/core/json/heavy.ir.json
+++ b/hvcc/core/json/heavy.ir.json
@@ -2792,6 +2792,34 @@
       "sse": 1
     }
   },
+  "__tabread_stoppable~f": {
+    "inlets": [
+      "-->",
+      "-->",
+      "-->"
+    ],
+    "ir": {
+      "control": true,
+      "signal": true,
+      "init": true
+    },
+    "outlets": [
+      "~f>",
+      "-->",
+      "-->"
+    ],
+    "args": [{
+      "default": null,
+      "value_type": "string",
+      "name": "table",
+      "description": "The name of the table to reference.",
+      "required": true
+    }],
+    "perf": {
+      "avx": 1,
+      "sse": 1
+    }
+  },
   "__tabread~if": {
     "inlets": [
       "~i>",
diff --git a/hvcc/generators/ir2c/SignalTabread.py b/hvcc/generators/ir2c/SignalTabread.py
index 9bf6b272..e9a1cfe7 100644
--- a/hvcc/generators/ir2c/SignalTabread.py
+++ b/hvcc/generators/ir2c/SignalTabread.py
@@ -20,7 +20,7 @@
 
 
 class SignalTabread(HeavyObject):
-    """Handles __tabread~if, __tabread~f, __tabreadu~f
+    """Handles __tabread~if, __tabread~f, __tabreadu~f, __tabread_stoppable~f
     """
 
     c_struct = "SignalTabread"
@@ -44,11 +44,11 @@ def get_C_init(cls, obj_type: str, obj_id: int, args: Dict) -> List[str]:
             "sTabread_init(&sTabread_{0}, &hTable_{1}, {2});".format(
                 obj_id,
                 args["table_id"],
-                "true" if obj_type == "__tabread~f" else "false")]
+                "true" if obj_type in {"__tabread~f", "__tabread_stoppable~f"} else "false")]
 
     @classmethod
     def get_C_onMessage(cls, obj_type: str, obj_id: int, inlet_index: int, args: Dict) -> List[str]:
-        if obj_type in ["__tabread~f", "__tabreadu~f"]:
+        if obj_type in ["__tabread~f", "__tabreadu~f", "__tabread_stoppable~f"]:
             return [
                 "sTabread_onMessage(_c, &Context(_c)->sTabread_{0}, {1}, m, &sTabread_{0}_sendMessage);".format(
                     obj_id,
@@ -71,6 +71,13 @@ def get_C_process(cls, process_dict: Dict, obj_type: str, obj_id: int, args: Dic
                     process_dict["id"],
                     ", ".join([f"VOf({cls._c_buffer(b)})" for b in process_dict["outputBuffers"]])
                 )]
+        elif obj_type == "__tabread_stoppable~f":
+            return [
+                "__hv_tabread_stoppable_f(this, &sTabread_{0}, {1}, &{2}_{0}_sendMessage);".format(
+                    process_dict["id"],
+                    ", ".join([f"VOf({cls._c_buffer(b)})" for b in process_dict["outputBuffers"]]),
+                    cls.preamble
+                )]
         elif obj_type == "__tabreadu~f":
             return [
                 "__hv_tabreadu_f(&sTabread_{0}, {1});".format(
diff --git a/hvcc/generators/ir2c/ir2c.py b/hvcc/generators/ir2c/ir2c.py
index 5671c65a..fb23277c 100644
--- a/hvcc/generators/ir2c/ir2c.py
+++ b/hvcc/generators/ir2c/ir2c.py
@@ -92,6 +92,7 @@ class ir2c:
         "__tabread~if": SignalTabread,
         "__tabread~f": SignalTabread,
         "__tabreadu~f": SignalTabread,
+        "__tabread_stoppable~f": SignalTabread,
         "__tabhead~f": SignalTabhead,
         "__tabwrite~f": SignalTabwrite,
         "__tabwrite_stoppable~f": SignalTabwrite,
diff --git a/hvcc/generators/ir2c/static/HvSignalTabread.c b/hvcc/generators/ir2c/static/HvSignalTabread.c
index 9197cf9b..f07f9eb9 100644
--- a/hvcc/generators/ir2c/static/HvSignalTabread.c
+++ b/hvcc/generators/ir2c/static/HvSignalTabread.c
@@ -19,7 +19,9 @@
 hv_size_t sTabread_init(SignalTabread *o, HvTable *table, bool forceAlignedLoads) {
   o->table = table;
   o->head = 0;
+  o->end = hTable_getSize(o->table);
   o->forceAlignedLoads = forceAlignedLoads;
+  o->playing = false;
   return 0;
 }
 
@@ -28,8 +30,13 @@ void sTabread_onMessage(HeavyContextInterface *_c, SignalTabread *o, int letIn,
   switch (letIn) {
     case 0: {
       if (o->table != NULL) {
+        o->end = hTable_getSize(o->table);
         switch (msg_getType(m,0)) {
-          case HV_MSG_BANG: o->head = 0; break;
+          case HV_MSG_BANG: {
+            o->head = 0;
+            o->playing = true;
+            break;
+          }
           case HV_MSG_FLOAT: {
             hv_uint32_t h = (hv_uint32_t) hv_abs_f(msg_getFloat(m,0));
             if (msg_getFloat(m,0) < 0.0f) {
@@ -37,6 +44,7 @@ void sTabread_onMessage(HeavyContextInterface *_c, SignalTabread *o, int letIn,
               h = hTable_getSize(o->table) - h;
             }
             o->head = o->forceAlignedLoads ? (h & ~HV_N_SIMD_MASK) : h;
+            o->playing = true;
 
             // output new head
             HvMessage *n = HV_MESSAGE_ON_STACK(1);
@@ -44,6 +52,13 @@ void sTabread_onMessage(HeavyContextInterface *_c, SignalTabread *o, int letIn,
             sendMessage(_c, 1, n);
             break;
           }
+          case HV_MSG_SYMBOL: {
+            if (msg_compareSymbol(m, 0, "stop")) {
+              o->head = 0;
+              o->playing = false;
+            }
+            break;
+          }
           default: break;
         }
       }
@@ -52,6 +67,21 @@ void sTabread_onMessage(HeavyContextInterface *_c, SignalTabread *o, int letIn,
     case 1: {
       if (msg_isHashLike(m,0)) {
         o->table = hv_table_get(_c, msg_getHash(m,0));
+        o->head = 0;
+        o->end = hTable_getSize(o->table);
+      }
+      break;
+    }
+    case 2: {
+      if (o->table != NULL) {
+        switch (msg_getType(m,0)) {
+          case HV_MSG_FLOAT: {
+            hv_uint32_t e = (hv_uint32_t) hv_abs_f(msg_getFloat(m,0));
+            o->end = o->head + e;
+            break;
+          }
+          default: break;
+        }
       }
       break;
     }
diff --git a/hvcc/generators/ir2c/static/HvSignalTabread.h b/hvcc/generators/ir2c/static/HvSignalTabread.h
index a3d330c0..62f50c12 100644
--- a/hvcc/generators/ir2c/static/HvSignalTabread.h
+++ b/hvcc/generators/ir2c/static/HvSignalTabread.h
@@ -26,7 +26,9 @@ extern "C" {
 typedef struct SignalTabread {
   HvTable *table; // the table to read
   hv_uint32_t head;
+  hv_uint32_t end;
   bool forceAlignedLoads; // false by default, true if using __hv_tabread_f
+  bool playing; // false by default, only __hv_tabread_stoppable_f may respond to it
 } SignalTabread;
 
 // random access to a table
@@ -43,30 +45,30 @@ static inline void __hv_tabread_if(SignalTabread *o, hv_bIni_t bIn, hv_bOutf_t b
 #if HV_SIMD_AVX
   const hv_int32_t *const i = (hv_int32_t *) &bIn;
 
-  hv_assert(i[0] >= 0 && i[0] < hTable_getAllocated(o->table));
-  hv_assert(i[1] >= 0 && i[1] < hTable_getAllocated(o->table));
-  hv_assert(i[2] >= 0 && i[2] < hTable_getAllocated(o->table));
-  hv_assert(i[3] >= 0 && i[3] < hTable_getAllocated(o->table));
-  hv_assert(i[4] >= 0 && i[4] < hTable_getAllocated(o->table));
-  hv_assert(i[5] >= 0 && i[5] < hTable_getAllocated(o->table));
-  hv_assert(i[6] >= 0 && i[6] < hTable_getAllocated(o->table));
-  hv_assert(i[7] >= 0 && i[7] < hTable_getAllocated(o->table));
+  hv_assert(i[0] >= 0 && (hv_uint32_t) i[0] < hTable_getAllocated(o->table));
+  hv_assert(i[1] >= 0 && (hv_uint32_t) i[1] < hTable_getAllocated(o->table));
+  hv_assert(i[2] >= 0 && (hv_uint32_t) i[2] < hTable_getAllocated(o->table));
+  hv_assert(i[3] >= 0 && (hv_uint32_t) i[3] < hTable_getAllocated(o->table));
+  hv_assert(i[4] >= 0 && (hv_uint32_t) i[4] < hTable_getAllocated(o->table));
+  hv_assert(i[5] >= 0 && (hv_uint32_t) i[5] < hTable_getAllocated(o->table));
+  hv_assert(i[6] >= 0 && (hv_uint32_t) i[6] < hTable_getAllocated(o->table));
+  hv_assert(i[7] >= 0 && (hv_uint32_t) i[7] < hTable_getAllocated(o->table));
 
   *bOut = _mm256_set_ps(b[i[7]], b[i[6]], b[i[5]], b[i[4]], b[i[3]], b[i[2]], b[i[1]], b[i[0]]);
 #elif HV_SIMD_SSE
   const hv_int32_t *const i = (hv_int32_t *) &bIn;
 
-  hv_assert(i[0] >= 0 && ((hv_uint32_t) i[0]) < hTable_getAllocated(o->table));
-  hv_assert(i[1] >= 0 && ((hv_uint32_t) i[1]) < hTable_getAllocated(o->table));
-  hv_assert(i[2] >= 0 && ((hv_uint32_t) i[2]) < hTable_getAllocated(o->table));
-  hv_assert(i[3] >= 0 && ((hv_uint32_t) i[3]) < hTable_getAllocated(o->table));
+  hv_assert(i[0] >= 0 && (hv_uint32_t) i[0] < hTable_getAllocated(o->table));
+  hv_assert(i[1] >= 0 && (hv_uint32_t) i[1] < hTable_getAllocated(o->table));
+  hv_assert(i[2] >= 0 && (hv_uint32_t) i[2] < hTable_getAllocated(o->table));
+  hv_assert(i[3] >= 0 && (hv_uint32_t) i[3] < hTable_getAllocated(o->table));
 
   *bOut = _mm_set_ps(b[i[3]], b[i[2]], b[i[1]], b[i[0]]);
 #elif HV_SIMD_NEON
-  hv_assert((bIn[0] >= 0) && (bIn[0] < hTable_getAllocated(o->table)));
-  hv_assert((bIn[1] >= 0) && (bIn[1] < hTable_getAllocated(o->table)));
-  hv_assert((bIn[2] >= 0) && (bIn[2] < hTable_getAllocated(o->table)));
-  hv_assert((bIn[3] >= 0) && (bIn[3] < hTable_getAllocated(o->table)));
+  hv_assert((bIn[0] >= 0) && (hv_uint32_t) bIn[0] < hTable_getAllocated(o->table));
+  hv_assert((bIn[1] >= 0) && (hv_uint32_t) bIn[1] < hTable_getAllocated(o->table));
+  hv_assert((bIn[2] >= 0) && (hv_uint32_t) bIn[2] < hTable_getAllocated(o->table));
+  hv_assert((bIn[3] >= 0) && (hv_uint32_t) bIn[3] < hTable_getAllocated(o->table));
 
   *bOut = (float32x4_t) {b[bIn[0]], b[bIn[1]], b[bIn[2]], b[bIn[3]]};
 #else // HV_SIMD_NONE
@@ -115,36 +117,35 @@ static inline void __hv_tabreadu_f(SignalTabread *o, hv_bOutf_t bOut) {
 }
 
 // this tabread can be instructed to stop. It is mainly intended for linear reads that only process a portion of a buffer.
-static inline void __hv_tabread_stoppable_f(SignalTabread *o, hv_bOutf_t bOut) {
-#if HV_SIMD_AVX
-  if (o->head == ~0x0) {
-    *bOut = _mm256_setzero_ps();
-  } else {
-    *bOut = _mm256_load_ps(hTable_getBuffer(o->table) + o->head);
-    o->head += HV_N_SIMD;
+static inline void __hv_tabread_stoppable_f(HeavyContextInterface *_c, SignalTabread *o, hv_bOutf_t bOut,
+  void (*sendMessage)(HeavyContextInterface *, int, const HvMessage *))
+{
+  hv_uint32_t head = o->head;
+  if (head >= hTable_getAllocated(o->table)) { // stop when we reach the table bounds
+    o->playing = false;
   }
-#elif HV_SIMD_SSE
-  if (o->head == ~0x0) {
-    *bOut = _mm_setzero_ps();
+  if (!o->playing) {
+    __hv_zero_f(bOut); // output silence when not playing
   } else {
-    *bOut = _mm_load_ps(hTable_getBuffer(o->table) + o->head);
-    o->head += HV_N_SIMD;
-  }
+    if (o->end > o->head) { // only play when end is further than play head
+#if HV_SIMD_AVX
+      *bOut = _mm256_loadu_ps(hTable_getBuffer(o->table) + head);
+#elif HV_SIMD_SSE
+      *bOut = _mm_load_ps(hTable_getBuffer(o->table) + head);
 #elif HV_SIMD_NEON
-  if (o->head == ~0x0) {
-    *bOut = vdupq_n_f32(0.0f);
-  } else {
-    *bOut = vld1q_f32(hTable_getBuffer(o->table) + o->head);
-    o->head += HV_N_SIMD;
-  }
+      *bOut = vld1q_f32(hTable_getBuffer(o->table) + head);
 #else // HV_SIMD_NONE
-  if (o->head == ~0x0) {
-    *bOut = 0.0f;
-  } else {
-    *bOut = *(hTable_getBuffer(o->table) + o->head);
-    o->head += HV_N_SIMD;
-  }
+      *bOut = *(hTable_getBuffer(o->table) + head);
 #endif
+      o->head = head + HV_N_SIMD;
+    } else {
+      o->playing = false;
+      __hv_zero_f(bOut); // output silence when not playing
+      HvMessage *n = HV_MESSAGE_ON_STACK(1);
+      msg_initWithBang(n, 0);
+      sendMessage(_c, 2, n);
+    }
+  }
 }
 
 void sTabread_onMessage(HeavyContextInterface *_c, SignalTabread *o, int letIn, const HvMessage *m,
diff --git a/hvcc/generators/ir2c/static/HvSignalTabwrite.h b/hvcc/generators/ir2c/static/HvSignalTabwrite.h
index cadf1e01..dfbce2ec 100644
--- a/hvcc/generators/ir2c/static/HvSignalTabwrite.h
+++ b/hvcc/generators/ir2c/static/HvSignalTabwrite.h
@@ -70,7 +70,7 @@ static inline void __hv_tabwriteu_f(SignalTabwrite *o, hv_bInf_t bIn) {
   hTable_setHead(o->table, head); // update remote write head
 }
 
-// this tabread can be instructed to stop. It is mainly intended for linear reads that only process a portion of a buffer.
+// this tabwrite can be instructed to stop. It is mainly intended for linear reads that only process a portion of a buffer.
 // Stores are unaligned, which can be slow but allows any indicies to be written to.
 // TODO(mhroth): this is not stopping!
 static inline void __hv_tabwrite_stoppable_f(SignalTabwrite *o, hv_bInf_t bIn) {
diff --git a/hvcc/interpreters/pd2hv/libs/pd/tabplay~.pd b/hvcc/interpreters/pd2hv/libs/pd/tabplay~.pd
new file mode 100644
index 00000000..45fba5c2
--- /dev/null
+++ b/hvcc/interpreters/pd2hv/libs/pd/tabplay~.pd
@@ -0,0 +1,69 @@
+#N canvas 303 459 461 431 10;
+#X obj 41 351 outlet~;
+#X text 38 375 @hv_arg \$1 table string "" true;
+#N canvas 0 22 450 300 @hv_obj 0;
+#X obj 146 77 inlet;
+#X obj 147 120 outlet;
+#X restore 112 150 pd @hv_obj system;
+#X msg 112 128 table \$1 size;
+#N canvas 0 22 450 300 @hv_obj 0;
+#X obj 178 47 inlet;
+#X obj 165 128 outlet;
+#X restore 43 55 pd @hv_obj slice 1 1;
+#N canvas 0 22 450 300 @hv_obj 0;
+#X obj 150 61 inlet;
+#X obj 147 118 outlet;
+#X restore 112 106 pd @hv_obj __var \$1;
+#X obj 112 174 - 1;
+#X obj 112 80 loadbang -1;
+#N canvas 0 23 450 300 @hv_obj 1;
+#X obj 169 185 outlet~;
+#X obj 249 48 inlet;
+#X obj 165 46 inlet;
+#X obj 337 48 inlet;
+#X obj 306 185 outlet;
+#X obj 374 185 outlet;
+#X restore 41 327 pd @hv_obj __tabread_stoppable~f \$1;
+#X obj 42 7 inlet -->;
+#X obj 41 270 max 0;
+#X obj 41 297 min 0;
+#X msg 253 131 \$1;
+#X obj 248 358 outlet;
+#N canvas 0 22 450 300 @hv_obj 0;
+#X obj 169 24 inlet;
+#X obj 273 82 outlet;
+#X obj 159 83 outlet;
+#X obj 333 82 outlet;
+#X obj 212 82 outlet;
+#X restore 42 34 pd @hv_obj __switchcase set bang stop;
+#X obj 41 221 loadbang 0;
+#N canvas 1361 151 307 222 @hv_obj 0;
+#X obj 37 37 inlet;
+#X obj 40 90 outlet;
+#X restore 41 244 pd @hv_obj __var \$1;
+#N canvas 0 22 450 300 @hv_obj 0;
+#X obj 178 47 inlet;
+#X obj 165 128 outlet;
+#X restore 287 132 pd @hv_obj slice 1 2;
+#X connect 2 0 6 0;
+#X connect 3 0 2 0;
+#X connect 4 0 8 1;
+#X connect 4 0 5 0;
+#X connect 5 0 3 0;
+#X connect 6 0 11 1;
+#X connect 7 0 5 0;
+#X connect 8 0 0 0;
+#X connect 8 2 13 0;
+#X connect 9 0 14 0;
+#X connect 10 0 11 0;
+#X connect 11 0 8 0;
+#X connect 12 0 16 0;
+#X connect 14 0 4 0;
+#X connect 14 1 16 0;
+#X connect 14 1 8 0;
+#X connect 14 2 8 0;
+#X connect 14 3 12 0;
+#X connect 14 3 17 0;
+#X connect 15 0 16 0;
+#X connect 16 0 10 0;
+#X connect 17 0 8 2;