CUDA: IR generation support for kernel call expressions

llvm-svn: 141300

clang/lib/CodeGen/CGCUDANV.cpp

@@ -0,0 +1,34 @@
+//===----- CGCUDANV.cpp - Interface to NVIDIA CUDA Runtime ----------------===//
+//
+//                     The LLVM Compiler Infrastructure
+//
+// This file is distributed under the University of Illinois Open Source
+// License. See LICENSE.TXT for details.
+//
+//===----------------------------------------------------------------------===//
+//
+// This provides a class for CUDA code generation targeting the NVIDIA CUDA
+// runtime library.
+//
+//===----------------------------------------------------------------------===//
+
+#include "CGCUDARuntime.h"
+
+using namespace clang;
+using namespace CodeGen;
+
+namespace {
+
+class CGNVCUDARuntime : public CGCUDARuntime {
+public:
+  CGNVCUDARuntime(CodeGenModule &CGM);
+};
+
+}
+
+CGNVCUDARuntime::CGNVCUDARuntime(CodeGenModule &CGM) : CGCUDARuntime(CGM) {
+}
+
+CGCUDARuntime *CodeGen::CreateNVCUDARuntime(CodeGenModule &CGM) {
+  return new CGNVCUDARuntime(CGM);
+}

clang/lib/CodeGen/CGCUDARuntime.cpp

@@ -0,0 +1,55 @@
+//===----- CGCUDARuntime.cpp - Interface to CUDA Runtimes -----------------===//
+//
+//                     The LLVM Compiler Infrastructure
+//
+// This file is distributed under the University of Illinois Open Source
+// License. See LICENSE.TXT for details.
+//
+//===----------------------------------------------------------------------===//
+//
+// This provides an abstract class for CUDA code generation.  Concrete
+// subclasses of this implement code generation for specific CUDA
+// runtime libraries.
+//
+//===----------------------------------------------------------------------===//
+
+#include "CGCUDARuntime.h"
+#include "clang/AST/Decl.h"
+#include "clang/AST/ExprCXX.h"
+#include "CGCall.h"
+#include "CodeGenFunction.h"
+
+using namespace clang;
+using namespace CodeGen;
+
+CGCUDARuntime::~CGCUDARuntime() {}
+
+RValue CGCUDARuntime::EmitCUDAKernelCallExpr(CodeGenFunction &CGF,
+                                             const CUDAKernelCallExpr *E,
+                                             ReturnValueSlot ReturnValue) {
+  llvm::BasicBlock *ConfigOKBlock = CGF.createBasicBlock("kcall.configok");
+  llvm::BasicBlock *ContBlock = CGF.createBasicBlock("kcall.end");
+
+  CodeGenFunction::ConditionalEvaluation eval(CGF);
+  CGF.EmitBranchOnBoolExpr(E->getConfig(), ContBlock, ConfigOKBlock);
+
+  eval.begin(CGF);
+  CGF.EmitBlock(ConfigOKBlock);
+
+  const Decl *TargetDecl = 0;
+  if (const ImplicitCastExpr *CE = dyn_cast<ImplicitCastExpr>(E->getCallee())) {
+    if (const DeclRefExpr *DRE = dyn_cast<DeclRefExpr>(CE->getSubExpr())) {
+      TargetDecl = DRE->getDecl();
+    }
+  }
+
+  llvm::Value *Callee = CGF.EmitScalarExpr(E->getCallee());
+  CGF.EmitCall(E->getCallee()->getType(), Callee, ReturnValue,
+               E->arg_begin(), E->arg_end(), TargetDecl);
+  CGF.EmitBranch(ContBlock);
+
+  CGF.EmitBlock(ContBlock);
+  eval.end(CGF);
+
+  return RValue::get(0);
+}

clang/lib/CodeGen/CGCUDARuntime.h

@@ -0,0 +1,50 @@
+//===----- CGCUDARuntime.h - Interface to CUDA Runtimes ---------*- C++ -*-===//
+//
+//                     The LLVM Compiler Infrastructure
+//
+// This file is distributed under the University of Illinois Open Source
+// License. See LICENSE.TXT for details.
+//
+//===----------------------------------------------------------------------===//
+//
+// This provides an abstract class for CUDA code generation.  Concrete
+// subclasses of this implement code generation for specific CUDA
+// runtime libraries.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef CLANG_CODEGEN_CUDARUNTIME_H
+#define CLANG_CODEGEN_CUDARUNTIME_H
+
+namespace clang {
+
+class CUDAKernelCallExpr;
+
+namespace CodeGen {
+
+class CodeGenFunction;
+class CodeGenModule;
+class ReturnValueSlot;
+class RValue;
+
+class CGCUDARuntime {
+protected:
+  CodeGenModule &CGM;
+
+public:
+  CGCUDARuntime(CodeGenModule &CGM) : CGM(CGM) {}
+  virtual ~CGCUDARuntime();
+
+  virtual RValue EmitCUDAKernelCallExpr(CodeGenFunction &CGF,
+                                        const CUDAKernelCallExpr *E,
+                                        ReturnValueSlot ReturnValue);
+  
+};
+
+/// Creates an instance of a CUDA runtime class.
+CGCUDARuntime *CreateNVCUDARuntime(CodeGenModule &CGM);
+
+}
+}
+
+#endif

clang/lib/CodeGen/CGExpr.cpp

