void DataChecks::CreateCrossCheck(IRBuilder<> &Builder, Value *V) {
IntegerType *PtrIntTy = Builder.getIntPtrTy(*DL);
// Cast the value to IntPtrType
auto *VTy = V->getType();
if (VTy->isPointerTy())
V = Builder.CreatePtrToInt(V, PtrIntTy);
else
V = Builder.CreateZExtOrTrunc(V, PtrIntTy);
// Call the check function.
// TODO: Optimize by inlining the call?
if (XCheckLog) {
DebugLoc Loc = Builder.GetInsertPoint()->getDebugLoc();
Value *caller = Builder.CreateGlobalStringPtr(Builder.GetInsertBlock()->getParent()->getName());
Value *line, *col, *file;
if (Loc) {
line = Builder.getInt32(Loc.getLine());
col = Builder.getInt32(Loc.getCol());
file = Builder.CreateGlobalStringPtr(Loc->getFilename());
} else {
line = Builder.getInt32(-1);
col = Builder.getInt32(-1);
file = Builder.CreateGlobalStringPtr("unknown");
}
Builder.CreateCall(CheckFnTy, CheckFn, { caller, file, line, col, V });
} else {
Builder.CreateCall(CheckFnTy, CheckFn, {V});
}
NumCrossChecks++;
}
Value*
code_emitter::emit_closure (filter_t *closure_filter, primary_t *args)
{
int num_args = compiler_num_filter_args(closure_filter) - 3;
Value *closure = NULL, *uservals;
userval_info_t *info;
int i;
g_assert(closure_filter->kind == FILTER_MATHMAP || closure_filter->kind == FILTER_NATIVE);
if (closure_filter->kind == FILTER_MATHMAP)
{
vector<Value*> args;
args.push_back(invocation_arg);
args.push_back(pools_arg);
args.push_back(make_int_const(num_args));
args.push_back(lookup_filter_function(module, closure_filter));
args.push_back(lookup_init_frame_function(module, closure_filter));
args.push_back(lookup_main_filter_function(module, closure_filter));
args.push_back(lookup_init_x_function(module, closure_filter));
args.push_back(lookup_init_y_function(module, closure_filter));
closure = builder->CreateCall(module->getFunction(string("alloc_closure_image")), args.begin(), args.end());
uservals = builder->CreateCall(module->getFunction(string("get_closure_uservals")), closure);
}
else
uservals = builder->CreateCall2(module->getFunction(string("alloc_uservals")),
pools_arg,
make_int_const(compiler_num_filter_args(closure_filter) - 3));
for (i = 0, info = closure_filter->userval_infos;
info != 0;
++i, info = info->next)
{
const char *set_func_name = get_userval_set_func_name(info->type);
Value *arg = emit_primary(&args[i]);
/* FIXME: remove this eventually - bool needs to be an int */
if (info->type == USERVAL_BOOL_CONST)
arg = promote(arg, TYPE_FLOAT);
builder->CreateCall3(module->getFunction(string(set_func_name)), uservals, make_int_const(i), arg);
}
g_assert(i == num_args);
if (closure_filter->kind == FILTER_MATHMAP)
{
builder->CreateCall3(module->getFunction(string("set_closure_pixel_size")),
closure, lookup_internal("__canvasPixelW"), lookup_internal("__canvasPixelH"));
return closure;
}
else
{
string filter_func_name = string("llvm_") + string(closure_filter->v.native.func_name);
return builder->CreateCall3(module->getFunction(filter_func_name),
invocation_arg, uservals, pools_arg);
}
}
void
code_emitter::setup_filter_function (bool is_main_filter_function)
{
Function *filter_function = is_main_filter_function
? lookup_main_filter_function(module, filter)
: lookup_filter_function(module, filter);
Function::arg_iterator args = filter_function->arg_begin();
Value *slice_arg = NULL;
if (is_main_filter_function)
{
slice_arg = args++;
slice_arg->setName("slice");
}
else
{
invocation_arg = args++;
invocation_arg->setName("invocation");
}
closure_arg = args++;
closure_arg->setName("closure");
if (is_main_filter_function)
{
x_vars_var = args++;
x_vars_var->setName("x_vars");
y_vars_var = args++;
y_vars_var->setName("y_vars");
}
set_internal(::lookup_internal(filter->v.mathmap.internals, "x", true), args++);
set_internal(::lookup_internal(filter->v.mathmap.internals, "y", true), args++);
set_internal(::lookup_internal(filter->v.mathmap.internals, "t", true), args++);
pools_arg = args++;
pools_arg->setName("pools");
BasicBlock *block = BasicBlock::Create("entry", filter_function);
builder = new IRBuilder<> (block);
if (is_main_filter_function)
{
x_vars_var = builder->CreateBitCast(x_vars_var, PointerType::getUnqual(x_vars_type));
y_vars_var = builder->CreateBitCast(y_vars_var, PointerType::getUnqual(y_vars_type));
frame_arg = builder->CreateCall(module->getFunction(string("get_slice_frame")), slice_arg);
invocation_arg = builder->CreateCall(module->getFunction(string("get_frame_invocation")), frame_arg);
}
set_internals_from_invocation(invocation_arg);
if (is_main_filter_function)
set_xy_vars_from_frame ();
else
setup_xy_vars_from_closure ();
alloc_complex_copy_var();
current_function = filter_function;
}
void
code_emitter::set_internals_from_invocation (Value *invocation_arg)
{
set_internal(::lookup_internal(filter->v.mathmap.internals, "__canvasPixelW", true),
builder->CreateCall(module->getFunction(string("get_invocation_img_width")), invocation_arg));
set_internal(::lookup_internal(filter->v.mathmap.internals, "__canvasPixelH", true),
builder->CreateCall(module->getFunction(string("get_invocation_img_height")), invocation_arg));
set_internal(::lookup_internal(filter->v.mathmap.internals, "__renderPixelW", true),
builder->CreateCall(module->getFunction(string("get_invocation_render_width")), invocation_arg));
set_internal(::lookup_internal(filter->v.mathmap.internals, "__renderPixelH", true),
builder->CreateCall(module->getFunction(string("get_invocation_render_height")), invocation_arg));
set_internal(::lookup_internal(filter->v.mathmap.internals, "R", true),
builder->CreateCall(module->getFunction(string("get_invocation_image_R")), invocation_arg));
}
Value* CallExprAST::Codegen()
{
//Look up the name in the global module
Function* CalleeF = theModule->getFunction(Callee);
if (CalleeF == 0)
{
#ifdef DEBUG
dumpVars();
#endif
cerr << "\033[31m ERROR: \033[37m Unknown Function Reference" << endl;
exit(EXIT_FAILURE);
}
if (CalleeF->arg_size() != Args.size())
{
#ifdef DEBUG
dumpVars();
#endif
cerr << "\033[31m ERROR: \033[37m Incorrect number of arguements" << endl;
exit(EXIT_FAILURE);
}
vector<Value*> ArgsV;
for (unsigned i = 0, e = Args.size(); i != e; ++i)
{
ArgsV.push_back(Args[i]->Codegen());
if (ArgsV.back() == 0)
return 0;
}
return Builder.CreateCall(CalleeF,ArgsV,"calltmp");
}
void CNodeCodeGenVisitor::Visit(CInput* s) {
IRBuilder<> builder = builders_.top();
Value* ptr_offset = builder.CreateGEP(ptr_, GetPtrOffset(s->GetOffset()));
Value* input = builder.CreateCall(get_char_);
builder.CreateStore(input, ptr_offset);
VisitNextCNode(s);
}
void CNodeCodeGenVisitor::Visit(COutput* s) {
IRBuilder<> builder = builders_.top();
Value* offset_ptr = builder.CreateGEP(ptr_, GetPtrOffset(s->GetOffset()));
Value* ptr_value = builder.CreateLoad(offset_ptr);
builder.CreateCall(put_char_, ptr_value);
VisitNextCNode(s);
}
Value*
code_emitter::promote (Value *val, int type)
{
switch (type)
{
case TYPE_FLOAT :
if (val->getType() == Type::Int32Ty)
val = builder->CreateCall(module->getFunction(string("promote_int_to_float")), val);
else
assert(val->getType() == Type::FloatTy);
break;
case TYPE_COMPLEX :
if (val->getType() == Type::Int32Ty)
val = builder->CreateCall(module->getFunction(string("promote_int_to_complex")), val);
else if (val->getType() == Type::FloatTy)
val = builder->CreateCall(module->getFunction(string("promote_float_to_complex")), val);
break;
}
return val;
}
/// compile_get - Emit code for ','
void BrainFTraceRecorder::compile_get(BrainFTraceNode *node,
IRBuilder<>& builder) {
Value *Ret = builder.CreateCall(getchar_func);
Value *Trunc =
builder.CreateTrunc(Ret, IntegerType::get(Ret->getContext(), 8));
builder.CreateStore(Ret, Trunc);
if (node->left != (BrainFTraceNode*)~0ULL)
compile_opcode(node->left, builder);
else {
HeaderPHI->addIncoming(DataPtr, builder.GetInsertBlock());
builder.CreateBr(Header);
}
}
/// compile_put - Emit code for '.'
void BrainFTraceRecorder::compile_put(BrainFTraceNode *node,
IRBuilder<>& builder) {
Value *Loaded = builder.CreateLoad(DataPtr);
Value *Print =
builder.CreateSExt(Loaded, IntegerType::get(Loaded->getContext(), 32));
builder.CreateCall(putchar_func, Print);
if (node->left != (BrainFTraceNode*)~0ULL)
compile_opcode(node->left, builder);
else {
HeaderPHI->addIncoming(DataPtr, builder.GetInsertBlock());
builder.CreateBr(Header);
}
}
// Calls to setjmp(p) are lowered to _setjmp3(p, 0) by the frontend.
// The idea behind _setjmp3 is that it takes an optional number of personality
// specific parameters to indicate how to restore the personality-specific frame
// state when longjmp is initiated. Typically, the current TryLevel is saved.
void WinEHStatePass::rewriteSetJmpCallSite(IRBuilder<> &Builder, Function &F,
CallSite CS, Value *State) {
// Don't rewrite calls with a weird number of arguments.
if (CS.getNumArgOperands() != 2)
return;
Instruction *Inst = CS.getInstruction();
SmallVector<OperandBundleDef, 1> OpBundles;
CS.getOperandBundlesAsDefs(OpBundles);
SmallVector<Value *, 3> OptionalArgs;
if (Personality == EHPersonality::MSVC_CXX) {
OptionalArgs.push_back(CxxLongjmpUnwind);
OptionalArgs.push_back(State);
OptionalArgs.push_back(emitEHLSDA(Builder, &F));
} else if (Personality == EHPersonality::MSVC_X86SEH) {
OptionalArgs.push_back(SehLongjmpUnwind);
OptionalArgs.push_back(State);
if (UseStackGuard)
OptionalArgs.push_back(Cookie);
} else {
llvm_unreachable("unhandled personality!");
}
SmallVector<Value *, 5> Args;
Args.push_back(
Builder.CreateBitCast(CS.getArgOperand(0), Builder.getInt8PtrTy()));
Args.push_back(Builder.getInt32(OptionalArgs.size()));
Args.append(OptionalArgs.begin(), OptionalArgs.end());
CallSite NewCS;
if (CS.isCall()) {
auto *CI = cast<CallInst>(Inst);
CallInst *NewCI = Builder.CreateCall(SetJmp3, Args, OpBundles);
NewCI->setTailCallKind(CI->getTailCallKind());
NewCS = NewCI;
} else {
auto *II = cast<InvokeInst>(Inst);
NewCS = Builder.CreateInvoke(
SetJmp3, II->getNormalDest(), II->getUnwindDest(), Args, OpBundles);
}
NewCS.setCallingConv(CS.getCallingConv());
NewCS.setAttributes(CS.getAttributes());
NewCS->setDebugLoc(CS->getDebugLoc());
Instruction *NewInst = NewCS.getInstruction();
NewInst->takeName(Inst);
Inst->replaceAllUsesWith(NewInst);
Inst->eraseFromParent();
}
Value*
code_emitter::setup_init_x_or_y_function (string function_name, const char *internal_name, StructType *vars_type)
{
current_function = module->getFunction(function_name);
Value *slice_arg;
Function::arg_iterator args = current_function->arg_begin();
slice_arg = args++;
slice_arg->setName("slice");
closure_arg = args++;
closure_arg->setName("closure");
set_internal(::lookup_internal(filter->v.mathmap.internals, internal_name, true), args++);
set_internal(::lookup_internal(filter->v.mathmap.internals, "t", true), args++);
BasicBlock *block = BasicBlock::Create("entry", current_function);
builder = new IRBuilder<> (block);
frame_arg = builder->CreateCall(module->getFunction(string("get_slice_frame")), slice_arg);
invocation_arg = builder->CreateCall(module->getFunction(string("get_frame_invocation")), frame_arg);
pools_arg = builder->CreateCall(module->getFunction("get_slice_pools"), slice_arg);
set_internals_from_invocation(invocation_arg);
set_xy_vars_from_frame();
ret_var = builder->CreateCall2(module->getFunction(string("_mathmap_pools_alloc")),
pools_arg, emit_sizeof(vars_type));
Value *vars_var = builder->CreateBitCast(ret_var, PointerType::getUnqual(vars_type));
alloc_complex_copy_var();
return vars_var;
}
/// EmitPutS - Emit a call to the puts function. This assumes that Str is
/// some pointer.
void llvm::EmitPutS(Value *Str, IRBuilder<> &B, const TargetData *TD) {
Module *M = B.GetInsertBlock()->getParent()->getParent();
AttributeWithIndex AWI[2];
AWI[0] = AttributeWithIndex::get(1, Attribute::NoCapture);
AWI[1] = AttributeWithIndex::get(~0u, Attribute::NoUnwind);
Value *PutS = M->getOrInsertFunction("puts", AttrListPtr::get(AWI, 2),
B.getInt32Ty(),
B.getInt8PtrTy(),
NULL);
CallInst *CI = B.CreateCall(PutS, CastToCStr(Str, B), "puts");
if (const Function *F = dyn_cast<Function>(PutS->stripPointerCasts()))
CI->setCallingConv(F->getCallingConv());
}
Value* MethodCall::codeGen(CodeGenContext &context) {
Function *function = context.module->getFunction(methodId.name.c_str());
assert( function != NULL );
std::vector<Value*> args;
IRBuilder<> *builder = context.currentBuilder();
ArgsList::const_iterator it;
for (it = argsList.begin(); it != argsList.end(); it++) {
args.push_back((**it).codeGen(context));
}
return builder->CreateCall(function, ArrayRef<Value*>(args)
, methodId.name.c_str());
}
Instruction *DIBuilder::insertDeclare(Value *Storage, DILocalVariable *VarInfo,
DIExpression *Expr, const DILocation *DL,
BasicBlock *InsertBB, Instruction *InsertBefore) {
assert(VarInfo && "empty or invalid DILocalVariable* passed to dbg.declare");
assert(DL && "Expected debug loc");
assert(DL->getScope()->getSubprogram() ==
VarInfo->getScope()->getSubprogram() &&
"Expected matching subprograms");
if (!DeclareFn)
DeclareFn = getDeclareIntrin(M);
trackIfUnresolved(VarInfo);
trackIfUnresolved(Expr);
Value *Args[] = {getDbgIntrinsicValueImpl(VMContext, Storage),
MetadataAsValue::get(VMContext, VarInfo),
MetadataAsValue::get(VMContext, Expr)};
IRBuilder<> B = getIRBForDbgInsertion(DL, InsertBB, InsertBefore);
return B.CreateCall(DeclareFn, Args);
}
static void IncrementTimeCounter(Value* Inc, Value* calle, unsigned Index, GlobalVariable* Counters, IRBuilder<>& Builder, Value* Point)
{
LLVMContext &Context = Inc->getContext();
//In order to insert instruction after Point, we use the nextIns function.
Value* nextIns = getNextIns(Point);
Builder.SetInsertPoint(dyn_cast<Instruction>(nextIns));
// Create the getelementptr constant expression
std::vector<Constant*> Indices(2);
Indices[0] = Constant::getNullValue(Type::getInt32Ty(Context));
Indices[1] = ConstantInt::get(Type::getInt32Ty(Context), Index);
Constant *ElementPtr =
ConstantExpr::getGetElementPtr(Counters, Indices);
// Load, increment and store the value back.
// Use this formula: a = a + end_time - start_time
ArrayRef<Value*> args;
CallInst* Inc_end = Builder.CreateCall(calle, args, "");
Value* OldVal = Builder.CreateLoad(ElementPtr, "OldTimeCounter");
Value* TmpVal = Builder.CreateFSub(OldVal, Inc, "TmpTimeCounter");
Value* NewVal = Builder.CreateFAdd(TmpVal, Inc_end, "NewTimeCounter");
Builder.CreateStore(NewVal, ElementPtr);
}
Value*
code_emitter::emit_rhs (rhs_t *rhs)
{
switch (rhs->kind) {
case RHS_PRIMARY :
return emit_primary(&rhs->v.primary);
case RHS_INTERNAL :
return lookup_internal(rhs->v.internal);
case RHS_OP :
{
operation_t *op = rhs->v.op.op;
type_t promotion_type = TYPE_NIL;
char *function_name = compiler_function_name_for_op_rhs(rhs, &promotion_type);
if (promotion_type == TYPE_NIL)
assert(op->type_prop == TYPE_PROP_CONST);
if (op->type_prop != TYPE_PROP_CONST)
assert(promotion_type != TYPE_NIL);
Function *func = module->getFunction(string(function_name));
g_assert(func);
vector<Value*> args;
args.push_back(invocation_arg);
args.push_back(closure_arg);
args.push_back(pools_arg);
for (int i = 0; i < rhs->v.op.op->num_args; ++i) {
type_t type = promotion_type == TYPE_NIL ? op->arg_types[i] : promotion_type;
Value *val = emit_primary(&rhs->v.op.args[i], type == TYPE_FLOAT);
val = promote(val, type);
#ifndef __MINGW32__
if (sizeof(gpointer) == 4 && val->getType() == llvm_type_for_type(module, TYPE_COMPLEX))
{
Value *copy = builder->CreateAlloca(llvm_type_for_type(module, TYPE_COMPLEX));
builder->CreateStore(val, copy);
val = copy;
}
#endif
#ifdef DEBUG_OUTPUT
val->dump();
#endif
args.push_back(val);
}
#ifdef DEBUG_OUTPUT
func->dump();
#endif
Value *result = builder->CreateCall(func, args.begin(), args.end());
/* FIXME: this is ugly - we should check for the type
of the operation or resulting value */
if (is_complex_return_type(result->getType()))
result = convert_complex_return_value(result);
return result;
}
case RHS_FILTER :
{
int num_args = compiler_num_filter_args(rhs->v.filter.filter);
Value *closure = emit_closure(rhs->v.filter.filter, rhs->v.filter.args);
Function *func = lookup_filter_function(module, rhs->v.filter.filter);
vector<Value*> args;
args.push_back(invocation_arg);
args.push_back(closure);
args.push_back(emit_primary(&rhs->v.filter.args[num_args - 3]));
args.push_back(emit_primary(&rhs->v.filter.args[num_args - 2]));
args.push_back(emit_primary(&rhs->v.filter.args[num_args - 1]));
args.push_back(pools_arg);
return builder->CreateCall(func, args.begin(), args.end());
}
case RHS_CLOSURE :
return emit_closure(rhs->v.closure.filter, rhs->v.closure.args);
case RHS_TUPLE :
case RHS_TREE_VECTOR :
{
Function *set_func = module->getFunction(string("tuple_set"));
Value *tuple = builder->CreateCall2(module->getFunction(string("alloc_tuple")),
pools_arg,
make_int_const(rhs->v.tuple.length));
int i;
for (i = 0; i < rhs->v.tuple.length; ++i)
{
Value *val = emit_primary(&rhs->v.tuple.args[i], true);
builder->CreateCall3(set_func, tuple, make_int_const(i), val);
}
if (rhs->kind == RHS_TREE_VECTOR)
{
return builder->CreateCall3(module->getFunction(string("alloc_tree_vector")),
pools_arg,
make_int_const(rhs->v.tuple.length),
tuple);
}
else
return tuple;
}
default :
g_assert_not_reached();
}
}
void
code_emitter::set_xy_vars_from_frame ()
{
Value *xy_vars_untyped = builder->CreateCall(module->getFunction(string("get_frame_xy_vars")), frame_arg);
xy_vars_var = builder->CreateBitCast(xy_vars_untyped, PointerType::getUnqual(xy_vars_type));
}
Value*
code_emitter::emit_primary (primary_t *primary, bool need_float)
{
switch (primary->kind)
{
case PRIMARY_VALUE :
if (primary->v.value->index < 0)
{
switch (primary->v.value->compvar->type)
{
case TYPE_INT :
return make_int_const(0);
case TYPE_FLOAT :
return make_float_const(0.0);
case TYPE_IMAGE :
return builder->CreateCall(module->getFunction(string("get_uninited_image")));
default :
g_assert_not_reached();
}
}
else
{
Value *val = lookup_value(primary->v.value);
if (need_float)
val = promote(val, TYPE_FLOAT);
return val;
}
case PRIMARY_CONST :
switch (primary->const_type) {
case TYPE_INT :
if (need_float)
return make_float_const((float)primary->v.constant.int_value);
else
return make_int_const(primary->v.constant.int_value);
case TYPE_FLOAT :
return make_float_const(primary->v.constant.float_value);
case TYPE_COMPLEX :
{
assert(!need_float);
Value *val = builder->CreateCall2(module->getFunction(string("make_complex")),
make_float_const(__real__ primary->v.constant.complex_value),
make_float_const(__imag__ primary->v.constant.complex_value));
return convert_complex_return_value(val);
}
case TYPE_COLOR :
assert(!need_float);
return builder->CreateCall4(module->getFunction(string("make_color")),
make_int_const(RED(primary->v.constant.color_value)),
make_int_const(GREEN(primary->v.constant.color_value)),
make_int_const(BLUE(primary->v.constant.color_value)),
make_int_const(ALPHA(primary->v.constant.color_value)));
default :
g_assert_not_reached();
}
default:
g_assert_not_reached();
}
}
int main(int argc, char **argv)
{
InitializeNativeTarget();
LLVMContext &Context = getGlobalContext();
Module *m = new Module("test", Context);
Module *TheM = LoadModule("struct.c.bc");
std::string ErrMsg;
if (Linker::LinkModules(m, TheM, Linker::DestroySource, &ErrMsg)) {
std::cout << "error" << ErrMsg << std::endl;
return 1;
}
StructType* ctxTy = StructType::create(Context, "struct.kcontext_t");
StructType* sfpTy = StructType::create(Context, "struct.ksfp_t");
Type *Int64Ty = Type::getInt64Ty(Context);
Type *ArgsTy[] = {
PointerType::get(ctxTy, 0),
PointerType::get(sfpTy, 0),
Int64Ty
};
Type *floatTy = Type::getDoubleTy(Context);
FunctionType *fnTy = FunctionType::get(floatTy, ArgsTy, false);
Function *Frand = CreateF(m, ctxTy, sfpTy);
Function *F = Function::Create(fnTy, GlobalValue::ExternalLinkage, "test", m);
BasicBlock *bb = BasicBlock::Create(Context, "EntryBlock", F);
IRBuilder<> *builder = new IRBuilder<>(bb);
Function::arg_iterator I = F->arg_begin();
Value *Ctx = I++;
Value *Sfp = I++;
Value *Rix = I;
Value *Args[] = {
Ctx, Sfp, Rix
};
Value *v = builder->CreateCall(Frand, Args);
builder->CreateRet(v);
std::cout << "before" << std::endl;
(*m).dump();
ExecutionEngine *ee = EngineBuilder(m).setEngineKind(EngineKind::JIT).create();
PassManager mpm;
mpm.add(createIPSCCPPass());
mpm.add(createFunctionInliningPass());
mpm.add(createLICMPass());
mpm.add(createGVNPass());
mpm.add(createGlobalDCEPass());
mpm.run(*m);
std::cout << std::endl << "before" << std::endl;
(*m).dump();
//{
// void *ptr = ee->getPointerToFunction(F);
// typedef struct ksfp_t {
// double v;
// void *p;
// } ksfp_t;
// typedef float (*F_t)(void *, ksfp_t *, long);
// ksfp_t sfp_[100] = {};
// F_t fptr = (F_t) ptr;
// std::cout << fptr(NULL, sfp_, 0) << std::endl;
// asm volatile("int3");
// std::cout << ((float (*)())ptr)() << std::endl;
//}
return 0;
}
std::pair<Value *, Value *>
AMDGPUPromoteAlloca::getLocalSizeYZ(IRBuilder<> &Builder) {
if (!IsAMDHSA) {
Function *LocalSizeYFn
= Intrinsic::getDeclaration(Mod, Intrinsic::r600_read_local_size_y);
Function *LocalSizeZFn
= Intrinsic::getDeclaration(Mod, Intrinsic::r600_read_local_size_z);
CallInst *LocalSizeY = Builder.CreateCall(LocalSizeYFn, {});
CallInst *LocalSizeZ = Builder.CreateCall(LocalSizeZFn, {});
LocalSizeY->setMetadata(LLVMContext::MD_range, MaxWorkGroupSizeRange);
LocalSizeZ->setMetadata(LLVMContext::MD_range, MaxWorkGroupSizeRange);
return std::make_pair(LocalSizeY, LocalSizeZ);
}
// We must read the size out of the dispatch pointer.
assert(IsAMDGCN);
// We are indexing into this struct, and want to extract the workgroup_size_*
// fields.
//
// typedef struct hsa_kernel_dispatch_packet_s {
// uint16_t header;
// uint16_t setup;
// uint16_t workgroup_size_x ;
// uint16_t workgroup_size_y;
// uint16_t workgroup_size_z;
// uint16_t reserved0;
// uint32_t grid_size_x ;
// uint32_t grid_size_y ;
// uint32_t grid_size_z;
//
// uint32_t private_segment_size;
// uint32_t group_segment_size;
// uint64_t kernel_object;
//
// #ifdef HSA_LARGE_MODEL
// void *kernarg_address;
// #elif defined HSA_LITTLE_ENDIAN
// void *kernarg_address;
// uint32_t reserved1;
// #else
// uint32_t reserved1;
// void *kernarg_address;
// #endif
// uint64_t reserved2;
// hsa_signal_t completion_signal; // uint64_t wrapper
// } hsa_kernel_dispatch_packet_t
//
Function *DispatchPtrFn
= Intrinsic::getDeclaration(Mod, Intrinsic::amdgcn_dispatch_ptr);
CallInst *DispatchPtr = Builder.CreateCall(DispatchPtrFn, {});
DispatchPtr->addAttribute(AttributeSet::ReturnIndex, Attribute::NoAlias);
DispatchPtr->addAttribute(AttributeSet::ReturnIndex, Attribute::NonNull);
// Size of the dispatch packet struct.
DispatchPtr->addDereferenceableAttr(AttributeSet::ReturnIndex, 64);
Type *I32Ty = Type::getInt32Ty(Mod->getContext());
Value *CastDispatchPtr = Builder.CreateBitCast(
DispatchPtr, PointerType::get(I32Ty, AMDGPUAS::CONSTANT_ADDRESS));
// We could do a single 64-bit load here, but it's likely that the basic
// 32-bit and extract sequence is already present, and it is probably easier
// to CSE this. The loads should be mergable later anyway.
Value *GEPXY = Builder.CreateConstInBoundsGEP1_64(CastDispatchPtr, 1);
LoadInst *LoadXY = Builder.CreateAlignedLoad(GEPXY, 4);
Value *GEPZU = Builder.CreateConstInBoundsGEP1_64(CastDispatchPtr, 2);
LoadInst *LoadZU = Builder.CreateAlignedLoad(GEPZU, 4);
MDNode *MD = llvm::MDNode::get(Mod->getContext(), None);
LoadXY->setMetadata(LLVMContext::MD_invariant_load, MD);
LoadZU->setMetadata(LLVMContext::MD_invariant_load, MD);
LoadZU->setMetadata(LLVMContext::MD_range, MaxWorkGroupSizeRange);
// Extract y component. Upper half of LoadZU should be zero already.
Value *Y = Builder.CreateLShr(LoadXY, 16);
return std::make_pair(Y, LoadZU);
}
void ASTCodeGenVisitor::Visit(Output* s) {
IRBuilder<> builder = builders_.top();
Value* output = builder.CreateLoad(ptr_);
builder.CreateCall(put_char_, output);
VisitNextASTNode(s);
}
// =============================================================================
// andOOPIsGone (formerly: createProcess)
//
// Formerly, OOP permitted the same SC_{METHOD,THREAD} functions to apply
// to each copy of a SC_MODULE. Aaaaand it's gone !
// (but OTOH we enable better optimizations)
// Creates a new C-style function that calls the old member function with the
// given sc_module. The call is then inlined.
// FIXME: assumes the method is non-virtual and that sc_module is the first
// inherited class of the SC_MODULE
// =============================================================================
Function *TwetoPassImpl::andOOPIsGone(Function * oldProc,
sc_core::sc_module * initiatorMod)
{
if (!oldProc)
return NULL;
// can't statically optimize if the address of the module isn't predictible
// TODO: also handle already-static variables, which also have
// fixed $pc-relative addresses
if (staticopt == optlevel && !permalloc::is_from (initiatorMod))
return NULL;
LLVMContext & context = getGlobalContext();
FunctionType *funType = oldProc->getFunctionType();
Type *type = funType->getParamType(0);
FunctionType *newProcType =
FunctionType::get(oldProc->getReturnType(),
ArrayRef < Type * >(), false);
// Create the new function
std::ostringstream id;
id << proc_counter++;
std::string name =
oldProc->getName().str() + std::string("_clone_") + id.str();
Function *newProc =
Function::Create(newProcType, Function::ExternalLinkage, name,
this->llvmMod);
assert(newProc->empty());
newProc->addFnAttr(Attribute::InlineHint);
// Create call to old function
BasicBlock *bb = BasicBlock::Create(context, "entry", newProc);
IRBuilder <> *irb = new IRBuilder <> (context);
irb->SetInsertPoint(bb);
Value* thisAddr = createRelocatablePointer (type, initiatorMod, irb);
CallInst *ci = irb->CreateCall(oldProc,
ArrayRef < Value * >(std::vector<Value*>(1,thisAddr)));
//bb->getInstList().insert(ci, thisAddr);
if (ci->getType()->isVoidTy())
irb->CreateRetVoid();
else
irb->CreateRet(ci);
// The function should be valid now
verifyFunction(*newProc);
{ // Inline the call
DataLayout *td = new DataLayout(this->llvmMod);
InlineFunctionInfo i(NULL, td);
bool success = InlineFunction(ci, i);
assert(success);
verifyFunction(*newProc);
}
// further optimize the function
inlineBasicIO (initiatorMod, newProc);
newProc->dump();
return newProc;
}
// =============================================================================
// createProcess
//
// Create a new function that contains a call to the old function.
// We inline the call in order to clone the old function's implementation.
// =============================================================================
Function *TLMBasicPassImpl::createProcess(Function *oldProc,
sc_core::sc_module *initiatorMod) {
LLVMContext &context = getGlobalContext();
IntegerType *intType;
if (this->is64Bit) {
intType = Type::getInt64Ty(context);
} else {
intType = Type::getInt32Ty(context);
}
// Retrieve a pointer to the initiator module
ConstantInt *initiatorModVal =
ConstantInt::getSigned(intType,reinterpret_cast<intptr_t>(initiatorMod));
FunctionType *funType = oldProc->getFunctionType();
Type *type = funType->getParamType(0);
IntToPtrInst *thisAddr =
new IntToPtrInst(initiatorModVal, type, "");
// Compute the type of the new function
FunctionType *oldProcType = oldProc->getFunctionType();
Value **argsBegin = new Value*[1];
Value **argsEnd = argsBegin;
*argsEnd++ = thisAddr;
const unsigned argsSize = argsEnd-argsBegin;
Value **args = argsBegin;
assert(oldProcType->getNumParams()==argsSize);
assert(!oldProc->isDeclaration());
std::vector<Type*> argTypes;
for (unsigned i = 0; i!=argsSize; ++i)
argTypes.push_back(oldProcType->getParamType(i));
FunctionType *newProcType =
FunctionType::get(oldProc->getReturnType(), ArrayRef<Type*>(argTypes), false);
// Create the new function
std::ostringstream id;
id << proc_counter++;
std::string name = oldProc->getName().str()+std::string("_clone_")+id.str();
Function *newProc =
Function::Create(newProcType, Function::ExternalLinkage, name, this->llvmMod);
assert(newProc->empty());
newProc->addFnAttr(Attributes::InlineHint);
{ // Set name of newfunc arguments and complete args
Function::arg_iterator nai = newProc->arg_begin();
Function::arg_iterator oai = oldProc->arg_begin();
for (unsigned i = 0; i!=argsSize; ++i, ++oai) {
nai->setName(oai->getName());
args[i] = nai;
++nai;
}
assert(nai==newProc->arg_end());
assert(oai==oldProc->arg_end());
}
// Create call to old function
BasicBlock *bb = BasicBlock::Create(context, "entry", newProc);
IRBuilder<> *irb = new IRBuilder<>(context);
irb->SetInsertPoint(bb);
CallInst *ci = irb->CreateCall(oldProc, ArrayRef<Value*>(argsBegin, argsEnd));
bb->getInstList().insert(ci, thisAddr);
if (ci->getType()->isVoidTy())
irb->CreateRetVoid();
else
irb->CreateRet(ci);
// The function should be valid now
verifyFunction(*newProc);
{ // Inline the call
DataLayout *td = new DataLayout(this->llvmMod);
InlineFunctionInfo i(NULL, td);
bool success = InlineFunction(ci, i);
assert(success);
verifyFunction(*newProc);
}
//newProc->dump();
return newProc;
}
codegen_value ast::distribution::createConstructor(Module *module, IRBuilder<> &builder,
const string &ctor_name,
Type *parameter_type,
const vector<type_spec> ¶m_type_list,
Value *eval, Value *sample, Value *pdf, Value *emit,
Function *dtor) {
//create function accepting parameters as arguments
vector<Type*> arg_types;
for (auto it = param_type_list.begin(); it != param_type_list.end(); ++it) arg_types.push_back((*it)->llvm_type());
FunctionType *ft = FunctionType::get(state->types["dfunc"]->llvm_type(), arg_types, false);
Function *f = Function::Create(ft, Function::ExternalLinkage, ctor_name, module);
BasicBlock *bb = BasicBlock::Create(getGlobalContext(), "func_entry", f);
builder.SetInsertPoint(bb);
//setup arguments for the alloc call
Value *gd_scene = module->getNamedGlobal(".__gd_scene");
assert(gd_scene != NULL);
Value *scene_ptr = builder.CreateLoad(gd_scene);
//compute the shader flags
codegen_value flag_val = codegen_all_flags(module, builder);
return errors::codegen_call(flag_val, [&] (Value *&flag_bitmask) -> codegen_value {
//get memory for a new distribution object
Value *dfunc_ptr = state->types["dfunc"]->allocate(module, builder);
//initialize the object and dynamically allocate parameter memory (calling a builtin function)
Type* int_ptr_ty = Type::getInt32Ty(getGlobalContext())->getPointerTo();
vector<Type*> alloc_arg_types({state->types["scene_ptr"]->llvm_type(),
Type::getInt32Ty(getGlobalContext()), state->types["shader_flag"]->llvm_type(),
int_ptr_ty, int_ptr_ty, int_ptr_ty, int_ptr_ty,
dtor->getType(), dfunc_ptr->getType()});
FunctionType *alloc_type = FunctionType::get(Type::getInt32PtrTy(getGlobalContext()), alloc_arg_types, false);
Function *alloc_func = GetExternalFunction(module, "gd_builtin_alloc_dfunc", alloc_type);
int param_data_size = DataLayout(module).getTypeAllocSize(parameter_type);
Constant *param_size_arg = ConstantInt::get(getGlobalContext(), APInt(8*sizeof(int), param_data_size));
vector<Value*> alloc_args({scene_ptr, param_size_arg, flag_bitmask,
builder.CreatePointerCast(eval, int_ptr_ty),
builder.CreatePointerCast(sample, int_ptr_ty),
builder.CreatePointerCast(pdf, int_ptr_ty),
builder.CreatePointerCast(emit, int_ptr_ty),
dtor, dfunc_ptr});
Value *param_ptr = builder.CreatePointerCast(builder.CreateCall(alloc_func, alloc_args),
parameter_type->getPointerTo(), "dfunc_param_ptr");
//set each parameter
auto arg_it = f->arg_begin();
unsigned int field_idx = 0;
for (auto it = param_type_list.begin(); it != param_type_list.end(); ++it, ++arg_it, ++field_idx) {
Value *param_copy = (*it)->copy(arg_it, module, builder);
(*it)->store(param_copy, builder.CreateStructGEP(param_ptr, field_idx), module, builder);
}
//return the object
Value *rt_val = builder.CreateLoad(dfunc_ptr, "dist_ref");
builder.CreateRet(rt_val);
return f;
});
}
void ASTCodeGenVisitor::Visit(GetInput* s) {
IRBuilder<> builder = builders_.top();
Value* input = builder.CreateCall(get_char_);
builder.CreateStore(input, ptr_);
VisitNextASTNode(s);
}