Skip to content

Commit 001fd5a

Browse files
whitneywhtsangwhitneywhtsang
authored
Reland disable vectorization for atomic_cas on all backends (#5024)
Fixes #4905 - Reland origin commit 2e1d322 - Skip tests for 16-bit values - Disable vectorization for atomic_cas on xpu backend
2 parents 386e419 + 1d456fd commit 001fd5a

File tree

11 files changed

+248
-99
lines changed

11 files changed

+248
-99
lines changed

‎python/test/unit/language/test_core.py‎

Lines changed: 18 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -1910,24 +1910,33 @@ def serialized_add(data, Lock, SEM: tl.constexpr):
19101910

19111911

19121912
@pytest.mark.interpreter
1913-
@pytest.mark.parametrize("sem", [None, 'acquire', 'release', 'acq_rel', 'relaxed'])
1913+
@pytest.mark.parametrize("sem", [None, "acquire", "release", "acq_rel", "relaxed"])
19141914
@pytest.mark.parametrize("num_ctas", num_ctas_list)
1915-
def test_tensor_atomic_cas(sem, num_ctas, device):
1915+
@pytest.mark.parametrize("size", [4, 128, 512])
1916+
@pytest.mark.parametrize("dtype_str", ['bfloat16', 'float16', 'float32', 'uint64', 'int64', 'float64'])
1917+
def test_tensor_atomic_cas(sem, size, dtype_str, num_ctas, device):
1918+
check_type_supported(dtype_str, device)
1919+
if "float" in dtype_str and is_hip():
1920+
pytest.skip("HIP does not support atomic cas with float types")
19161921

19171922
@triton.jit
1918-
def change_value(X, BLOCK_SIZE: tl.constexpr, sem: tl.constexpr):
1923+
def change_value(X, BLOCK_SIZE: tl.constexpr, sem: tl.constexpr, dtype: tl.constexpr):
19191924
pid = tl.program_id(axis=0)
19201925
block_start = pid * BLOCK_SIZE
19211926
offsets = block_start + tl.arange(0, BLOCK_SIZE)
1922-
t1 = tl.full((BLOCK_SIZE, ), 0, dtype=tl.int64)
1923-
t2 = tl.full((BLOCK_SIZE, ), 2, dtype=tl.int64)
1927+
t1 = tl.full((BLOCK_SIZE, ), 0, dtype=dtype)
1928+
t2 = tl.full((BLOCK_SIZE, ), 2, dtype=dtype)
19241929
tl.atomic_cas(X + offsets, t1, t2, sem=sem)
19251930

1926-
X = torch.tensor([0, 1, 0, 1, 0, 1, 0, 1], device=device, dtype=torch.int64)
1927-
Y = torch.tensor([2, 1, 2, 1, 2, 1, 2, 1], device=device, dtype=torch.int64)
1931+
torch_dtype = getattr(torch, dtype_str)
1932+
X = torch.zeros((size, ), device=device, dtype=torch_dtype)
1933+
X[1::2] = 1
1934+
Y = X.clone()
1935+
Y[0::2] = 2
19281936

1929-
change_value[(2, )](X, 4, sem)
1930-
assert (torch.equal(X, Y))
1937+
tl_dtype = getattr(tl, dtype_str)
1938+
change_value[(2, )](X, BLOCK_SIZE=size // 2, sem=sem, dtype=tl_dtype)
1939+
assert torch.equal(X, Y)
19311940

19321941

19331942
@pytest.mark.interpreter

‎scripts/skiplist/a770/language.txt‎

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -755,3 +755,33 @@ python/test/unit/language/test_core.py::test_tensor_atomic_add_access_patterns[s
755755
python/test/unit/language/test_core.py::test_tensor_atomic_use_result[cas-1-float16]
756756
python/test/unit/language/test_core.py::test_tensor_atomic_use_result[cas-4-float16]
757757
python/test/unit/language/test_core.py::test_tensor_atomic_use_result[cas-16-float16]
758+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-None]
759+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-acquire]
760+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-release]
761+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-acq_rel]
762+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-relaxed]
763+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-None]
764+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-acquire]
765+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-release]
766+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-acq_rel]
767+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-relaxed]
768+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-None]
769+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-acquire]
770+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-release]
771+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-acq_rel]
772+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-relaxed]
773+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-None]
774+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-acquire]
775+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-release]
776+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-acq_rel]
777+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-relaxed]
778+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-None]
779+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-acquire]
780+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-release]
781+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-acq_rel]
782+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-relaxed]
783+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-None]
784+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-acquire]
785+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-release]
786+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-acq_rel]
787+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-relaxed]

‎scripts/skiplist/arl-h/language.txt‎

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -613,3 +613,33 @@ python/test/unit/language/test_tensor_descriptor.py::test_tensor_descriptor_redu
613613
python/test/unit/language/test_core.py::test_tensor_atomic_use_result[cas-1-float16]
614614
python/test/unit/language/test_core.py::test_tensor_atomic_use_result[cas-4-float16]
615615
python/test/unit/language/test_core.py::test_tensor_atomic_use_result[cas-16-float16]
616+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-None]
617+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-acquire]
618+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-release]
619+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-acq_rel]
620+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-relaxed]
621+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-None]
622+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-acquire]
623+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-release]
624+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-acq_rel]
625+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-relaxed]
626+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-None]
627+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-acquire]
628+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-release]
629+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-acq_rel]
630+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-relaxed]
631+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-None]
632+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-acquire]
633+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-release]
634+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-acq_rel]
635+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-relaxed]
636+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-None]
637+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-acquire]
638+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-release]
639+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-acq_rel]
640+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-relaxed]
641+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-None]
642+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-acquire]
643+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-release]
644+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-acq_rel]
645+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-relaxed]

