From 5c2be8525ead0ddd5ac3c431d0a550018a31cdfd Mon Sep 17 00:00:00 2001 From: Matt Arsenault Date: Wed, 26 Jun 2024 18:11:06 +0200 Subject: [PATCH] Use order and scope macros in test --- .../test/CodeGenOpenCL/builtins-amdgcn-vi.cl | 44 +++++++++---------- 1 file changed, 22 insertions(+), 22 deletions(-) diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl index 4e3b1e14ead62b..5bd8f77a5930c4 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl @@ -183,19 +183,19 @@ void test_ds_fminf(__attribute__((address_space(3))) float *out, float src) { *out = __builtin_amdgcn_ds_fminf(out, src, 0, 0, true); // Test all orders. - *out = __builtin_amdgcn_ds_fminf(out, src, 1, 0, false); - *out = __builtin_amdgcn_ds_fminf(out, src, 2, 0, false); - *out = __builtin_amdgcn_ds_fminf(out, src, 3, 0, false); - *out = __builtin_amdgcn_ds_fminf(out, src, 4, 0, false); - *out = __builtin_amdgcn_ds_fminf(out, src, 5, 0, false); - *out = __builtin_amdgcn_ds_fminf(out, src, 5, 0, false); // invalid + *out = __builtin_amdgcn_ds_fminf(out, src, __ATOMIC_CONSUME, __MEMORY_SCOPE_SYSTEM, false); + *out = __builtin_amdgcn_ds_fminf(out, src, __ATOMIC_ACQUIRE, __MEMORY_SCOPE_SYSTEM, false); + *out = __builtin_amdgcn_ds_fminf(out, src, __ATOMIC_RELEASE, __MEMORY_SCOPE_SYSTEM, false); + *out = __builtin_amdgcn_ds_fminf(out, src, __ATOMIC_ACQ_REL, __MEMORY_SCOPE_SYSTEM, false); + *out = __builtin_amdgcn_ds_fminf(out, src, __ATOMIC_SEQ_CST, __MEMORY_SCOPE_SYSTEM, false); + *out = __builtin_amdgcn_ds_fminf(out, src, __ATOMIC_SEQ_CST, __MEMORY_SCOPE_SYSTEM, false); // invalid // Test all syncscopes. - *out = __builtin_amdgcn_ds_fminf(out, src, 0, 1, false); - *out = __builtin_amdgcn_ds_fminf(out, src, 0, 2, false); - *out = __builtin_amdgcn_ds_fminf(out, src, 0, 3, false); - *out = __builtin_amdgcn_ds_fminf(out, src, 0, 4, false); - *out = __builtin_amdgcn_ds_fminf(out, src, 0, 5, false); // invalid + *out = __builtin_amdgcn_ds_fminf(out, src, __ATOMIC_RELAXED, __MEMORY_SCOPE_DEVICE, false); + *out = __builtin_amdgcn_ds_fminf(out, src, __ATOMIC_RELAXED, __MEMORY_SCOPE_WRKGRP, false); + *out = __builtin_amdgcn_ds_fminf(out, src, __ATOMIC_RELAXED, __MEMORY_SCOPE_WVFRNT, false); + *out = __builtin_amdgcn_ds_fminf(out, src, __ATOMIC_RELAXED, __MEMORY_SCOPE_SINGLE, false); + *out = __builtin_amdgcn_ds_fminf(out, src, __ATOMIC_RELAXED, 5, false); // invalid } // CHECK-LABEL: @test_ds_fmax @@ -224,19 +224,19 @@ void test_ds_fmaxf(__attribute__((address_space(3))) float *out, float src) { *out = __builtin_amdgcn_ds_fmaxf(out, src, 0, 0, true); // Test all orders. - *out = __builtin_amdgcn_ds_fmaxf(out, src, 1, 0, false); - *out = __builtin_amdgcn_ds_fmaxf(out, src, 2, 0, false); - *out = __builtin_amdgcn_ds_fmaxf(out, src, 3, 0, false); - *out = __builtin_amdgcn_ds_fmaxf(out, src, 4, 0, false); - *out = __builtin_amdgcn_ds_fmaxf(out, src, 5, 0, false); - *out = __builtin_amdgcn_ds_fmaxf(out, src, 5, 0, false); // invalid + *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_CONSUME, __MEMORY_SCOPE_SYSTEM, false); + *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_ACQUIRE, __MEMORY_SCOPE_SYSTEM, false); + *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_RELEASE, __MEMORY_SCOPE_SYSTEM, false); + *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_ACQ_REL, __MEMORY_SCOPE_SYSTEM, false); + *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_SEQ_CST, __MEMORY_SCOPE_SYSTEM, false); + *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_SEQ_CST, __MEMORY_SCOPE_SYSTEM, false); // invalid // Test all syncscopes. - *out = __builtin_amdgcn_ds_fmaxf(out, src, 0, 1, false); - *out = __builtin_amdgcn_ds_fmaxf(out, src, 0, 2, false); - *out = __builtin_amdgcn_ds_fmaxf(out, src, 0, 3, false); - *out = __builtin_amdgcn_ds_fmaxf(out, src, 0, 4, false); - *out = __builtin_amdgcn_ds_fmaxf(out, src, 0, 5, false); // invalid + *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_RELAXED, __MEMORY_SCOPE_DEVICE, false); + *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_RELAXED, __MEMORY_SCOPE_WRKGRP, false); + *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_RELAXED, __MEMORY_SCOPE_WVFRNT, false); + *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_RELAXED, __MEMORY_SCOPE_SINGLE, false); + *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_RELAXED, 5, false); // invalid } // CHECK-LABEL: @test_s_memtime