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



_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits

Reply via email to