Repository: google/jax
Pull request #21634 (2024-06-04 15:54:49)

Moved Pallas GPU ops into pallas/ops/gpu
Moved Pallas GPU ops into pallas/ops/gpu



Patch:
diff --git a/jax/BUILD b/jax/BUILD
index 5097d9ffc..ff5003c28 100644
--- a/jax/BUILD
+++ b/jax/BUILD
@@ -601,6 +601,7 @@ pytype_strict_library(
         exclude = [
             "experimental/pallas/gpu.py",
             "experimental/pallas/tpu.py",
+            "experimental/pallas/ops/gpu/**/*.py",
             "experimental/pallas/ops/tpu/**/*.py",
         ],
     ),
diff --git a/jax/experimental/pallas/ops/__init__.py b/jax/experimental/pallas/ops/__init__.py
index 017372046..132d3839e 100644
--- a/jax/experimental/pallas/ops/__init__.py
+++ b/jax/experimental/pallas/ops/__init__.py
@@ -12,13 +12,6 @@
 # See the License for the specific language governing permissions and
 # limitations under the License.
 
-
-from jax.experimental.pallas.ops import attention
-from jax.experimental.pallas.ops import layer_norm
-from jax.experimental.pallas.ops import rms_norm
-from jax.experimental.pallas.ops import softmax
-
-
 # All files within ops should be treated as user code.
 import os
 import jax._src.source_info_util
diff --git a/jax/experimental/pallas/ops/gpu/__init__.py b/jax/experimental/pallas/ops/gpu/__init__.py
new file mode 100644
index 000000000..862a661e2
--- /dev/null
+++ b/jax/experimental/pallas/ops/gpu/__init__.py
@@ -0,0 +1,13 @@
+# Copyright 2024 The JAX Authors.
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#     https://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.


Test patch:
diff --git a/jax/experimental/pallas/ops/attention.py b/jax/experimental/pallas/ops/gpu/attention.py
similarity index 100%
rename from jax/experimental/pallas/ops/attention.py
rename to jax/experimental/pallas/ops/gpu/attention.py
diff --git a/jax/experimental/pallas/ops/layer_norm.py b/jax/experimental/pallas/ops/gpu/layer_norm.py
similarity index 100%
rename from jax/experimental/pallas/ops/layer_norm.py
rename to jax/experimental/pallas/ops/gpu/layer_norm.py
diff --git a/jax/experimental/pallas/ops/rms_norm.py b/jax/experimental/pallas/ops/gpu/rms_norm.py
similarity index 100%
rename from jax/experimental/pallas/ops/rms_norm.py
rename to jax/experimental/pallas/ops/gpu/rms_norm.py
diff --git a/jax/experimental/pallas/ops/softmax.py b/jax/experimental/pallas/ops/gpu/softmax.py
similarity index 100%
rename from jax/experimental/pallas/ops/softmax.py
rename to jax/experimental/pallas/ops/gpu/softmax.py
diff --git a/tests/pallas/BUILD b/tests/pallas/BUILD
index 441cacecc..69cf99b26 100644
--- a/tests/pallas/BUILD
+++ b/tests/pallas/BUILD
@@ -58,6 +58,7 @@ jax_test(
     deps = [
         "//jax:pallas",
         "//jax:pallas_gpu",
+        "//jax:pallas_gpu_ops",
     ] + py_deps("absl/testing") + py_deps("numpy"),
 )
 
diff --git a/tests/pallas/pallas_test.py b/tests/pallas/pallas_test.py
index a0f32a785..80018a3d7 100644
--- a/tests/pallas/pallas_test.py
+++ b/tests/pallas/pallas_test.py
@@ -32,10 +32,10 @@ from jax._src import test_util as jtu
 from jax._src.lax.control_flow.for_loop import for_loop
 from jax._src.pallas.pallas_call import _trace_to_jaxpr
 from jax.experimental import pallas as pl
-from jax.experimental.pallas.ops import attention
-from jax.experimental.pallas.ops import layer_norm
-from jax.experimental.pallas.ops import rms_norm
-from jax.experimental.pallas.ops import softmax
+from jax.experimental.pallas.ops.gpu import attention
+from jax.experimental.pallas.ops.gpu import layer_norm
+from jax.experimental.pallas.ops.gpu import rms_norm
+from jax.experimental.pallas.ops.gpu import softmax
 from jax.interpreters import partial_eval as pe
 import jax.numpy as jnp
 import numpy as np