[PATCH] D44747: Set calling convention for CUDA kernel

[John McCall via Phabricator via cfe-commits]

rjmccall added inline comments.


================
Comment at: lib/AST/Type.cpp:2762
+  case CC_CUDAKernel:
+    return "cuda_kernel";
   }
----------------
For consistency with the rest of this switch, please put the return on the same line as its case.


================
Comment at: lib/AST/TypePrinter.cpp:781
     case CC_OpenCLKernel:
+    case CC_CUDAKernel:
       // Do nothing. These CCs are not available as attributes.
----------------
I think the spelling for this is `__global__`.  You might need to adjust printing because this isn't the right place to print it, of course.


================
Comment at: lib/CodeGen/CGCall.cpp:68
+  case CC_CUDAKernel:
+    return CGM.getTargetCodeGenInfo().getCUDAKernelCallingConv();
   }
----------------
For consistency with the rest of this switch, please put the return on the same line as its case.


================
Comment at: lib/Sema/SemaOverload.cpp:1492
+      Changed = true;
+    }
+
----------------
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?


================
Comment at: tools/libclang/CXType.cpp:630
+    case CC_CUDAKernel:
+      return CXCallingConv_Unexposed;
       break;
----------------
Formatting.


https://reviews.llvm.org/D44747