diff --git a/CMakeLists.txt b/CMakeLists.txt
index d1dea4432..013dab9ee 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -237,6 +237,11 @@ if(USE_GCC OR USE_CLANG OR USE_INTELCC OR USE_QCC)
   set(OPT_DEF_GCC_ATOMICS ON)
 endif()
 
+if(WIIU)
+  # Prefer coreinit atomics on Wii U due to a hardware bug in load-exclusive and store-exclusive instructions
+  set(OPT_DEF_GCC_ATOMICS OFF)
+endif()
+
 # Default option knobs
 if(UNIX OR MINGW OR MSYS OR (USE_CLANG AND NOT WINDOWS) OR VITA OR PSP OR PS2 OR N3DS OR WIIU)
   set(OPT_DEF_LIBC ON)
diff --git a/include/SDL_atomic.h b/include/SDL_atomic.h
index 1dd816a38..26b5fd2ef 100644
--- a/include/SDL_atomic.h
+++ b/include/SDL_atomic.h
@@ -230,6 +230,10 @@ typedef void (*SDL_KernelMemoryBarrierFunc)();
 #include <mbarrier.h>
 #define SDL_MemoryBarrierRelease()  __machine_rel_barrier()
 #define SDL_MemoryBarrierAcquire()  __machine_acq_barrier()
+#elif defined(__WIIU__)
+#include <coreinit/cache.h>
+#define SDL_MemoryBarrierRelease()  OSMemoryBarrier()
+#define SDL_MemoryBarrierAcquire()  OSMemoryBarrier()
 #else
 /* This is correct for the x86 and x64 CPUs, and we'll expand this over time. */
 #define SDL_MemoryBarrierRelease()  SDL_CompilerBarrier()
@@ -242,7 +246,7 @@ typedef void (*SDL_KernelMemoryBarrierFunc)();
     #define SDL_CPUPauseInstruction() __asm__ __volatile__("pause\n")  /* Some assemblers can't do REP NOP, so go with PAUSE. */
 #elif (defined(__arm__) && defined(__ARM_ARCH) && __ARM_ARCH >= 7) || defined(__aarch64__)
     #define SDL_CPUPauseInstruction() __asm__ __volatile__("yield" ::: "memory")
-#elif (defined(__powerpc__) || defined(__powerpc64__))
+#elif (defined(__powerpc__) || defined(__powerpc64__) || defined(__WIIU__))
     #define SDL_CPUPauseInstruction() __asm__ __volatile__("or 27,27,27");
 #elif defined(_MSC_VER) && (defined(_M_IX86) || defined(_M_X64))
     #define SDL_CPUPauseInstruction() _mm_pause()  /* this is actually "rep nop" and not a SIMD instruction. No inline asm in MSVC x86-64! */
diff --git a/src/atomic/SDL_atomic.c b/src/atomic/SDL_atomic.c
index 4e10d52db..0a3c85dd2 100644
--- a/src/atomic/SDL_atomic.c
+++ b/src/atomic/SDL_atomic.c
@@ -35,6 +35,10 @@
 #include <atomic.h>
 #endif
 
+#if defined(__WIIU__)
+#include <coreinit/atomic.h>
+#endif
+
 /* The __atomic_load_n() intrinsic showed up in different times for different compilers. */
 #if defined(__clang__)
 #if __has_builtin(__atomic_load_n) || defined(HAVE_GCC_ATOMICS)
@@ -45,9 +49,9 @@
 #endif
 #endif
 #elif defined(__GNUC__)
-#if (__GNUC__ >= 5)
-#define HAVE_ATOMIC_LOAD_N 1
-#endif
+#   if (__GNUC__ >= 5) && defined(HAVE_GCC_ATOMICS)
+#     define HAVE_ATOMIC_LOAD_N 1
+#   endif
 #endif
 
 /* *INDENT-OFF* */ /* clang-format off */
@@ -102,7 +106,7 @@ extern __inline int _SDL_xadd_watcom(volatile int *a, int v);
   Contributed by Bob Pendleton, bob@pendleton.com
 */
 
-#if !defined(HAVE_MSC_ATOMICS) && !defined(HAVE_GCC_ATOMICS) && !defined(__MACOSX__) && !defined(__SOLARIS__) && !defined(HAVE_WATCOM_ATOMICS)
+#if !defined(HAVE_MSC_ATOMICS) && !defined(HAVE_GCC_ATOMICS) && !defined(__MACOSX__) && !defined(__SOLARIS__) && !defined(HAVE_WATCOM_ATOMICS) && !defined(__WIIU__)
 #define EMULATE_CAS 1
 #endif
 
