diff --git a/esp-hal-common/src/pulse_control.rs b/esp-hal-common/src/pulse_control.rs
index 31f7caf1b1d..98d44fff01b 100644
--- a/esp-hal-common/src/pulse_control.rs
+++ b/esp-hal-common/src/pulse_control.rs
@@ -65,9 +65,9 @@
 //! // Create pulse sequence
 //! let mut seq = [PulseCode {
 //!     level1: true,
-//!     length1: 1,
+//!     length1: 10u32.nanos(),
 //!     level2: false,
-//!     length2: 9,
+//!     length2: 90u32.nanos(),
 //! }; 288];
 //!
 //! // Send sequence
@@ -284,13 +284,6 @@ macro_rules! channel_instance {
                             // Configure memory block size
                             w.mem_size()
                                 .bits(1)
-                                // Set config bit
-                                .conf_update()
-                                .set_bit()
-                                // Enable wrap mode (this is enabled globally for
-                                // the ESP32 and ESP32-S2)
-                                .mem_tx_wrap_en()
-                                .set_bit()
                         });
                     }
                     else {
@@ -588,7 +581,7 @@ macro_rules! output_channel {
                     self.write_sequence(&mut sequence_iter, CHANNEL_RAM_SIZE);
                 } else {
                     // Write whole sequence to FIFO RAM
-                    self.write_sequence(&mut sequence_iter, CHANNEL_RAM_SIZE - 1);
+                    self.write_sequence(&mut sequence_iter, CHANNEL_RAM_SIZE);
                 }
 
                 // Clear the relevant interrupts
@@ -639,6 +632,19 @@ macro_rules! output_channel {
                     }
                 }
 
