gaspersic/triton
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

[General] LLVM-9 -> LLVM-10

Browse Source
This commit is contained in:
Philippe Tillet
2020-11-07 22:35:23 -05:00
committed by Philippe Tillet
parent 8f3ee53f24
commit 50587bbf4b
9 changed files with 21 additions and 31 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
lib/codegen/selection/generator.cc
View File

@@ -695,7 +695,7 @@ void generator::visit_atomic_add_inst(ir::atomic_add_inst* add) {
builder_->CreateCondBr(rmw_msk, mask_then_bb, mask_done_bb);
builder_->SetInsertPoint(mask_then_bb);
builder_->CreateAtomicRMW(AtomicRMWInst::FAdd, rmw_ptr, rmw_val,
AtomicOrdering::Monotonic,
AtomicOrdering::Unordered,
SyncScope::System);
builder_->CreateBr(mask_done_bb);
builder_->SetInsertPoint(mask_done_bb);

10
lib/codegen/target.cc
View File

@@ -2,6 +2,8 @@
#include "llvm/IR/IRBuilder.h"
#include "llvm/IR/Function.h"
#include "llvm/IR/Intrinsics.h"
#include "llvm/IR/IntrinsicsNVPTX.h"
#include "llvm/IR/IntrinsicsAMDGPU.h"
#include "llvm/IR/Value.h"
#include "llvm/IR/IRBuilder.h"
#include <iostream>
@@ -49,13 +51,7 @@ Value* amd_cl_target::get_block_id(Module *module, IRBuilder<>& builder, unsigne
}
Value* amd_cl_target::get_num_blocks(Module *module, IRBuilder<>& builder, unsigned ax) {
static std::array<Intrinsic::ID, 3> ids = {
Intrinsic::r600_read_ngroups_x,
Intrinsic::r600_read_ngroups_y,
Intrinsic::r600_read_ngroups_z
};
Value* get_num_group = Intrinsic::getDeclaration(module, ids[ax]);
return builder.CreateCall(get_num_group, {});
throw std::runtime_error("not implemented on AMD");
}
Value* amd_cl_target::get_local_id(Module *module, IRBuilder<>& builder, unsigned ax) {

8
lib/driver/module.cc
View File

@@ -28,6 +28,8 @@
#include "llvm/IR/Verifier.h"
#include "llvm/IR/IRPrintingPasses.h"
#include "llvm/IR/Module.h"
#include "llvm/Support/CodeGen.h"
#include "llvm/Support/CommandLine.h"
#include "llvm/Support/SourceMgr.h"
#include "llvm/Support/raw_ostream.h"
#include "llvm/Support/TargetRegistry.h"
@@ -115,8 +117,8 @@ void module::compile_llvm_module(std::unique_ptr<llvm::Module> module, const std
// convert triton file type to llvm file type
auto ll_file_type = [&](module::file_type_t type){
if(type == Object)
return llvm::TargetMachine::CGFT_ObjectFile;
return llvm::TargetMachine::CGFT_AssemblyFile;
return llvm::CodeGenFileType::CGFT_ObjectFile;
return llvm::CodeGenFileType::CGFT_AssemblyFile;
};
// emit
machine->addPassesToEmitFile(pass, stream, nullptr, ll_file_type(ft));
@@ -198,7 +200,7 @@ host_module::host_module(driver::context * context, std::unique_ptr<llvm::Module
llvm::EngineBuilder builder(std::move(src));
builder.setErrorStr(&hst_->error);
builder.setMCJITMemoryManager(llvm::make_unique<llvm::SectionMemoryManager>());
builder.setMCJITMemoryManager(std::make_unique<llvm::SectionMemoryManager>());
builder.setOptLevel(llvm::CodeGenOpt::Aggressive);
builder.setEngineKind(llvm::EngineKind::JIT);
hst_->engine = builder.create();

5
lib/lang/code_gen.cc
View File

@@ -278,10 +278,7 @@ void Generator::VisitFuncCall(FuncCall* funcCall) {
ir::value* val = ret_;
return set_ret(bld_->create_atomic_exch(ptr, val));
}
if(name == "f32_atomic_add" ||
name == "atomic_add_32x32" || name == "atomic_add_32x64" || name == "atomic_add_32x128" ||
name == "atomic_add_64x32" || name == "atomic_add_64x64" || name == "atomic_add_64x128" ||
name == "atomic_add_128x32"|| name == "atomic_add_128x64"|| name == "atomic_add_128x128"){
if(name.substr(0, 10) == "atomic_add"){
VisitExpr(funcCall->Args()->at(0));
ir::value* ptr = ret_;
VisitExpr(funcCall->Args()->at(1));

16
lib/runtime/function.cc
View File

@@ -359,17 +359,11 @@ std::string function::preheader() {
#define PASTER(a, b, _) a ## _ ## b
#define EVALUATOR(a, b, _) PASTER(a, b, _)
#define atomic_add(TM, TN) EVALUATOR(atomic_add, EVALUATOR(TM, TN, x), _)
extern void atomic_add_64(float*[64], float[64], bool[64]);
extern void atomic_add_32x32(float*[32, 32], float[32, 32], bool[32, 32]);
extern void atomic_add_32x64(float*[32, 64], float[32, 64], bool[32, 64]);
extern void atomic_add_32x128(float*[32, 128], float[32, 128], bool[32, 128]);
extern void atomic_add_64x32(float*[64, 32], float[64, 32], bool[64, 32]);
extern void atomic_add_64x64(float*[64, 64], float[64, 64], bool[64, 64]);
extern void atomic_add_64x128(float*[64, 128], float[64, 128], bool[64, 128]);
extern void atomic_add_128x32(float*[128, 32], float[128, 32], bool[128, 32]);
extern void atomic_add_128x64(float*[128, 64], float[128, 64], bool[128, 64]);
extern void atomic_add_128x128(float*[128, 128], float[128, 128], bool[128, 128]);
#define atomic_add(TYPE, TM, TN) EVALUATOR(atomic_add, EVALUATOR(TYPE, EVALUATOR(TM, TN, x), _), _)
#define DECLARATION(TYPE, TM, TN) extern void atomic_add(TYPE, TM, TN)(TYPE*[TM, TN], TYPE[TM, TN], bool[TM, TN])
DECLARATION(float, 64, 64);
DECLARATION(half , 64, 64);
extern int atomic_cas(int*, int, int);
extern int atomic_xchg(int*, int);