Add SMEM as a supported Pallas output memory space.

PiperOrigin-RevId: 712144883

Author: Google-ML-Automation

jax/_src/pallas/mosaic/pallas_call_registration.py

@@ -84,6 +84,8 @@ def _get_memory_space_from_aval(
       return None
     case tpu_core.TPUMemorySpace.VMEM:
       return tpu_custom_call.MemorySpace.VMEM
+    case tpu_core.TPUMemorySpace.SMEM:
+      return tpu_custom_call.MemorySpace.SMEM
     case tpu_core.TPUMemorySpace.SEMAPHORE:
       return tpu_custom_call.MemorySpace.SEMAPHORE_MEM
   return None

jax/_src/tpu_custom_call.py

@@ -83,6 +83,7 @@ class MemorySpace(enum.Enum):
   HBM = enum.auto()
   VMEM = enum.auto()
   SEMAPHORE_MEM = enum.auto()
+  SMEM = enum.auto()
 
   @property
   def color(self) -> int:
@@ -92,6 +93,8 @@ def color(self) -> int:
       return 1
     elif self == MemorySpace.SEMAPHORE_MEM:
       return 2
+    elif self == MemorySpace.SMEM:
+      return 4
     else:
       raise ValueError("invalid memory space: " + str(self))