From 6f23d8d54ad3edd780295d0feeb2462cb5175558 Mon Sep 17 00:00:00 2001 From: Jules-Kong Date: Thu, 13 Mar 2025 17:32:54 +0800 Subject: [PATCH] [VENTUS][NFC] Removed old implementations of some builtin functions --- clang/include/clang/Basic/BuiltinsRISCV.def | 17 ----------- clang/lib/CodeGen/CGBuiltin.cpp | 32 --------------------- 2 files changed, 49 deletions(-) diff --git a/clang/include/clang/Basic/BuiltinsRISCV.def b/clang/include/clang/Basic/BuiltinsRISCV.def index 35042db78043..d2a9c40ce3b8 100644 --- a/clang/include/clang/Basic/BuiltinsRISCV.def +++ b/clang/include/clang/Basic/BuiltinsRISCV.def @@ -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. //===----------------------------------------------------------------------===// diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index d9595d8e7ca6..7214de6bc967 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -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: