[cfe-commits] r141300 - in /cfe/trunk: lib/CodeGen/CGCUDANV.cpp lib/CodeGen/CGCUDARuntime.cpp lib/CodeGen/CGCUDARuntime.h lib/CodeGen/CGExpr.cpp lib/CodeGen/CGExprCXX.cpp lib/CodeGen/CMakeLists.txt lib/CodeGen/CodeGenFunction.h lib/CodeGen/CodeGenModule.cpp lib/CodeGen/CodeGenModule.h test/CodeGenCUDA/kernel-call.cu
Peter Collingbourne
peter at pcc.me.uk
Thu Oct 6 11:29:37 PDT 2011
Author: pcc
Date: Thu Oct 6 13:29:37 2011
New Revision: 141300
URL: http://llvm.org/viewvc/llvm-project?rev=141300&view=rev
Log:
CUDA: IR generation support for kernel call expressions
Added:
cfe/trunk/lib/CodeGen/CGCUDANV.cpp
cfe/trunk/lib/CodeGen/CGCUDARuntime.cpp
cfe/trunk/lib/CodeGen/CGCUDARuntime.h
cfe/trunk/test/CodeGenCUDA/kernel-call.cu
Modified:
cfe/trunk/lib/CodeGen/CGExpr.cpp
cfe/trunk/lib/CodeGen/CGExprCXX.cpp
cfe/trunk/lib/CodeGen/CMakeLists.txt
cfe/trunk/lib/CodeGen/CodeGenFunction.h
cfe/trunk/lib/CodeGen/CodeGenModule.cpp
cfe/trunk/lib/CodeGen/CodeGenModule.h
Added: cfe/trunk/lib/CodeGen/CGCUDANV.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGCUDANV.cpp?rev=141300&view=auto
==============================================================================
--- cfe/trunk/lib/CodeGen/CGCUDANV.cpp (added)
+++ cfe/trunk/lib/CodeGen/CGCUDANV.cpp Thu Oct 6 13:29:37 2011
@@ -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);
+}
Added: cfe/trunk/lib/CodeGen/CGCUDARuntime.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGCUDARuntime.cpp?rev=141300&view=auto
==============================================================================
--- cfe/trunk/lib/CodeGen/CGCUDARuntime.cpp (added)
+++ cfe/trunk/lib/CodeGen/CGCUDARuntime.cpp Thu Oct 6 13:29:37 2011
@@ -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);
+}
Added: cfe/trunk/lib/CodeGen/CGCUDARuntime.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGCUDARuntime.h?rev=141300&view=auto
==============================================================================
--- cfe/trunk/lib/CodeGen/CGCUDARuntime.h (added)
+++ cfe/trunk/lib/CodeGen/CGCUDARuntime.h Thu Oct 6 13:29:37 2011
@@ -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
Modified: cfe/trunk/lib/CodeGen/CGExpr.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGExpr.cpp?rev=141300&r1=141299&r2=141300&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGExpr.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGExpr.cpp Thu Oct 6 13:29:37 2011
@@ -2192,6 +2192,9 @@
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())
Modified: cfe/trunk/lib/CodeGen/CGExprCXX.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGExprCXX.cpp?rev=141300&r1=141299&r2=141300&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGExprCXX.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGExprCXX.cpp Thu Oct 6 13:29:37 2011
@@ -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 @@
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) {
Modified: cfe/trunk/lib/CodeGen/CMakeLists.txt
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CMakeLists.txt?rev=141300&r1=141299&r2=141300&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CMakeLists.txt (original)
+++ cfe/trunk/lib/CodeGen/CMakeLists.txt Thu Oct 6 13:29:37 2011
@@ -14,6 +14,8 @@
CGBuiltin.cpp
CGCall.cpp
CGClass.cpp
+ CGCUDANV.cpp
+ CGCUDARuntime.cpp
CGCXX.cpp
CGCXXABI.cpp
CGCleanup.cpp
Modified: cfe/trunk/lib/CodeGen/CodeGenFunction.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenFunction.h?rev=141300&r1=141299&r2=141300&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CodeGenFunction.h (original)
+++ cfe/trunk/lib/CodeGen/CodeGenFunction.h Thu Oct 6 13:29:37 2011
@@ -2076,6 +2076,9 @@
const CXXMethodDecl *MD,
ReturnValueSlot ReturnValue);
+ RValue EmitCUDAKernelCallExpr(const CUDAKernelCallExpr *E,
+ ReturnValueSlot ReturnValue);
+
RValue EmitBuiltinExpr(const FunctionDecl *FD,
unsigned BuiltinID, const CallExpr *E);
Modified: cfe/trunk/lib/CodeGen/CodeGenModule.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenModule.cpp?rev=141300&r1=141299&r2=141300&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CodeGenModule.cpp (original)
+++ cfe/trunk/lib/CodeGen/CodeGenModule.cpp Thu Oct 6 13:29:37 2011
@@ -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 @@
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 @@
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() {
delete ObjCRuntime;
delete OpenCLRuntime;
+ delete CUDARuntime;
delete &ABI;
delete TBAA;
delete DebugInfo;
@@ -131,6 +135,10 @@
OpenCLRuntime = new CGOpenCLRuntime(*this);
}
+void CodeGenModule::createCUDARuntime() {
+ CUDARuntime = CreateNVCUDARuntime(*this);
+}
+
void CodeGenModule::Release() {
EmitDeferred();
EmitCXXGlobalInitFunc();
Modified: cfe/trunk/lib/CodeGen/CodeGenModule.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenModule.h?rev=141300&r1=141299&r2=141300&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CodeGenModule.h (original)
+++ cfe/trunk/lib/CodeGen/CodeGenModule.h Thu Oct 6 13:29:37 2011
@@ -76,6 +76,7 @@
class CGDebugInfo;
class CGObjCRuntime;
class CGOpenCLRuntime;
+ class CGCUDARuntime;
class BlockFieldFlags;
class FunctionArgList;
@@ -228,6 +229,7 @@
CGObjCRuntime* ObjCRuntime;
CGOpenCLRuntime* OpenCLRuntime;
+ CGCUDARuntime* CUDARuntime;
CGDebugInfo* DebugInfo;
ARCEntrypoints *ARCData;
RREntrypoints *RRData;
@@ -320,6 +322,7 @@
void createObjCRuntime();
void createOpenCLRuntime();
+ void createCUDARuntime();
llvm::LLVMContext &VMContext;
@@ -361,12 +364,18 @@
/// 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; }
Added: cfe/trunk/test/CodeGenCUDA/kernel-call.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/kernel-call.cu?rev=141300&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/kernel-call.cu (added)
+++ cfe/trunk/test/CodeGenCUDA/kernel-call.cu Thu Oct 6 13:29:37 2011
@@ -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);
+}
More information about the cfe-commits
mailing list