OSchip
/
llvm-project
mirror of https://github.com/THU-DSP-LAB/llvm-project.git
13
0
Fork 2
Code Issues Packages Projects Releases Wiki Activity

[VENTUS][NFC] Removed old implementations of some builtin functions

This commit is contained in:
Jules-Kong 2025-03-13 17:32:54 +08:00
parent f8c76c62ae
commit 6f23d8d54a
2 changed files with 0 additions and 49 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

17
clang/include/clang/Basic/BuiltinsRISCV.def
View File

@ -22,24 +22,7 @@
// Barrier function builtins
TARGET_BUILTIN(barrier, "vi.", "n", "32bit")
TARGET_BUILTIN(work_group_barrier, "vi.", "n", "32bit")
// WORKAROUND: Disabled for now.
/*
BUILTIN(__builtin_riscv_workgroup_id_x, "Ui", "nc")
BUILTIN(__builtin_riscv_workgroup_id_y, "Ui", "nc")
BUILTIN(__builtin_riscv_workgroup_id_z, "Ui", "nc")
BUILTIN(__builtin_riscv_workitem_id_x, "Ui", "nc")
BUILTIN(__builtin_riscv_workitem_id_y, "Ui", "nc")
BUILTIN(__builtin_riscv_workitem_id_z, "Ui", "nc")
BUILTIN(__builtin_riscv_workgroup_size_x, "Us", "nc")
BUILTIN(__builtin_riscv_workgroup_size_y, "Us", "nc")
BUILTIN(__builtin_riscv_workgroup_size_z, "Us", "nc")
BUILTIN(__builtin_riscv_grid_size_x, "Ui", "nc")
BUILTIN(__builtin_riscv_grid_size_y, "Ui", "nc")
BUILTIN(__builtin_riscv_grid_size_z, "Ui", "nc")
*/
//===----------------------------------------------------------------------===//
// Standard RISCV instruction builtins.
//===----------------------------------------------------------------------===//

32
clang/lib/CodeGen/CGBuiltin.cpp
View File

@ -19461,34 +19461,6 @@ Value *CodeGenFunction::EmitRISCVBuiltinExpr(unsigned BuiltinID,
switch (BuiltinID) {
default: llvm_unreachable("unexpected builtin ID");
// WORKAROUND: Disabled for now.
/*
case RISCV::BI__builtin_riscv_workitem_id_x:
return emitRangedBuiltin(*this, Intrinsic::riscv_workitem_id_x, 0, 1024);
case RISCV::BI__builtin_riscv_workitem_id_y:
return emitRangedBuiltin(*this, Intrinsic::riscv_workitem_id_y, 0, 1024);
case RISCV::BI__builtin_riscv_workitem_id_z:
return emitRangedBuiltin(*this, Intrinsic::riscv_workitem_id_z, 0, 1024);
// FIXME: Use AMDGPU function here temporarily
// Ventus GPGPU workgroup size
case RISCV::BI__builtin_riscv_workgroup_size_x:
return EmitAMDGPUWorkGroupSize(*this, 0);
case RISCV::BI__builtin_riscv_workgroup_size_y:
return EmitAMDGPUWorkGroupSize(*this, 1);
case RISCV::BI__builtin_riscv_workgroup_size_z:
return EmitAMDGPUWorkGroupSize(*this, 2);
// Ventus GPGPU grid size
case RISCV::BI__builtin_riscv_grid_size_x:
return EmitAMDGPUGridSize(*this, 0);
case RISCV::BI__builtin_riscv_grid_size_y:
return EmitAMDGPUGridSize(*this, 1);
case RISCV::BI__builtin_riscv_grid_size_z:
return EmitAMDGPUGridSize(*this, 2);
*/
// Ventus GPGPU workitem
case RISCV::BIbarrier:
case RISCV::BIwork_group_barrier: {
@ -19506,10 +19478,6 @@ Value *CodeGenFunction::EmitRISCVBuiltinExpr(unsigned BuiltinID,
break;
}
// case RISCV::BIwork_group_barrier_scope:
// ID = Intrinsic::riscv_ventus_barrier_with_scope;
// break;
case RISCV::BI__builtin_riscv_orc_b_32:
case RISCV::BI__builtin_riscv_orc_b_64:
case RISCV::BI__builtin_riscv_clz_32: