lib/SpecConstant.cpp - master - redgpu/clspv

lib/SpecConstant.cpp
183 строки7 KB
James Price
Use distinct nodes for clspv.next_spec_constant_id (#696)
07 янв 2021, 19:44
07 янв 2021, 19:442cf925c
// Copyright 2020 The Clspv Authors. All rights reserved.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
//     http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.

#include "llvm/IR/Constants.h"

#include "Constants.h"
#include "SpecConstant.h"

using namespace llvm;

namespace {

void InitSpecConstantMetadata(Module *module) {
  auto next_spec_id_md =
      module->getOrInsertNamedMetadata(clspv::NextSpecConstantMetadataName());
  next_spec_id_md->clearOperands();

  // Start at 3 to accommodate workgroup size ids.
  const uint32_t first_spec_id = 3;
  auto id_const = ValueAsMetadata::getConstant(ConstantInt::get(
      IntegerType::get(module->getContext(), 32), first_spec_id));
  auto id_md = MDTuple::getDistinct(module->getContext(), {id_const});
  next_spec_id_md->addOperand(id_md);

  auto spec_constant_list_md =
      module->getOrInsertNamedMetadata(clspv::SpecConstantMetadataName());
  spec_constant_list_md->clearOperands();
}

} // namespace

namespace clspv {

const char *GetSpecConstantName(SpecConstant kind) {
  switch (kind) {
  case SpecConstant::kWorkgroupSizeX:
    return "workgroup_size_x";
  case SpecConstant::kWorkgroupSizeY:
    return "workgroup_size_y";
  case SpecConstant::kWorkgroupSizeZ:
    return "workgroup_size_z";
  case SpecConstant::kLocalMemorySize:
    return "local_memory_size";
  case SpecConstant::kWorkDim:
    return "work_dim";
  case SpecConstant::kGlobalOffsetX:
    return "global_offset_x";
  case SpecConstant::kGlobalOffsetY:
    return "global_offset_y";
  case SpecConstant::kGlobalOffsetZ:
    return "global_offset_z";
  }
  llvm::errs() << "Unhandled case in clspv::GetSpecConstantName: " << int(kind)
               << "\n";
  return "";
}

SpecConstant GetSpecConstantFromName(const std::string &name) {
  if (name == "workgroup_size_x")
    return SpecConstant::kWorkgroupSizeX;
  else if (name == "workgroup_size_y")
    return SpecConstant::kWorkgroupSizeY;
  else if (name == "workgroup_size_z")
    return SpecConstant::kWorkgroupSizeZ;
  else if (name == "local_memory_size")
    return SpecConstant::kLocalMemorySize;
  else if (name == "work_dim")
    return SpecConstant::kWorkDim;
  else if (name == "global_offset_x")
    return SpecConstant::kGlobalOffsetX;
  else if (name == "global_offset_y")
    return SpecConstant::kGlobalOffsetY;
  else if (name == "global_offset_z")
    return SpecConstant::kGlobalOffsetZ;

  llvm::errs() << "Unhandled csae in clspv::GetSpecConstantFromName: " << name
               << "\n";
  return SpecConstant::kWorkgroupSizeX;
}

void AddWorkgroupSpecConstants(Module *module) {
  auto spec_constant_list_md =
      module->getNamedMetadata(SpecConstantMetadataName());
  if (!spec_constant_list_md) {
    InitSpecConstantMetadata(module);
    spec_constant_list_md =
        module->getNamedMetadata(SpecConstantMetadataName());
  }

  // Workgroup size spec constants always occupy ids 0, 1 and 2.
  auto enum_const = ValueAsMetadata::getConstant(
      ConstantInt::get(IntegerType::get(module->getContext(), 32),
                       static_cast<uint64_t>(SpecConstant::kWorkgroupSizeX)));
  auto id_const = ValueAsMetadata::getConstant(
      ConstantInt::get(IntegerType::get(module->getContext(), 32), 0));
  auto wg_md = MDTuple::get(module->getContext(), {enum_const, id_const});
  spec_constant_list_md->addOperand(wg_md);

  enum_const = ValueAsMetadata::getConstant(
      ConstantInt::get(IntegerType::get(module->getContext(), 32),
                       static_cast<uint64_t>(SpecConstant::kWorkgroupSizeY)));
  id_const = ValueAsMetadata::getConstant(
      ConstantInt::get(IntegerType::get(module->getContext(), 32), 1));
  wg_md = MDTuple::get(module->getContext(), {enum_const, id_const});
  spec_constant_list_md->addOperand(wg_md);

  enum_const = ValueAsMetadata::getConstant(
      ConstantInt::get(IntegerType::get(module->getContext(), 32),
                       static_cast<uint64_t>(SpecConstant::kWorkgroupSizeZ)));
  id_const = ValueAsMetadata::getConstant(
      ConstantInt::get(IntegerType::get(module->getContext(), 32), 2));
  wg_md = MDTuple::get(module->getContext(), {enum_const, id_const});
  spec_constant_list_md->addOperand(wg_md);
}

uint32_t AllocateSpecConstant(Module *module, SpecConstant kind) {
  auto spec_constant_id_md =
      module->getNamedMetadata(NextSpecConstantMetadataName());
  if (!spec_constant_id_md) {
    InitSpecConstantMetadata(module);
    spec_constant_id_md =
        module->getNamedMetadata(NextSpecConstantMetadataName());
  }

  auto value_md = spec_constant_id_md->getOperand(0);
  auto value = cast<ConstantInt>(
      dyn_cast<ValueAsMetadata>(value_md->getOperand(0))->getValue());
  uint32_t next_id = static_cast<uint32_t>(value->getZExtValue());
  // Update the next available id.
  value_md->replaceOperandWith(
      0, ValueAsMetadata::getConstant(ConstantInt::get(
             IntegerType::get(module->getContext(), 32), next_id + 1)));

  // Add the allocation to the metadata list.
  auto spec_constant_list_md =
      module->getNamedMetadata(SpecConstantMetadataName());
  auto enum_const = ValueAsMetadata::getConstant(ConstantInt::get(
      IntegerType::get(module->getContext(), 32), static_cast<uint64_t>(kind)));
  auto id_const = ValueAsMetadata::getConstant(
      ConstantInt::get(IntegerType::get(module->getContext(), 32), next_id));
  auto wg_md = MDTuple::get(module->getContext(), {enum_const, id_const});
  spec_constant_list_md->addOperand(wg_md);

  return next_id;
}

std::vector<std::pair<SpecConstant, uint32_t>>
GetSpecConstants(Module *module) {
  std::vector<std::pair<SpecConstant, uint32_t>> spec_constants;
  auto spec_constant_md =
      module->getNamedMetadata(clspv::SpecConstantMetadataName());
  if (!spec_constant_md)
    return spec_constants;

  for (auto pair : spec_constant_md->operands()) {
    // Metadata is formatted as pairs of <SpecConstant, id>.
    auto kind = static_cast<SpecConstant>(
        cast<ConstantInt>(
            cast<ValueAsMetadata>(pair->getOperand(0))->getValue())
            ->getZExtValue());

    uint32_t spec_id = static_cast<uint32_t>(
        cast<ConstantInt>(
            cast<ValueAsMetadata>(pair->getOperand(1))->getValue())
            ->getZExtValue());
    spec_constants.emplace_back(kind, spec_id);
  }

  return spec_constants;
}

} // namespace clspv