// RUN: %clang_cc1 -fsyntax-only -verify %s

#include "Inputs/cuda.h"

__global__ void g1(int x) {}

template <typename T> void t1(T arg) {
  g1<<<arg, arg>>>(1);
}

void h1(int x) {}
int h2(int x) { return 1; }

int main(void) {
  g1<<<1, 1>>>(42);
  g1(42); // expected-error {{call to global function 'g1' not configured}}
  g1<<<1>>>(42); // expected-error {{too few execution configuration arguments to kernel function call}}
  g1<<<1, 1, 0, 0, 0>>>(42); // expected-error {{too many execution configuration arguments to kernel function call}}

  t1(1);

  h1<<<1, 1>>>(42); // expected-error {{kernel call to non-global function 'h1'}}

  int (*fp)(int) = h2;
  fp<<<1, 1>>>(42); // expected-error {{must have void return type}}

  g1<<<undeclared, 1>>>(42); // expected-error {{use of undeclared identifier 'undeclared'}}
}

// Make sure we can call static member kernels.
template <typename > struct a0 {
  template <typename T> static __global__ void Call(T);
};
struct a1 {
  template <typename T> static __global__ void Call(T);
};
template <typename T> struct a2 {
  static __global__ void Call(T);
};
struct a3 {
  static __global__ void Call(int);
  static __global__ void Call(void*);
};

struct b {
  template <typename c> void d0(c arg) {
    a0<c>::Call<<<0, 0>>>(arg);
    a1::Call<<<0,0>>>(arg);
    a2<c>::Call<<<0,0>>>(arg);
    a3::Call<<<0, 0>>>(arg);
  }
  void d1(void* arg) {
    a0<void*>::Call<<<0, 0>>>(arg);
    a1::Call<<<0,0>>>(arg);
    a2<void*>::Call<<<0,0>>>(arg);
    a3::Call<<<0, 0>>>(arg);
  }
  void e() { d0(1); }
};