+                // always enable tx wrap
+                #[cfg(any(feature = "esp32c3", feature = "esp32s3"))]
+                unsafe { &*RMT::PTR }.ch_tx_conf0[$num].modify(|_, w| {
+                    w.mem_tx_wrap_en()
+                        .set_bit()
+                });
+
+                // apply configuration updates
+                #[cfg(any(feature = "esp32c3", feature = "esp32s3"))]
+                unsafe { &*RMT::PTR }.ch_tx_conf0[$num].modify(|_, w| {
+                    w.conf_update()
+                        .set_bit()
+                });
 
                 // Depending on the variant, other registers have to be used here
                 cfg_if::cfg_if! {
@@ -929,7 +935,7 @@ macro_rules! rmt {
             self.reg.int_ena.write(|w| unsafe { w.bits(0) });
 
             // Clear all interrupts
-            self.reg.int_clr.write(|w| unsafe { w.bits(0) });
+            self.reg.int_clr.write(|w| unsafe { w.bits(0xffffffff) });
 
             Ok(())
         }
diff --git a/esp32-hal/examples/pulse_control.rs b/esp32-hal/examples/pulse_control.rs
new file mode 100644
index 00000000000..9586b94c053
--- /dev/null
+++ b/esp32-hal/examples/pulse_control.rs
@@ -0,0 +1,78 @@
+//! This demos basic usage of RMT / PulseControl
+//! Use a logic analyzer to see the generated pulses.
+//! The correct output is only achieved when running in release mode.
+
+#![no_std]
+#![no_main]
+
+use esp32_hal::{
+    clock::ClockControl,
+    gpio::IO,
+    pac::Peripherals,
+    prelude::*,
+    pulse_control::{OutputChannel, PulseCode, RepeatMode},
+    timer::TimerGroup,
+    PulseControl,
+    Rtc,
+};
+use panic_halt as _;
+use xtensa_lx_rt::entry;
+
+#[entry]
+fn main() -> ! {
+    let peripherals = Peripherals::take().unwrap();
+    let mut system = peripherals.DPORT.split();
+    let clocks = ClockControl::boot_defaults(system.clock_control).freeze();
+
+    let timer_group0 = TimerGroup::new(peripherals.TIMG0, &clocks);
+    let mut wdt = timer_group0.wdt;
+    let mut rtc = Rtc::new(peripherals.RTC_CNTL);
+
+    // Disable MWDT and RWDT (Watchdog) flash boot protection
+    wdt.disable();
+    rtc.rwdt.disable();
+
+    let io = IO::new(peripherals.GPIO, peripherals.IO_MUX);
+
+    // Configure RMT peripheral globally
+    let pulse = PulseControl::new(peripherals.RMT, &mut system.peripheral_clock_control).unwrap();
+
+    let mut rmt_channel0 = pulse.channel0;
+
+    rmt_channel0
+        .set_idle_output_level(false)
+        .set_carrier_modulation(false)
+        .set_channel_divider(1)
+        .set_idle_output(true);
+
+    // Assign GPIO pin where pulses should be sent to
+    rmt_channel0.assign_pin(io.pins.gpio4);
+
+    // Create pulse sequence
+    let mut seq = [PulseCode {
+        level1: true,
+        length1: 0u32.nanos(),
+        level2: false,
+        length2: 0u32.nanos(),
+    }; 128];
+
+    // -1 to make sure that the last element is a transmission end marker (i.e.
+    // lenght 0)
+    for i in 0..(seq.len() - 1) {
+        seq[i] = PulseCode {
+            level1: true,
+            length1: (10u32 * (i as u32 + 1u32)).nanos(),
+            level2: false,
+            length2: 60u32.nanos(),
+        };
+    }
+
+    esp_println::println!("Start");
+
+    loop {
+        // Send sequence
+        rmt_channel0
+            .send_pulse_sequence(RepeatMode::SingleShot, &seq)
+            .unwrap();
+    }
+}
diff --git a/esp32c3-hal/examples/pulse_control.rs b/esp32c3-hal/examples/pulse_control.rs
new file mode 100644
index 00000000000..9f3d0d282cc
--- /dev/null
+++ b/esp32c3-hal/examples/pulse_control.rs
@@ -0,0 +1,93 @@
+//! This demos basic usage of RMT / PulseControl
+//! Use a logic analyzer to see the generated pulses.
+//! The correct output is only achieved when running in release mode.
+
+#![no_std]
+#![no_main]
+
+use esp32c3_hal::{
+    clock::ClockControl,
+    gpio::IO,
+    pac::Peripherals,
+    prelude::*,
+    pulse_control::{ClockSource, OutputChannel, PulseCode, RepeatMode},
+    system::SystemExt,
+    timer::TimerGroup,
+    PulseControl,
+    Rtc,
+};
+use panic_halt as _;
+use riscv_rt::entry;
+
+#[entry]
+fn main() -> ! {
+    let peripherals = Peripherals::take().unwrap();
+    let mut system = peripherals.SYSTEM.split();
+    let clocks = ClockControl::boot_defaults(system.clock_control).freeze();
+
+    // Disable the watchdog timers. For the ESP32-C3, this includes the Super WDT,
+    // the RTC WDT, and the TIMG WDTs.
+    let mut rtc = Rtc::new(peripherals.RTC_CNTL);
+    let timer_group0 = TimerGroup::new(peripherals.TIMG0, &clocks);
+    let mut wdt0 = timer_group0.wdt;
+    let timer_group1 = TimerGroup::new(peripherals.TIMG1, &clocks);
+    let mut wdt1 = timer_group1.wdt;
+
+    rtc.swd.disable();
+    rtc.rwdt.disable();
+    wdt0.disable();
+    wdt1.disable();
+
+    let io = IO::new(peripherals.GPIO, peripherals.IO_MUX);
+
+    // Configure RMT peripheral globally
+    let pulse = PulseControl::new(
+        peripherals.RMT,
+        &mut system.peripheral_clock_control,
+        ClockSource::APB,
+        0,
+        0,
+        0,
+    )
+    .unwrap();
+
+    let mut rmt_channel0 = pulse.channel0;
+
+    // Set up channel
+    rmt_channel0
+        .set_idle_output_level(false)
+        .set_carrier_modulation(false)
+        .set_channel_divider(1)
+        .set_idle_output(true);
+
+    // Assign GPIO pin where pulses should be sent to
+    rmt_channel0.assign_pin(io.pins.gpio4);
+
+    // Create pulse sequence
+    let mut seq = [PulseCode {
+        level1: true,
+        length1: 0u32.nanos(),
+        level2: false,
+        length2: 0u32.nanos(),
+    }; 128];
+
+    // -1 to make sure that the last element is a transmission end marker (i.e.
+    // lenght 0)
+    for i in 0..(seq.len() - 1) {
+        seq[i] = PulseCode {
+            level1: true,
+            length1: (10u32 * (i as u32 + 1u32)).nanos(),
+            level2: false,
+            length2: 60u32.nanos(),
+        };
+    }
+
+    esp_println::println!("Start");
+
+    loop {
+        // Send sequence
+        rmt_channel0
+            .send_pulse_sequence(RepeatMode::SingleShot, &seq)
+            .unwrap();
+    }
+}
diff --git a/esp32c3-hal/examples/rtc_watchdog.rs b/esp32c3-hal/examples/rtc_watchdog.rs
index a5781617e84..84746f6af09 100644
--- a/esp32c3-hal/examples/rtc_watchdog.rs
+++ b/esp32c3-hal/examples/rtc_watchdog.rs
@@ -1,7 +1,7 @@
 //! This demos the RTC Watchdog Timer (RWDT).
-//! The RWDT is initially configured to trigger an interrupt after a given timeout.
-//! Then, upon expiration, the RWDT is restarted and then reconfigured to reset both the main
-//! system and the RTC.
+//! The RWDT is initially configured to trigger an interrupt after a given
+//! timeout. Then, upon expiration, the RWDT is restarted and then reconfigured
+//! to reset both the main system and the RTC.
 
 #![no_std]
 #![no_main]
@@ -9,13 +9,13 @@
 use core::cell::RefCell;
 
 use bare_metal::Mutex;
-
 use esp32c3_hal::{
     clock::ClockControl,
     interrupt,
     pac::{self, Peripherals},
     prelude::*,
-    Rtc, Rwdt,
+    Rtc,
+    Rwdt,
 };
 use panic_halt as _;
 use riscv_rt::entry;
diff --git a/esp32c3-hal/examples/watchdog.rs b/esp32c3-hal/examples/watchdog.rs
index 3c6ba8d3f51..c2aa552b94f 100644
--- a/esp32c3-hal/examples/watchdog.rs
+++ b/esp32c3-hal/examples/watchdog.rs
@@ -8,7 +8,12 @@
 use core::fmt::Write;
 
 use esp32c3_hal::{
-    clock::ClockControl, pac::Peripherals, prelude::*, timer::TimerGroup, Rtc, Serial,
+    clock::ClockControl,
+    pac::Peripherals,
+    prelude::*,
+    timer::TimerGroup,
+    Rtc,
+    Serial,
 };
 use nb::block;
 use panic_halt as _;
diff --git a/esp32c3-hal/src/lib.rs b/esp32c3-hal/src/lib.rs
index 0f606c5363d..cfcfeafab37 100644
--- a/esp32c3-hal/src/lib.rs
+++ b/esp32c3-hal/src/lib.rs
@@ -4,8 +4,29 @@ use core::arch::global_asm;
 
 pub use embedded_hal as ehal;
 pub use esp_hal_common::{
-    clock, efuse, gpio as gpio_types, i2c, interrupt, ledc, macros, pac, prelude, pulse_control,
-    serial, spi, system, systimer, timer, utils, Cpu, Delay, PulseControl, Rng, Rtc, Rwdt, Serial,
+    clock,
+    efuse,
+    gpio as gpio_types,
+    i2c,
+    interrupt,
+    ledc,
+    macros,
+    pac,
+    prelude,
+    pulse_control,
+    serial,
+    spi,
+    system,
+    systimer,
+    timer,
+    utils,
+    Cpu,
+    Delay,
+    PulseControl,
+    Rng,
+    Rtc,
+    Rwdt,
+    Serial,
     UsbSerialJtag,
 };
 #[cfg(feature = "direct-boot")]
diff --git a/esp32s2-hal/examples/pulse_control.rs b/esp32s2-hal/examples/pulse_control.rs
new file mode 100644
index 00000000000..9baa72d4f3b
--- /dev/null
+++ b/esp32s2-hal/examples/pulse_control.rs
@@ -0,0 +1,79 @@
+//! This demos basic usage of RMT / PulseControl
+//! Use a logic analyzer to see the generated pulses.
+//! The correct output is only achieved when running in release mode.
+
+#![no_std]
+#![no_main]
+
+use esp32s2_hal::{
+    clock::ClockControl,
+    gpio::IO,
+    pac::Peripherals,
+    prelude::*,
+    pulse_control::{OutputChannel, PulseCode, RepeatMode},
+    timer::TimerGroup,
+    PulseControl,
+    Rtc,
+};
+use panic_halt as _;
+use xtensa_lx_rt::entry;
+
+#[entry]
+fn main() -> ! {
+    let peripherals = Peripherals::take().unwrap();
+    let mut system = peripherals.SYSTEM.split();
+    let clocks = ClockControl::boot_defaults(system.clock_control).freeze();
+
+    let timer_group0 = TimerGroup::new(peripherals.TIMG0, &clocks);
+    let mut wdt = timer_group0.wdt;
+    let mut rtc = Rtc::new(peripherals.RTC_CNTL);
+
+    // Disable MWDT and RWDT (Watchdog) flash boot protection
+    wdt.disable();
+    rtc.rwdt.disable();
+
+    let io = IO::new(peripherals.GPIO, peripherals.IO_MUX);
+
+    // Configure RMT peripheral globally
+    let pulse = PulseControl::new(peripherals.RMT, &mut system.peripheral_clock_control).unwrap();
+
+    let mut rmt_channel0 = pulse.channel0;
+
+    // Set up channel
+    rmt_channel0
+        .set_idle_output_level(false)
+        .set_carrier_modulation(false)
+        .set_channel_divider(1)
+        .set_idle_output(true);
+
+    // Assign GPIO pin where pulses should be sent to
+    rmt_channel0.assign_pin(io.pins.gpio4);
+
+    // Create pulse sequence
+    let mut seq = [PulseCode {
+        level1: true,
+        length1: 0u32.nanos(),
+        level2: false,
+        length2: 0u32.nanos(),
+    }; 128];
+
+    // -1 to make sure that the last element is a transmission end marker (i.e.
+    // lenght 0)
+    for i in 0..(seq.len() - 1) {
+        seq[i] = PulseCode {
+            level1: true,
+            length1: (10u32 * (i as u32 + 1u32)).nanos(),
+            level2: false,
+            length2: 60u32.nanos(),
+        };
+    }
+
+    esp_println::println!("Start");
+
+    loop {
+        // Send sequence
+        rmt_channel0
+            .send_pulse_sequence(RepeatMode::SingleShot, &seq)
+            .unwrap();
+    }
+}
diff --git a/esp32s3-hal/examples/pulse_control.rs b/esp32s3-hal/examples/pulse_control.rs
new file mode 100644
index 00000000000..5e77a4f4456
--- /dev/null
+++ b/esp32s3-hal/examples/pulse_control.rs
@@ -0,0 +1,87 @@
+//! This demos basic usage of RMT / PulseControl
+//! Use a logic analyzer to see the generated pulses.
+//! The correct output is only achieved when running in release mode.
+
+#![no_std]
+#![no_main]
+
+use esp32s3_hal::{
+    clock::ClockControl,
+    gpio::IO,
+    pac::Peripherals,
+    prelude::*,
+    pulse_control::{ClockSource, OutputChannel, PulseCode, RepeatMode},
+    timer::TimerGroup,
+    PulseControl,
+    Rtc,
+};
+use panic_halt as _;
+use xtensa_lx_rt::entry;
+
+#[entry]
+fn main() -> ! {
+    let peripherals = Peripherals::take().unwrap();
+    let mut system = peripherals.SYSTEM.split();
+    let clocks = ClockControl::boot_defaults(system.clock_control).freeze();
+
+    let timer_group0 = TimerGroup::new(peripherals.TIMG0, &clocks);
+    let mut wdt = timer_group0.wdt;
+    let mut rtc = Rtc::new(peripherals.RTC_CNTL);
+
+    // Disable MWDT and RWDT (Watchdog) flash boot protection
+    wdt.disable();
+    rtc.rwdt.disable();
+
+    let io = IO::new(peripherals.GPIO, peripherals.IO_MUX);
+
+    // Configure RMT peripheral globally
+    let pulse = PulseControl::new(
+        peripherals.RMT,
+        &mut system.peripheral_clock_control,
+        ClockSource::APB,
+        0,
+        0,
+        0,
+    )
+    .unwrap();
+
+    let mut rmt_channel0 = pulse.channel0;
+
+    // Set up channel
+    rmt_channel0
+        .set_idle_output_level(false)
+        .set_carrier_modulation(false)
+        .set_channel_divider(1)
+        .set_idle_output(true);
+
+    // Assign GPIO pin where pulses should be sent to
+    rmt_channel0.assign_pin(io.pins.gpio4);
+
+    // Create pulse sequence
+    let mut seq = [PulseCode {
+        level1: true,
+        length1: 0u32.nanos(),
+        level2: false,
+        length2: 0u32.nanos(),
+    }; 128];
+
+    // -1 to make sure that the last element is a transmission end marker (i.e.
+    // lenght 0)
+    for i in 0..(seq.len() - 1) {
+        seq[i] = PulseCode {
+            level1: true,
+            length1: (10u32 * (i as u32 + 1u32)).nanos(),
+            level2: false,
+            length2: 60u32.nanos(),
+        };
+    }
+
+    esp_println::println!("Start");
+
+    loop {
+        // Send sequence
+        rmt_channel0
+            .send_pulse_sequence(RepeatMode::SingleShot, &seq)
+            .unwrap();
+    }
+}