// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -fcuda-is-device \
// RUN: -fsyntax-only -verify %s
#include "Inputs/cuda.h"
// expected-no-diagnostics
__device__ void __threadfence_system() {
// This shouldn't produce an error, since __nvvm_membar_sys should be
// __device__ and thus callable from device code.
__nvvm_membar_sys();
}