137 lines · plain
1// RUN: %clang_cc1 -ast-dump %s | FileCheck %s2// RUN: %clang_cc1 -ast-dump -fcuda-is-device %s | FileCheck %s3// RUN: %clang_cc1 -ast-dump -fcuda-is-device %s \4// RUN: -fatomic-fine-grained-memory -fatomic-ignore-denormal-mode \5// RUN: | FileCheck %s6 7#include "Inputs/cuda.h"8 9// CHECK-LABEL: FunctionDecl {{.*}} test_default10// CHECK-NOT: AttributedStmt11// CHECK-NOT: AtomicAttr12// CHECK: CompoundStmt13// CHECK-NEXT: `-AtomicExpr14__device__ __host__ void test_default(float *a) {15 __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);16}17 18// CHECK-LABEL: FunctionDecl {{.*}} test_one19// CHECK: `-AttributedStmt20// CHECK-NEXT: |-AtomicAttr {{.*}} no_remote_memory{{$}}21// CHECK-NEXT: `-CompoundStmt22// CHECK-NEXT: `-AtomicExpr23__device__ __host__ void test_one(float *a) {24 [[clang::atomic(no_remote_memory)]] {25 __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);26 }27}28 29// CHECK-LABEL: FunctionDecl {{.*}} test_two30// CHECK: `-AttributedStmt31// CHECK-NEXT: |-AtomicAttr {{.*}} remote_memory ignore_denormal_mode{{$}}32// CHECK-NEXT: `-CompoundStmt33// CHECK-NEXT: `-AtomicExpr34__device__ __host__ void test_two(float *a) {35 [[clang::atomic(remote_memory, ignore_denormal_mode)]] {36 __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);37 }38}39 40// CHECK-LABEL: FunctionDecl {{.*}} test_three41// CHECK: `-AttributedStmt42// CHECK-NEXT: |-AtomicAttr {{.*}} no_remote_memory fine_grained_memory no_ignore_denormal_mode{{$}}43// CHECK-NEXT: `-CompoundStmt44// CHECK-NEXT: `-AtomicExpr45__device__ __host__ void test_three(float *a) {46 [[clang::atomic(no_remote_memory, fine_grained_memory, no_ignore_denormal_mode)]] {47 __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);48 }49}50 51// CHECK-LABEL: FunctionDecl {{.*}} test_duplicate52// CHECK: `-AttributedStmt53// CHECK-NEXT: |-AtomicAttr {{.*}} no_remote_memory{{$}}54// CHECK-NEXT: `-CompoundStmt55// CHECK-NEXT: `-AtomicExpr56__device__ __host__ void test_duplicate(float *a) {57 [[clang::atomic(no_remote_memory, no_remote_memory)]] {58 __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);59 }60}61 62// CHECK-LABEL: FunctionDecl {{.*}} test_conflict63// CHECK: `-AttributedStmt64// CHECK-NEXT: |-AtomicAttr {{.*}} remote_memory{{$}}65// CHECK-NEXT: `-CompoundStmt66// CHECK-NEXT: `-AtomicExpr67__device__ __host__ void test_conflict(float *a) {68 [[clang::atomic(no_remote_memory, remote_memory)]] {69 __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);70 }71}72 73// CHECK-LABEL: FunctionDecl {{.*}} test_multiple_attrs74// CHECK: `-AttributedStmt75// CHECK-NEXT: |-AtomicAttr {{.*}} no_remote_memory{{$}}76// CHECK-NEXT: |-AtomicAttr {{.*}} remote_memory{{$}}77// CHECK-NEXT: `-CompoundStmt78// CHECK-NEXT: `-AtomicExpr79__device__ __host__ void test_multiple_attrs(float *a) {80 [[clang::atomic(no_remote_memory)]] [[clang::atomic(remote_memory)]] {81 __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);82 }83}84 85// CHECK-LABEL: FunctionDecl {{.*}} test_nested86// CHECK: CompoundStmt87// CHECK: |-AtomicExpr88// CHECK: `-AttributedStmt89// CHECK-NEXT: |-AtomicAttr {{.*}} remote_memory fine_grained_memory no_ignore_denormal_mode{{$}}90// CHECK-NEXT: `-CompoundStmt91// CHECK: |-AtomicExpr92// CHECK: |-AttributedStmt93// CHECK-NEXT: |-AtomicAttr {{.*}} no_remote_memory{{$}}94// CHECK-NEXT: `-CompoundStmt95// CHECK-NEXT: `-AtomicExpr96// CHECK: `-AttributedStmt97// CHECK-NEXT: |-AtomicAttr {{.*}} no_fine_grained_memory{{$}}98// CHECK-NEXT: `-CompoundStmt99// CHECK-NEXT: `-AtomicExpr100__device__ __host__ void test_nested(float *a) {101 __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);102 [[clang::atomic(remote_memory, fine_grained_memory, no_ignore_denormal_mode)]] {103 __scoped_atomic_fetch_max(a, 2, __ATOMIC_SEQ_CST, __MEMORY_SCOPE_DEVICE);104 [[clang::atomic(no_remote_memory)]] {105 __scoped_atomic_fetch_min(a, 3, __ATOMIC_ACQUIRE, __MEMORY_SCOPE_WRKGRP);106 }107 [[clang::atomic(no_fine_grained_memory)]] {108 __scoped_atomic_fetch_sub(a, 4, __ATOMIC_RELEASE, __MEMORY_SCOPE_WVFRNT);109 }110 }111}112 113// CHECK-LABEL: FunctionTemplateDecl {{.*}} test_template114// CHECK: |-FunctionDecl {{.*}} test_template 'void (T *)'115// CHECK: | |-CompoundStmt116// CHECK: | | `-AttributedStmt117// CHECK: | | |-AtomicAttr {{.*}} no_remote_memory fine_grained_memory no_ignore_denormal_mode{{$}}118// CHECK: | | `-CompoundStmt119// CHECK: | | `-CallExpr {{.*}} '<dependent type>'120// CHECK: `-FunctionDecl {{.*}} used test_template 'void (float *)' implicit_instantiation121// CHECK: |-CompoundStmt122// CHECK: | `-AttributedStmt123// CHECK: | |-AtomicAttr {{.*}} no_remote_memory fine_grained_memory no_ignore_denormal_mode{{$}}124// CHECK: | `-CompoundStmt125// CHECK: | `-AtomicExpr {{.*}} 'float'126template<typename T>127__device__ __host__ void test_template(T *a) {128 [[clang::atomic(no_remote_memory, fine_grained_memory, no_ignore_denormal_mode)]] {129 __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);130 }131}132 133__device__ __host__ void test_template_caller() {134 float *p;135 test_template(p);136}137