@@ -2192,6 +2192,9 @@ RValue CodeGenFunction::EmitCallExpr(const CallExpr *E,
   if (const CXXMemberCallExpr *CE = dyn_cast<CXXMemberCallExpr>(E))
     return EmitCXXMemberCallExpr(CE, ReturnValue);
 
+  if (const CUDAKernelCallExpr *CE = dyn_cast<CUDAKernelCallExpr>(E))
+    return EmitCUDAKernelCallExpr(CE, ReturnValue);
+
   const Decl *TargetDecl = E->getCalleeDecl();
   if (const FunctionDecl *FD = dyn_cast_or_null<FunctionDecl>(TargetDecl)) {
     if (unsigned builtinID = FD->getBuiltinID())

clang/lib/CodeGen/CGExprCXX.cpp

@@ -13,6 +13,7 @@
 
 #include "clang/Frontend/CodeGenOptions.h"
 #include "CodeGenFunction.h"
+#include "CGCUDARuntime.h"
 #include "CGCXXABI.h"
 #include "CGObjCRuntime.h"
 #include "CGDebugInfo.h"
@@ -347,6 +348,11 @@ CodeGenFunction::EmitCXXOperatorMemberCallExpr(const CXXOperatorCallExpr *E,
                            E->arg_begin() + 1, E->arg_end());
 }
 
+RValue CodeGenFunction::EmitCUDAKernelCallExpr(const CUDAKernelCallExpr *E,
+                                               ReturnValueSlot ReturnValue) {
+  return CGM.getCUDARuntime().EmitCUDAKernelCallExpr(*this, E, ReturnValue);
+}
+
 void
 CodeGenFunction::EmitCXXConstructExpr(const CXXConstructExpr *E,
                                       AggValueSlot Dest) {

clang/lib/CodeGen/CMakeLists.txt

@@ -14,6 +14,8 @@ add_clang_library(clangCodeGen
   CGBuiltin.cpp
   CGCall.cpp
   CGClass.cpp
+  CGCUDANV.cpp
+  CGCUDARuntime.cpp
   CGCXX.cpp
   CGCXXABI.cpp
   CGCleanup.cpp

clang/lib/CodeGen/CodeGenFunction.h

@@ -2076,6 +2076,9 @@ public:
                                        const CXXMethodDecl *MD,
                                        ReturnValueSlot ReturnValue);
 
+  RValue EmitCUDAKernelCallExpr(const CUDAKernelCallExpr *E,
+                                ReturnValueSlot ReturnValue);
+
 
   RValue EmitBuiltinExpr(const FunctionDecl *FD,
                          unsigned BuiltinID, const CallExpr *E);

clang/lib/CodeGen/CodeGenModule.cpp

@@ -16,6 +16,7 @@
 #include "CodeGenFunction.h"
 #include "CodeGenTBAA.h"
 #include "CGCall.h"
+#include "CGCUDARuntime.h"
 #include "CGCXXABI.h"
 #include "CGObjCRuntime.h"
 #include "CGOpenCLRuntime.h"
@@ -66,9 +67,9 @@ CodeGenModule::CodeGenModule(ASTContext &C, const CodeGenOptions &CGO,
     ABI(createCXXABI(*this)), 
     Types(C, M, TD, getTargetCodeGenInfo().getABIInfo(), ABI, CGO),
     TBAA(0),
-    VTables(*this), ObjCRuntime(0), OpenCLRuntime(0), DebugInfo(0), ARCData(0),
-    RRData(0), CFConstantStringClassRef(0), ConstantStringClassRef(0),
-    NSConstantStringType(0),
+    VTables(*this), ObjCRuntime(0), OpenCLRuntime(0), CUDARuntime(0),
+    DebugInfo(0), ARCData(0), RRData(0), CFConstantStringClassRef(0),
+    ConstantStringClassRef(0), NSConstantStringType(0),
     VMContext(M.getContext()),
     NSConcreteGlobalBlock(0), NSConcreteStackBlock(0),
     BlockObjectAssign(0), BlockObjectDispose(0),
@@ -77,6 +78,8 @@ CodeGenModule::CodeGenModule(ASTContext &C, const CodeGenOptions &CGO,
     createObjCRuntime();
   if (Features.OpenCL)
     createOpenCLRuntime();
+  if (Features.CUDA)
+    createCUDARuntime();
 
   // Enable TBAA unless it's suppressed.
   if (!CodeGenOpts.RelaxedAliasing && CodeGenOpts.OptimizationLevel > 0)
@@ -113,6 +116,7 @@ CodeGenModule::CodeGenModule(ASTContext &C, const CodeGenOptions &CGO,
 CodeGenModule::~CodeGenModule() {
   delete ObjCRuntime;
   delete OpenCLRuntime;
+  delete CUDARuntime;
   delete &ABI;
   delete TBAA;
   delete DebugInfo;
@@ -131,6 +135,10 @@ void CodeGenModule::createOpenCLRuntime() {
   OpenCLRuntime = new CGOpenCLRuntime(*this);
 }
 
+void CodeGenModule::createCUDARuntime() {
+  CUDARuntime = CreateNVCUDARuntime(*this);
+}
+
 void CodeGenModule::Release() {
   EmitDeferred();
   EmitCXXGlobalInitFunc();

clang/lib/CodeGen/CodeGenModule.h

@@ -76,6 +76,7 @@ namespace CodeGen {
   class CGDebugInfo;
   class CGObjCRuntime;
   class CGOpenCLRuntime;
+  class CGCUDARuntime;
   class BlockFieldFlags;
   class FunctionArgList;
   
@@ -228,6 +229,7 @@ class CodeGenModule : public CodeGenTypeCache {
 
   CGObjCRuntime* ObjCRuntime;
   CGOpenCLRuntime* OpenCLRuntime;
+  CGCUDARuntime* CUDARuntime;
   CGDebugInfo* DebugInfo;
   ARCEntrypoints *ARCData;
   RREntrypoints *RRData;
@@ -320,6 +322,7 @@ class CodeGenModule : public CodeGenTypeCache {
   void createObjCRuntime();
 
   void createOpenCLRuntime();
+  void createCUDARuntime();
 
   llvm::LLVMContext &VMContext;
 
@@ -361,12 +364,18 @@ public:
   /// been configured.
   bool hasObjCRuntime() { return !!ObjCRuntime; }
 
-  /// getObjCRuntime() - Return a reference to the configured OpenCL runtime.
+  /// getOpenCLRuntime() - Return a reference to the configured OpenCL runtime.
   CGOpenCLRuntime &getOpenCLRuntime() {
     assert(OpenCLRuntime != 0);
     return *OpenCLRuntime;
   }
 
+  /// getCUDARuntime() - Return a reference to the configured CUDA runtime.
+  CGCUDARuntime &getCUDARuntime() {
+    assert(CUDARuntime != 0);
+    return *CUDARuntime;
+  }
+
   /// getCXXABI() - Return a reference to the configured C++ ABI.
   CGCXXABI &getCXXABI() { return ABI; }
 

clang/test/CodeGenCUDA/kernel-call.cu

@@ -0,0 +1,13 @@
+// RUN: %clang_cc1 -emit-llvm %s -o - | FileCheck %s
+
+#include "../SemaCUDA/cuda.h"
+
+__global__ void g1(int x) {}
+
+int main(void) {
+  // CHECK: call{{.*}}cudaConfigureCall
+  // CHECK: icmp
+  // CHECK: br
+  // CHECK: call{{.*}}g1
+  g1<<<1, 1>>>(42);
+}