These should have been removed before with r128225.

// RUN: %clang_cc1 -emit-pch -o %t %s

// RUN: %clang_cc1 -include-pch %t -fsyntax-only %s 

#ifndef HEADER

#define HEADER

// Header.

#include "../SemaCUDA/cuda.h"

void kcall(void (*kp)()) {

  kp<<<1, 1>>>();

}

__global__ void kern() {

}

#else

// Using the header.

void test() {

  kcall(kern);

  kern<<<1, 1>>>();

}

#endif

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

void foo(void) {

  foo<<<1;      // expected-error {{expected '>>>'}} expected-note {{to match this '<<<'}}

  foo<<<1,1>>>; // expected-error {{expected '('}}

  foo<<<>>>();  // expected-error {{expected expression}}

}

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

void cudaConfigureCall(unsigned gridSize, unsigned blockSize); // expected-error {{must have scalar return type}}

/* Minimal declarations for CUDA support.  Testing purposes only. */

#include <stddef.h>

#define __constant__ __attribute__((constant))

#define __device__ __attribute__((device))

#define __global__ __attribute__((global))

#define __host__ __attribute__((host))

#define __shared__ __attribute__((shared))

struct dim3 {

  unsigned x, y, z;

  dim3(unsigned x, unsigned y = 1, unsigned z = 1) : x(x), y(y), z(z) {}

};

typedef struct cudaStream *cudaStream_t;

int cudaConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0,

                      cudaStream_t stream = 0);

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

#include "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);

  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}}

}