[PATCH] D44747: Set calling convention for CUDA kernel

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl marked 3 inline comments as done.
yaxunl added inline comments.


================
Comment at: lib/Sema/SemaOverload.cpp:1492
+      Changed = true;
+    }
+
----------------
rjmccall wrote:
> It's cheaper not to check the CUDA language mode here; pulling the CC out of the FPT is easy.
> 
> Why is this necessary, anyway?  From the spec, it doesn't look to me like kernel function pointers can be converted to ordinary function pointers.  A kernel function pointer is supposed to be declared with something like `__global__ void (*fn)(void)`.  You'll need to change your patch to SemaType to apply the CC even when compiling for the host, of course.
> 
> I was going to say that you should use this CC in your validation that calls with execution configurations go to kernel functions, but... I can't actually find where you do that validation.
> 
> Do you need these function pointers to be a different size from the host function pointer?
In CUDA, `__global__` can only be used with function declaration or definition. Using it in function pointer declaration will result in a warning: 'global' attribute only applies to functions.

Also, there is this lit test in SemaCUDA:

```
__global__ void kernel() {}

typedef void (*fn_ptr_t)();

__host__ fn_ptr_t get_ptr_h() {
  return kernel;
}

```
It allows implicit conversion of `__global__ void()` to void(*)(), therefore I need the above change to drop the CUDA kernel calling convention in such implicit conversion.


https://reviews.llvm.org/D44747