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



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

Reply via email to