void FindActionParameters::postorder(const IR::ActionListElement* element) {
auto path = element->getPath();
auto decl = refMap->getDeclaration(path, true);
BUG_CHECK(decl->is<IR::P4Action>(), "%1%: not an action", element);
BUG_CHECK(element->expression->is<IR::MethodCallExpression>(),
"%1%: expected a method call", element->expression);
invocations->bind(decl->to<IR::P4Action>(),
element->expression->to<IR::MethodCallExpression>(), false);
}
const IR::Node* DoRemoveActionParameters::postorder(IR::P4Action* action) {
LOG1("Visiting " << dbp(action));
BUG_CHECK(getParent<IR::P4Control>() || getParent<IR::P4Program>(),
"%1%: unexpected parent %2%", getOriginal(), getContext()->node);
auto result = new IR::IndexedVector<IR::Declaration>();
auto leftParams = new IR::IndexedVector<IR::Parameter>();
auto initializers = new IR::IndexedVector<IR::StatOrDecl>();
auto postamble = new IR::IndexedVector<IR::StatOrDecl>();
auto invocation = invocations->get(getOriginal<IR::P4Action>());
if (invocation == nullptr)
return action;
auto args = invocation->arguments;
ParameterSubstitution substitution;
substitution.populate(action->parameters, args);
bool removeAll = invocations->removeAllParameters(getOriginal<IR::P4Action>());
for (auto p : action->parameters->parameters) {
if (p->direction == IR::Direction::None && !removeAll) {
leftParams->push_back(p);
} else {
auto decl = new IR::Declaration_Variable(p->srcInfo, p->name,
p->annotations, p->type, nullptr);
LOG3("Added declaration " << decl << " annotations " << p->annotations);
result->push_back(decl);
auto arg = substitution.lookup(p);
if (arg == nullptr) {
::error("action %1%: parameter %2% must be bound", invocation, p);
continue;
}
if (p->direction == IR::Direction::In ||
p->direction == IR::Direction::InOut ||
p->direction == IR::Direction::None) {
auto left = new IR::PathExpression(p->name);
auto assign = new IR::AssignmentStatement(arg->srcInfo, left, arg->expression);
initializers->push_back(assign);
}
if (p->direction == IR::Direction::Out ||
p->direction == IR::Direction::InOut) {
auto right = new IR::PathExpression(p->name);
auto assign = new IR::AssignmentStatement(arg->srcInfo, arg->expression, right);
postamble->push_back(assign);
}
}
}
if (result->empty())
return action;
initializers->append(action->body->components);
initializers->append(*postamble);
action->parameters = new IR::ParameterList(action->parameters->srcInfo, *leftParams);
action->body = new IR::BlockStatement(action->body->srcInfo, *initializers);
LOG1("To replace " << dbp(action));
result->push_back(action);
return result;
}
bool
ErrorCodesVisitor::preorder(const IR::Type_Error* errors) {
auto &map = *errorCodesMap;
for (auto m : *errors->getDeclarations()) {
BUG_CHECK(map.find(m) == map.end(), "Duplicate error");
map[m] = map.size();
}
return false;
}
const IR::Node* DoConvertEnums::preorder(IR::Type_Enum* type) {
bool convert = policy->convert(type);
if (!convert)
return type;
unsigned long long count = type->members.size();
unsigned long long width = policy->enumSize(count);
LOG2("Converting enum " << type->name << " to " << "bit<" << width << ">");
BUG_CHECK(count <= (1ULL << width),
"%1%: not enough bits to represent %2%", width, type);
auto r = new EnumRepresentation(type->srcInfo, width);
auto canontype = typeMap->getTypeType(getOriginal(), true);
BUG_CHECK(canontype->is<IR::Type_Enum>(),
"canon type of enum %s is non enum %s?", type, canontype);
repr.emplace(canontype->to<IR::Type_Enum>(), r);
for (auto d : type->members)
r->add(d->name.name);
return nullptr; // delete the declaration
}
/// Given an expression and a destination type, convert ListExpressions
/// that occur within expression to StructInitializerExpression if the
/// destination type matches.
const IR::Expression*
convert(const IR::Expression* expression, const IR::Type* type) {
bool modified = false;
if (auto st = type->to<IR::Type_StructLike>()) {
auto si = new IR::IndexedVector<IR::NamedExpression>();
if (auto le = expression->to<IR::ListExpression>()) {
size_t index = 0;
for (auto f : st->fields) {
auto expr = le->components.at(index);
auto conv = convert(expr, f->type);
auto ne = new IR::NamedExpression(conv->srcInfo, f->name, conv);
si->push_back(ne);
index++;
}
auto result = new IR::StructInitializerExpression(
expression->srcInfo, st->name, *si, st->is<IR::Type_Header>());
return result;
} else if (auto sli = expression->to<IR::StructInitializerExpression>()) {
for (auto f : st->fields) {
auto ne = sli->components.getDeclaration<IR::NamedExpression>(f->name.name);
BUG_CHECK(ne != nullptr, "%1%: no initializer for %2%", expression, f);
auto convNe = convert(ne->expression, f->type);
if (convNe != ne->expression)
modified = true;
ne = new IR::NamedExpression(ne->srcInfo, f->name, convNe);
si->push_back(ne);
}
if (modified) {
auto result = new IR::StructInitializerExpression(
expression->srcInfo, st->name, *si, st->is<IR::Type_Header>());
return result;
}
}
} else if (auto tup = type->to<IR::Type_Tuple>()) {
auto le = expression->to<IR::ListExpression>();
if (le == nullptr)
return expression;
auto vec = new IR::Vector<IR::Expression>();
for (size_t i = 0; i < le->size(); i++) {
auto expr = le->components.at(i);
auto type = tup->components.at(i);
auto conv = convert(expr, type);
vec->push_back(conv);
modified |= (conv != expr);
}
if (modified) {
auto result = new IR::ListExpression(expression->srcInfo, *vec);
return result;
}
}
return expression;
}
void ControlBodyTranslator::compileEmit(const IR::Vector<IR::Argument>* args) {
BUG_CHECK(args->size() == 1, "%1%: expected 1 argument for emit", args);
auto expr = args->at(0)->expression;
auto type = typeMap->getType(expr);
auto ht = type->to<IR::Type_Header>();
if (ht == nullptr) {
::error("Cannot emit a non-header type %1%", expr);
return;
}
auto program = control->program;
builder->emitIndent();
builder->append("if (");
visit(expr);
builder->append(".ebpf_valid) ");
builder->blockStart();
unsigned width = ht->width_bits();
builder->emitIndent();
builder->appendFormat("if (%s < %s + BYTES(%s + %d)) ",
program->packetEndVar.c_str(),
program->packetStartVar.c_str(),
program->offsetVar.c_str(), width);
builder->blockStart();
builder->emitIndent();
builder->appendFormat("%s = %s;", program->errorVar.c_str(),
p4lib.packetTooShort.str());
builder->newline();
builder->emitIndent();
builder->appendFormat("return %s;", builder->target->abortReturnCode().c_str());
builder->newline();
builder->blockEnd(true);
unsigned alignment = 0;
for (auto f : ht->fields) {
auto ftype = typeMap->getType(f);
auto etype = EBPFTypeFactory::instance->create(ftype);
auto et = dynamic_cast<IHasWidth*>(etype);
if (et == nullptr) {
::error("Only headers with fixed widths supported %1%", f);
return;
}
compileEmitField(expr, f->name, alignment, etype);
alignment += et->widthInBits();
alignment %= 8;
}
builder->blockEnd(true);
return;
}
bool ControlBodyTranslator::preorder(const IR::SwitchStatement* statement) {
cstring newName = control->program->refMap->newName("action_run");
saveAction.push_back(newName);
// This must be a table.apply().action_run
auto mem = statement->expression->to<IR::Member>();
BUG_CHECK(mem != nullptr,
"%1%: Unexpected expression in switch statement", statement->expression);
visit(mem->expr);
saveAction.pop_back();
saveAction.push_back(nullptr);
builder->emitIndent();
builder->append("switch (");
builder->append(newName);
builder->append(") ");
builder->blockStart();
for (auto c : statement->cases) {
builder->emitIndent();
if (c->label->is<IR::DefaultExpression>()) {
builder->append("default");
} else {
builder->append("case ");
auto pe = c->label->to<IR::PathExpression>();
auto decl = control->program->refMap->getDeclaration(pe->path, true);
BUG_CHECK(decl->is<IR::P4Action>(), "%1%: expected an action", pe);
auto act = decl->to<IR::P4Action>();
cstring name = EBPFObject::externalName(act);
builder->append(name);
}
builder->append(":");
builder->newline();
builder->emitIndent();
visit(c->statement);
builder->newline();
builder->emitIndent();
builder->appendLine("break;");
}
builder->blockEnd(false);
saveAction.pop_back();
return false;
}
const IR::IDeclaration* ReferenceMap::getDeclaration(const IR::Path* path, bool notNull) const {
CHECK_NULL(path);
auto result = get(pathToDeclaration, path);
if (result)
LOG1("Looking up " << path << " found " << result->getNode());
else
LOG1("Looking up " << path << " found nothing");
if (notNull)
BUG_CHECK(result != nullptr, "Cannot find declaration for %1%", path);
return result;
}
void DoCheckConstants::postorder(const IR::MethodCallExpression* expression) {
auto mi = MethodInstance::resolve(expression, refMap, typeMap);
if (auto bi = mi->to<BuiltInMethod>()) {
if (bi->name == IR::Type_Stack::push_front ||
bi->name == IR::Type_Stack::pop_front) {
BUG_CHECK(expression->arguments->size() == 1,
"Expected 1 argument for %1%", expression);
auto arg0 = expression->arguments->at(0)->expression;
if (!arg0->is<IR::Constant>())
::error("%1%: argument must be a constant", arg0);
}
}
}
const IR::Node* TypeVariableSubstitutionVisitor::preorder(IR::TypeParameters *tps) {
// remove all variables that were substituted
for (auto it = tps->parameters.begin(); it != tps->parameters.end();) {
const IR::Type* type = bindings->lookup(*it);
if (type != nullptr && !replace) {
LOG3("Removing from generic parameters " << *it);
it = tps->parameters.erase(it);
} else {
if (type != nullptr)
BUG_CHECK(type->is<IR::Type_Var>(),
"cannot replace a type parameter %1% with %2%:", *it, type);
++it;
}
}
return tps;
}
void EBPFControl::emitDeclaration(CodeBuilder* builder, const IR::Declaration* decl) {
if (decl->is<IR::Declaration_Variable>()) {
auto vd = decl->to<IR::Declaration_Variable>();
auto etype = EBPFTypeFactory::instance->create(vd->type);
builder->emitIndent();
etype->declare(builder, vd->name, false);
builder->endOfStatement(true);
BUG_CHECK(vd->initializer == nullptr,
"%1%: declarations with initializers not supported", decl);
return;
} else if (decl->is<IR::P4Table>() ||
decl->is<IR::P4Action>() ||
decl->is<IR::Declaration_Instance>()) {
return;
}
BUG("%1%: not yet handled", decl);
}
const IR::StructInitializerExpression* StructTypeReplacement::explode(
const IR::Expression *root, cstring prefix) {
auto vec = new IR::IndexedVector<IR::NamedExpression>();
auto fieldType = ::get(structFieldMap, prefix);
BUG_CHECK(fieldType, "No field for %1%", prefix);
for (auto f : fieldType->fields) {
cstring fieldName = prefix + "." + f->name.name;
auto newFieldname = ::get(fieldNameRemap, fieldName);
const IR::Expression* expr;
if (!newFieldname.isNullOrEmpty()) {
expr = new IR::Member(root, newFieldname);
} else {
expr = explode(root, fieldName);
}
vec->push_back(new IR::NamedExpression(f->name, expr));
}
return new IR::StructInitializerExpression(fieldType->name, *vec, false);
}
const IR::Node* DoResetHeaders::postorder(IR::Declaration_Variable* decl) {
if (findContext<IR::ParserState>() == nullptr)
return decl;
if (decl->initializer != nullptr)
return decl;
auto resets = new IR::Vector<IR::StatOrDecl>();
resets->push_back(decl);
BUG_CHECK(getContext()->node->is<IR::Vector<IR::StatOrDecl>>() ||
getContext()->node->is<IR::ParserState>() ||
getContext()->node->is<IR::BlockStatement>(),
"%1%: parent is not Vector<StatOrDecl>, but %2%",
decl, getContext()->node);
auto type = typeMap->getType(getOriginal(), true);
auto path = new IR::PathExpression(decl->getName());
generateResets(typeMap, type, path, resets);
if (resets->size() == 1)
return decl;
return resets;
}
const IR::Node* DoBindTypeVariables::postorder(IR::MethodCallExpression* expression) {
if (!expression->typeArguments->empty())
return expression;
auto type = typeMap->getType(expression->method, true);
BUG_CHECK(type->is<IR::IMayBeGenericType>(), "%1%: unexpected type %2% for method",
expression->method, type);
auto mt = type->to<IR::IMayBeGenericType>();
if (mt->getTypeParameters()->empty())
return expression;
auto typeArgs = new IR::Vector<IR::Type>();
for (auto p : mt->getTypeParameters()->parameters) {
auto type = getVarValue(p, expression);
if (type == nullptr)
return expression;
typeArgs->push_back(type);
}
expression->typeArguments = typeArgs;
return expression;
}
const IR::Node* DoBindTypeVariables::postorder(IR::Declaration_Instance* decl) {
if (decl->type->is<IR::Type_Specialized>())
return decl;
auto type = typeMap->getType(getOriginal(), true);
BUG_CHECK(type->is<IR::IMayBeGenericType>(), "%1%: unexpected type %2% for declaration",
decl, type);
auto mt = type->to<IR::IMayBeGenericType>();
if (mt->getTypeParameters()->empty())
return decl;
auto typeArgs = new IR::Vector<IR::Type>();
for (auto p : mt->getTypeParameters()->parameters) {
auto type = getVarValue(p, decl);
if (type == nullptr)
return decl;
typeArgs->push_back(type);
}
decl->type = new IR::Type_Specialized(
decl->type->srcInfo, decl->type->to<IR::Type_Name>(), typeArgs);
return decl;
}
const IR::Node* DoBindTypeVariables::postorder(IR::ConstructorCallExpression* expression) {
if (expression->constructedType->is<IR::Type_Specialized>())
return expression;
auto type = typeMap->getType(getOriginal(), true);
BUG_CHECK(type->is<IR::IMayBeGenericType>(), "%1%: unexpected type %2% for expression",
expression, type);
auto mt = type->to<IR::IMayBeGenericType>();
if (mt->getTypeParameters()->empty())
return expression;
auto typeArgs = new IR::Vector<IR::Type>();
for (auto p : mt->getTypeParameters()->parameters) {
auto type = getVarValue(p, expression);
if (type == nullptr)
return expression;
typeArgs->push_back(type);
}
expression->constructedType = new IR::Type_Specialized(
expression->constructedType->srcInfo,
expression->constructedType->to<IR::Type_Name>(), typeArgs);
return expression;
}
clfftStatus FFTAction::enqueue(clfftPlanHandle plHandle,
clfftDirection dir,
cl_uint numQueuesAndEvents,
cl_command_queue* commQueues,
cl_uint numWaitEvents,
const cl_event* waitEvents,
cl_event* outEvents,
cl_mem* clInputBuffers,
cl_mem* clOutputBuffers)
{
FFTRepo & fftRepo = FFTRepo::getInstance();
std::vector< cl_mem > inputBuff;
std::vector< cl_mem > outputBuff;
clfftStatus status = selectBufferArguments(this->plan,
clInputBuffers, clOutputBuffers,
inputBuff, outputBuff);
if (status != CLFFT_SUCCESS)
{
return status;
}
// TODO: In the case of length == 1, FFT is a trivial NOP, but we still need to apply the forward and backwards tranforms
// TODO: Are map lookups expensive to call here? We can cache a pointer to the cl_program/cl_kernel in the plan
// Translate the user plan into the structure that we use to map plans to clPrograms
cl_program prog;
cl_kernel kern;
OPENCL_V( fftRepo.getclProgram( this->getGenerator(), this->getSignatureData(), prog, this->plan->bakeDevice, this->plan->context ), _T( "fftRepo.getclProgram failed" ) );
OPENCL_V( fftRepo.getclKernel( prog, dir, kern ), _T( "fftRepo.getclKernels failed" ) );
cl_uint uarg = 0;
if (!this->plan->transflag && !(this->plan->gen == Copy))
{
// ::clSetKernelArg() is not thread safe, according to the openCL spec for the same cl_kernel object
// TODO: Need to verify that two different plans (which would get through our lock above) with exactly the same
// parameters would NOT share the same cl_kernel objects
/* constant buffer */
OPENCL_V( clSetKernelArg( kern, uarg++, sizeof( cl_mem ), (void*)&this->plan->const_buffer ), _T( "clSetKernelArg failed" ) );
}
// Input buffer(s)
// Input may be 1 buffer (CLFFT_COMPLEX_INTERLEAVED)
// or 2 buffers (CLFFT_COMPLEX_PLANAR)
for (size_t i = 0; i < inputBuff.size(); ++i)
{
OPENCL_V( clSetKernelArg( kern, uarg++, sizeof( cl_mem ), (void*)&inputBuff[i] ), _T( "clSetKernelArg failed" ) );
}
// Output buffer(s)
// Output may be 0 buffers (CLFFT_INPLACE)
// or 1 buffer (CLFFT_COMPLEX_INTERLEAVED)
// or 2 buffers (CLFFT_COMPLEX_PLANAR)
for (size_t o = 0; o < outputBuff.size(); ++o)
{
OPENCL_V( clSetKernelArg( kern, uarg++, sizeof( cl_mem ), (void*)&outputBuff[o] ), _T( "clSetKernelArg failed" ) );
}
//If callback function is set for the plan, pass the appropriate aruments
if (this->plan->hasPreCallback || this->plan->hasPostCallback)
{
if (this->plan->hasPreCallback)
{
OPENCL_V( clSetKernelArg( kern, uarg++, sizeof( cl_mem ), (void*)&this->plan->precallUserData ), _T( "clSetKernelArg failed" ) );
}
//If post-callback function is set for the plan, pass the appropriate aruments
if (this->plan->hasPostCallback)
{
OPENCL_V( clSetKernelArg( kern, uarg++, sizeof( cl_mem ), (void*)&this->plan->postcallUserData ), _T( "clSetKernelArg failed" ) );
}
//Pass LDS size arument if set
if ((this->plan->hasPreCallback && this->plan->preCallback.localMemSize > 0) ||
(this->plan->hasPostCallback && this->plan->postCallbackParam.localMemSize > 0))
{
int localmemSize = 0;
if (this->plan->hasPreCallback && this->plan->preCallback.localMemSize > 0)
localmemSize = this->plan->preCallback.localMemSize;
if (this->plan->hasPostCallback && this->plan->postCallbackParam.localMemSize > 0)
localmemSize += this->plan->postCallbackParam.localMemSize;
OPENCL_V( clSetKernelArg( kern, uarg++, localmemSize, NULL ), _T( "clSetKernelArg failed" ) );
}
}
std::vector< size_t > gWorkSize;
std::vector< size_t > lWorkSize;
clfftStatus result = this->getWorkSizes (gWorkSize, lWorkSize);
// TODO: if getWorkSizes returns CLFFT_INVALID_GLOBAL_WORK_SIZE, that means
// that this multidimensional input data array is too large to be transformed
// with a single call to clEnqueueNDRangeKernel. For now, we will just return
// the error code back up the call stack.
// The *correct* course of action would be to split the work into mutliple
// calls to clEnqueueNDRangeKernel.
if (CLFFT_INVALID_GLOBAL_WORK_SIZE == result)
{
OPENCL_V( result, _T("Work size too large for clEnqueNDRangeKernel()"));
}
else
{
OPENCL_V( result, _T("FFTAction::getWorkSizes failed"));
}
BUG_CHECK (gWorkSize.size() == lWorkSize.size());
cl_int call_status = clEnqueueNDRangeKernel( *commQueues, kern, static_cast< cl_uint >( gWorkSize.size( ) ),
NULL, &gWorkSize[ 0 ], &lWorkSize[ 0 ], numWaitEvents, waitEvents, outEvents );
OPENCL_V( call_status, _T( "clEnqueueNDRangeKernel failed" ) );
if( fftRepo.pStatTimer )
{
fftRepo.pStatTimer->AddSample( plHandle, this->plan, kern, numQueuesAndEvents, outEvents, gWorkSize, lWorkSize );
}
return CLFFT_SUCCESS;
}
void ControlBodyTranslator::processApply(const P4::ApplyMethod* method) {
builder->emitIndent();
auto table = control->getTable(method->object->getName().name);
BUG_CHECK(table != nullptr, "No table for %1%", method->expr);
P4::ParameterSubstitution binding;
cstring actionVariableName;
if (!saveAction.empty()) {
actionVariableName = saveAction.at(saveAction.size() - 1);
if (!actionVariableName.isNullOrEmpty()) {
builder->appendFormat("enum %s %s;\n",
table->actionEnumName.c_str(), actionVariableName.c_str());
builder->emitIndent();
}
}
builder->blockStart();
BUG_CHECK(method->expr->arguments->size() == 0, "%1%: table apply with arguments", method);
cstring keyname = "key";
if (table->keyGenerator != nullptr) {
builder->emitIndent();
builder->appendLine("/* construct key */");
builder->emitIndent();
builder->appendFormat("struct %s %s = {}", table->keyTypeName.c_str(), keyname.c_str());
builder->endOfStatement(true);
table->emitKey(builder, keyname);
}
builder->emitIndent();
builder->appendLine("/* value */");
builder->emitIndent();
cstring valueName = "value";
builder->appendFormat("struct %s *%s = NULL", table->valueTypeName.c_str(), valueName.c_str());
builder->endOfStatement(true);
if (table->keyGenerator != nullptr) {
builder->emitIndent();
builder->appendLine("/* perform lookup */");
builder->emitIndent();
builder->target->emitTableLookup(builder, table->dataMapName, keyname, valueName);
builder->endOfStatement(true);
}
builder->emitIndent();
builder->appendFormat("if (%s == NULL) ", valueName.c_str());
builder->blockStart();
builder->emitIndent();
builder->appendLine("/* miss; find default action */");
builder->emitIndent();
builder->appendFormat("%s = 0", control->hitVariable.c_str());
builder->endOfStatement(true);
builder->emitIndent();
builder->target->emitTableLookup(builder, table->defaultActionMapName,
control->program->zeroKey, valueName);
builder->endOfStatement(true);
builder->blockEnd(false);
builder->append(" else ");
builder->blockStart();
builder->emitIndent();
builder->appendFormat("%s = 1", control->hitVariable.c_str());
builder->endOfStatement(true);
builder->blockEnd(true);
builder->emitIndent();
builder->appendFormat("if (%s != NULL) ", valueName.c_str());
builder->blockStart();
builder->emitIndent();
builder->appendLine("/* run action */");
table->emitAction(builder, valueName);
if (!actionVariableName.isNullOrEmpty()) {
builder->emitIndent();
builder->appendFormat("%s = %s->action",
actionVariableName.c_str(), valueName.c_str());
builder->endOfStatement(true);
}
toDereference.clear();
builder->blockEnd(true);
builder->emitIndent();
builder->appendFormat("else return %s", builder->target->abortReturnCode().c_str());
builder->endOfStatement(true);
builder->blockEnd(true);
}
void ControlBodyTranslator::compileEmitField(const IR::Expression* expr, cstring field,
unsigned alignment, EBPFType* type) {
unsigned widthToEmit = dynamic_cast<IHasWidth*>(type)->widthInBits();
cstring swap = "";
if (widthToEmit == 16)
swap = "htons";
else if (widthToEmit == 32)
swap = "htonl";
if (!swap.isNullOrEmpty()) {
builder->emitIndent();
visit(expr);
builder->appendFormat(".%s = %s(", field.c_str(), swap);
visit(expr);
builder->appendFormat(".%s)", field.c_str());
builder->endOfStatement(true);
}
auto program = control->program;
unsigned bitsInFirstByte = widthToEmit % 8;
if (bitsInFirstByte == 0) bitsInFirstByte = 8;
unsigned bitsInCurrentByte = bitsInFirstByte;
unsigned left = widthToEmit;
for (unsigned i=0; i < (widthToEmit + 7) / 8; i++) {
builder->emitIndent();
builder->appendFormat("%s = ((char*)(&", program->byteVar.c_str());
visit(expr);
builder->appendFormat(".%s))[%d]", field.c_str(), i);
builder->endOfStatement(true);
unsigned freeBits = alignment == 0 ? (8 - alignment) : 8;
unsigned bitsToWrite = bitsInCurrentByte > freeBits ? freeBits : bitsInCurrentByte;
BUG_CHECK((bitsToWrite > 0) && (bitsToWrite <= 8), "invalid bitsToWrite %d", bitsToWrite);
builder->emitIndent();
if (alignment == 0)
builder->appendFormat("write_byte(%s, BYTES(%s) + %d, (%s) << %d)",
program->packetStartVar.c_str(), program->offsetVar.c_str(), i,
program->byteVar.c_str(), 8 - bitsToWrite);
else
builder->appendFormat("write_partial(%s + BYTES(%s) + %d, %d, (%s) << %d)",
program->packetStartVar.c_str(), program->offsetVar.c_str(), i,
alignment,
program->byteVar.c_str(), 8 - bitsToWrite);
builder->endOfStatement(true);
left -= bitsToWrite;
bitsInCurrentByte -= bitsToWrite;
if (bitsInCurrentByte > 0) {
builder->emitIndent();
builder->appendFormat(
"write_byte(%s, BYTES(%s) + %d + 1, (%s << %d))",
program->packetStartVar.c_str(),
program->offsetVar.c_str(), i, program->byteVar.c_str(), 8 - alignment % 8);
builder->endOfStatement(true);
left -= bitsInCurrentByte;
}
alignment = (alignment + bitsToWrite) % 8;
bitsInCurrentByte = left >= 8 ? 8 : left;
}
builder->emitIndent();
builder->appendFormat("%s += %d", program->offsetVar.c_str(), widthToEmit);
builder->endOfStatement(true);
}
bool ControlBodyTranslator::preorder(const IR::MethodCallExpression* expression) {
builder->append("/* ");
visit(expression->method);
builder->append("(");
bool first = true;
for (auto a : *expression->arguments) {
if (!first)
builder->append(", ");
first = false;
visit(a);
}
builder->append(")");
builder->append("*/");
builder->newline();
auto mi = P4::MethodInstance::resolve(expression,
control->program->refMap,
control->program->typeMap);
auto apply = mi->to<P4::ApplyMethod>();
if (apply != nullptr) {
processApply(apply);
return false;
}
auto ef = mi->to<P4::ExternFunction>();
if (ef != nullptr) {
processFunction(ef);
return false;
}
auto ext = mi->to<P4::ExternMethod>();
if (ext != nullptr) {
processMethod(ext);
return false;
}
auto bim = mi->to<P4::BuiltInMethod>();
if (bim != nullptr) {
builder->emitIndent();
if (bim->name == IR::Type_Header::isValid) {
visit(bim->appliedTo);
builder->append(".ebpf_valid");
return false;
} else if (bim->name == IR::Type_Header::setValid) {
visit(bim->appliedTo);
builder->append(".ebpf_valid = true");
return false;
} else if (bim->name == IR::Type_Header::setInvalid) {
visit(bim->appliedTo);
builder->append(".ebpf_valid = false");
return false;
}
}
auto ac = mi->to<P4::ActionCall>();
if (ac != nullptr) {
// Action arguments have been eliminated by the mid-end.
BUG_CHECK(expression->arguments->size() == 0,
"%1%: unexpected arguments for action call", expression);
visit(ac->action->body);
return false;
}
::error("Unsupported method invocation %1%", expression);
return false;
}
const IR::Node* DoMoveActionsToTables::postorder(IR::MethodCallStatement* statement) {
auto mi = MethodInstance::resolve(statement, refMap, typeMap);
if (!mi->is<ActionCall>())
return statement;
auto ac = mi->to<ActionCall>();
auto action = ac->action;
auto directionArgs = new IR::Vector<IR::Argument>();
auto mc = statement->methodCall;
// TODO: should use argument names
auto it = action->parameters->parameters.begin();
auto arg = mc->arguments->begin();
for (; it != action->parameters->parameters.end(); ++it) {
auto p = *it;
if (p->direction == IR::Direction::None)
break;
directionArgs->push_back(*arg);
++arg;
}
// Action invocation
BUG_CHECK(ac->expr->method->is<IR::PathExpression>(),
"%1%: Expected a PathExpression", ac->expr->method);
auto actionPath = new IR::PathExpression(IR::ID(mc->srcInfo, ac->action->name));
auto call = new IR::MethodCallExpression(mc->srcInfo, actionPath,
new IR::Vector<IR::Type>(), directionArgs);
auto actinst = new IR::ActionListElement(statement->srcInfo, call);
// Action list property
auto actlist = new IR::ActionList({actinst});
auto prop = new IR::Property(
IR::ID(IR::TableProperties::actionsPropertyName, nullptr),
actlist, false);
// default action property
auto otherArgs = new IR::Vector<IR::Argument>();
for (; it != action->parameters->parameters.end(); ++it) {
otherArgs->push_back(*arg);
++arg;
}
BUG_CHECK(arg == mc->arguments->end(), "%1%: mismatched arguments", mc);
auto amce = new IR::MethodCallExpression(mc->srcInfo, mc->method, mc->typeArguments, otherArgs);
auto defactval = new IR::ExpressionValue(amce);
auto defprop = new IR::Property(
IR::ID(IR::TableProperties::defaultActionPropertyName, nullptr),
defactval, true);
// List of table properties
auto props = new IR::TableProperties({ prop, defprop });
// Synthesize a new table
cstring tblName = IR::ID(refMap->newName(cstring("tbl_") + ac->action->name.name), nullptr);
auto annos = new IR::Annotations();
annos->add(new IR::Annotation(IR::Annotation::hiddenAnnotation, {}));
auto tbl = new IR::P4Table(tblName, annos, props);
tables.push_back(tbl);
// Table invocation statement
auto tblpath = new IR::PathExpression(tblName);
auto method = new IR::Member(tblpath, IR::IApply::applyMethodName);
auto mce = new IR::MethodCallExpression(
statement->srcInfo, method, new IR::Vector<IR::Type>(),
new IR::Vector<IR::Argument>());
auto stat = new IR::MethodCallStatement(mce->srcInfo, mce);
return stat;
}
