18e2d82297
Summary: Some function calls in CUDA are allowed to appear in semantically-correct programs but are an error if they're ever codegen'ed. Specifically, a host+device function may call a host function, but it's an error if such a function is ever codegen'ed in device mode (and vice versa). Previously, clang made no attempt to catch these errors. For the most part, they would be caught by ptxas, and reported as "call to unknown function 'foo'". Now we catch these errors and report them the same as we report other illegal calls (e.g. a call from a host function to a device function). This has a small change in error-message behavior for calls that were previously disallowed (e.g. calls from a host to a device function). Previously, we'd catch disallowed calls fairly early, before doing additional semantic checking e.g. of the call's arguments. Now we catch these illegal calls at the very end of our semantic checks, so we'll only emit a "illegal CUDA call" error if the call is otherwise well-formed. Reviewers: tra, rnk Subscribers: cfe-commits Differential Revision: https://reviews.llvm.org/D23242 llvm-svn: 278759
30 lines
958 B
C
30 lines
958 B
C
/* Minimal declarations for CUDA support. Testing purposes only. */
|
|
|
|
#include <stddef.h>
|
|
|
|
// Make this file work with nvcc, for testing compatibility.
|
|
|
|
#ifndef __NVCC__
|
|
#define __constant__ __attribute__((constant))
|
|
#define __device__ __attribute__((device))
|
|
#define __global__ __attribute__((global))
|
|
#define __host__ __attribute__((host))
|
|
#define __shared__ __attribute__((shared))
|
|
#define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__)))
|
|
|
|
struct dim3 {
|
|
unsigned x, y, z;
|
|
__host__ __device__ 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);
|
|
|
|
// Device-side placement new overloads.
|
|
__device__ void *operator new(__SIZE_TYPE__, void *p) { return p; }
|
|
__device__ void *operator new[](__SIZE_TYPE__, void *p) { return p; }
|
|
|
|
#endif // !__NVCC__
|