‎scripts/skiplist/conda/language.txt‎

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -239,3 +239,33 @@ python/test/unit/language/test_core.py::test_constexpr_if_return
239239
python/test/unit/language/test_core.py::test_tensor_atomic_use_result[cas-1-float16]
240240
python/test/unit/language/test_core.py::test_tensor_atomic_use_result[cas-4-float16]
241241
python/test/unit/language/test_core.py::test_tensor_atomic_use_result[cas-16-float16]
242+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-None]
243+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-acquire]
244+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-release]
245+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-acq_rel]
246+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-relaxed]
247+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-None]
248+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-acquire]
249+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-release]
250+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-acq_rel]
251+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-relaxed]
252+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-None]
253+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-acquire]
254+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-release]
255+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-acq_rel]
256+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-relaxed]
257+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-None]
258+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-acquire]
259+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-release]
260+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-acq_rel]
261+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-relaxed]
262+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-None]
263+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-acquire]
264+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-release]
265+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-acq_rel]
266+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-relaxed]
267+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-None]
268+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-acquire]
269+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-release]
270+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-acq_rel]
271+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-relaxed]

‎scripts/skiplist/default/language.txt‎

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -90,3 +90,33 @@ python/test/unit/language/test_tensor_descriptor.py::test_tensor_descriptor_redu
9090
python/test/unit/language/test_core.py::test_tensor_atomic_use_result[cas-1-float16]
9191
python/test/unit/language/test_core.py::test_tensor_atomic_use_result[cas-4-float16]
9292
python/test/unit/language/test_core.py::test_tensor_atomic_use_result[cas-16-float16]
93+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-None]
94+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-acquire]
95+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-release]
96+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-acq_rel]
97+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-relaxed]
98+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-None]
99+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-acquire]
100+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-release]
101+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-acq_rel]
102+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-relaxed]
103+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-None]
104+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-acquire]
105+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-release]
106+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-acq_rel]
107+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-relaxed]
108+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-None]
109+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-acquire]
110+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-release]
111+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-acq_rel]
112+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-relaxed]
113+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-None]
114+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-acquire]
115+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-release]
116+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-acq_rel]
117+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-relaxed]
118+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-None]
119+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-acquire]
120+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-release]
121+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-acq_rel]
122+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-relaxed]

‎scripts/skiplist/lts/language.txt‎

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -340,3 +340,33 @@ python/test/unit/language/test_tensor_descriptor.py::test_tensor_descriptor_redu
340340
python/test/unit/language/test_core.py::test_tensor_atomic_use_result[cas-1-float16]
341341
python/test/unit/language/test_core.py::test_tensor_atomic_use_result[cas-4-float16]
342342
python/test/unit/language/test_core.py::test_tensor_atomic_use_result[cas-16-float16]
343+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-None]
344+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-acquire]
345+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-release]
346+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-acq_rel]
347+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-relaxed]
348+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-None]
349+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-acquire]
350+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-release]
351+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-acq_rel]
352+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-relaxed]
353+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-None]
354+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-acquire]
355+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-release]
356+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-acq_rel]
357+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-relaxed]
358+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-None]
359+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-acquire]
360+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-release]
361+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-acq_rel]
362+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-relaxed]
363+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-None]
364+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-acquire]
365+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-release]
366+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-acq_rel]
367+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-relaxed]
368+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-None]
369+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-acquire]
370+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-release]
371+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-acq_rel]
372+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-relaxed]

‎scripts/skiplist/mtl/language.txt‎

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -390,3 +390,33 @@ python/test/unit/language/test_tensor_descriptor.py::test_tensor_descriptor_redu
390390
python/test/unit/language/test_core.py::test_tensor_atomic_use_result[cas-1-float16]
391391
python/test/unit/language/test_core.py::test_tensor_atomic_use_result[cas-4-float16]
392392
python/test/unit/language/test_core.py::test_tensor_atomic_use_result[cas-16-float16]
393+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-None]
394+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-acquire]
395+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-release]
396+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-acq_rel]
397+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-4-1-relaxed]
398+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-None]
399+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-acquire]
400+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-release]
401+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-acq_rel]
402+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-128-1-relaxed]
403+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-None]
404+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-acquire]
405+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-release]
406+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-acq_rel]
407+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[bfloat16-512-1-relaxed]
408+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-None]
409+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-acquire]
410+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-release]
411+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-acq_rel]
412+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-4-1-relaxed]
413+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-None]
414+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-acquire]
415+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-release]
416+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-acq_rel]
417+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-128-1-relaxed]
418+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-None]
419+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-acquire]
420+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-release]
421+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-acq_rel]
422+
python/test/unit/language/test_core.py::test_tensor_atomic_cas[float16-512-1-relaxed]

0 commit comments

Comments
 (0)