76 lines · plain
1// RUN: %clang_cc1 -fsyntax-only -verify %s2// RUN: %clang_cc1 -fsyntax-only -verify -fcuda-is-device %s3// RUN: %clang_cc1 -fsyntax-only -verify -fcuda-is-device %s \4// RUN: -fatomic-fine-grained-memory -fatomic-ignore-denormal-mode5 6#include "Inputs/cuda.h"7 8#if !__has_extension(clang_atomic_attributes)9#error "We should have atomic attributes support"10#endif11 12[[clang::atomic(!no_remote_memory)]] // expected-error {{use of undeclared identifier 'no_remote_memory'}}13__device__ __host__ void test_location(float *a) {14 __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);15 [[clang::atomic(!no_remote_memory)]] int x; // expected-error {{use of undeclared identifier 'no_remote_memory'}}16}17 18__device__ __host__ void test_invalid_option(float *a) {19 [[clang::atomic(fast)]] { // expected-error {{invalid argument 'fast' to atomic attribute; valid options are: 'remote_memory', 'fine_grained_memory', 'ignore_denormal_mode' (optionally prefixed with 'no_')}}20 __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);21 }22}23 24__device__ __host__ void test_invalid_value(float *a) {25 [[clang::atomic(no_remote_memory(default))]] { // expected-error2 {{expected ','}} expected-error {{expected ')'}}26 __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);27 }28}29 30__device__ __host__ void test_invalid_format(float *a) {31 [[clang::atomic(no_remote_memory=on)]] { // expected-error2 {{expected ','}} expected-error {{expected ')'}}32 __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);33 }34}35 36[[clang::atomic(no_remote_memory)]] // expected-error {{'clang::atomic' attribute cannot be applied to a declaration}}37__device__ __host__ void test_not_compound_stmt(float *a) {38 __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);39}40 41__device__ __host__ void test_quoted(float *a) {42 [[clang::atomic("no_remote_memory", "remote_memory")]] { // expected-error {{'clang::atomic' attribute requires an identifier}}43 __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);44 }45}46 47__device__ __host__ void test_one_value(float *a) {48 [[clang::atomic(no_remote_memory)]] {49 __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);50 }51}52 53__device__ __host__ void test_multiple_value(float *a) {54 [[clang::atomic(no_remote_memory, fine_grained_memory)]] {55 __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);56 }57}58 59__device__ __host__ void test_duplicate_value(float *a) {60 [[clang::atomic(no_remote_memory, no_remote_memory)]] {61 __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);62 }63}64 65__device__ __host__ void test_conflict_value(float *a) {66 [[clang::atomic(no_remote_memory, remote_memory)]] {67 __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);68 }69}70 71__device__ __host__ void test_multiple_attrs(float *a) {72 [[clang::atomic(no_remote_memory)]] [[clang::atomic(remote_memory)]] {73 __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);74 }75}76