@@ -131,6 +135,8 @@ SDL_bool SDL_AtomicCAS(SDL_atomic_t *a, int oldval, int newval)
     return _InterlockedCompareExchange((long *)&a->value, (long)newval, (long)oldval) == (long)oldval;
 #elif defined(HAVE_WATCOM_ATOMICS)
     return (SDL_bool)_SDL_cmpxchg_watcom(&a->value, newval, oldval);
+#elif defined(__WIIU__)
+    return OSCompareAndSwapAtomic((volatile uint32_t *)&a->value, (uint32_t)oldval, (uint32_t)newval);
 #elif defined(HAVE_GCC_ATOMICS)
     return (SDL_bool) __sync_bool_compare_and_swap(&a->value, oldval, newval);
 #elif defined(__MACOSX__)  /* this is deprecated in 10.12 sdk; favor gcc atomics. */
@@ -159,6 +165,8 @@ SDL_bool SDL_AtomicCASPtr(void **a, void *oldval, void *newval)
     return _InterlockedCompareExchangePointer(a, newval, oldval) == oldval;
 #elif defined(HAVE_WATCOM_ATOMICS)
     return (SDL_bool)_SDL_cmpxchg_watcom((int *)a, (long)newval, (long)oldval);
+#elif defined(__WIIU__)
+    return OSCompareAndSwapAtomic((volatile uint32_t *)a, (uint32_t)oldval, (uint32_t)newval);
 #elif defined(HAVE_GCC_ATOMICS)
     return __sync_bool_compare_and_swap(a, oldval, newval);
 #elif defined(__MACOSX__) && defined(__LP64__)  /* this is deprecated in 10.12 sdk; favor gcc atomics. */
@@ -190,6 +198,8 @@ int SDL_AtomicSet(SDL_atomic_t *a, int v)
     return _InterlockedExchange((long *)&a->value, v);
 #elif defined(HAVE_WATCOM_ATOMICS)
     return _SDL_xchg_watcom(&a->value, v);
+#elif defined(__WIIU__)
+    return (int) OSSwapAtomic((volatile uint32_t *)&a->value, (uint32_t)v);
 #elif defined(HAVE_GCC_ATOMICS)
     return __sync_lock_test_and_set(&a->value, v);
 #elif defined(__SOLARIS__)
@@ -209,6 +219,8 @@ void *SDL_AtomicSetPtr(void **a, void *v)
     return _InterlockedExchangePointer(a, v);
 #elif defined(HAVE_WATCOM_ATOMICS)
     return (void *)_SDL_xchg_watcom((int *)a, (long)v);
+#elif defined(__WIIU__)
+    return (void *) OSSwapAtomic((volatile uint32_t *)a, (uint32_t)v);
 #elif defined(HAVE_GCC_ATOMICS)
     return __sync_lock_test_and_set(a, v);
 #elif defined(__SOLARIS__)
@@ -229,6 +241,8 @@ int SDL_AtomicAdd(SDL_atomic_t *a, int v)
     return _InterlockedExchangeAdd((long *)&a->value, v);
 #elif defined(HAVE_WATCOM_ATOMICS)
     return _SDL_xadd_watcom(&a->value, v);
+#elif defined(__WIIU__)
+    return OSAddAtomic((volatile int32_t *)&a->value, v);
 #elif defined(HAVE_GCC_ATOMICS)
     return __sync_fetch_and_add(&a->value, v);
 #elif defined(__SOLARIS__)
diff --git a/src/atomic/SDL_spinlock.c b/src/atomic/SDL_spinlock.c
index 420d49623..4fdc1024c 100644
--- a/src/atomic/SDL_spinlock.c
+++ b/src/atomic/SDL_spinlock.c
@@ -37,7 +37,8 @@
 #endif
 
 #if defined(__WIIU__)
-#include <stdatomic.h>
+#include <coreinit/cache.h>
+#include <coreinit/atomic.h>
 #endif
 
 #if defined(_MSC_VER) && (defined(_M_IX86) || defined(_M_X64))
@@ -168,8 +169,7 @@ SDL_bool SDL_AtomicTryLock(SDL_SpinLock *lock)
     }
     return res;
 #elif defined(__WIIU__)
-    uint64_t val = 0;
-    return (SDL_bool) atomic_compare_exchange_strong((volatile _Atomic uint64_t*)lock, &val, 1);
+    return OSCompareAndSwapAtomic((volatile uint32_t *)lock, 0u, 1u);
 
 #else
 #error Please implement for your platform.
@@ -213,6 +213,10 @@ void SDL_AtomicUnlock(SDL_SpinLock *lock)
     *lock = 0;
     membar_producer();
 
+#elif defined(__WIIU__)
+    *lock = 0;
+    OSMemoryBarrier();
+
 #else
     *lock = 0;
 #endif