Skip to content

Commit

Permalink
Use order and scope macros in test
Browse files Browse the repository at this point in the history
  • Loading branch information
arsenm committed Jun 26, 2024
1 parent 0d9ab2b commit 5c2be85
Showing 1 changed file with 22 additions and 22 deletions.
44 changes: 22 additions & 22 deletions clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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
Expand Down

0 comments on commit 5c2be85

Please sign in to comment.