Skip to content

Rename FusionExector to KernelExecutor, fe to ke, fec to executor_cache #3349

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 7 commits into from
Nov 7, 2024
Merged

Rename FusionExector to KernelExecutor, fe to ke, fec to executor_cache #3349

Show file tree
Hide file tree
Changes from 3 commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
b115851
Rename FusionExecutor to KernelExecutor
naoyam Nov 5, 2024
58207c5
rename fe to ke
naoyam Nov 5, 2024
6019626
Rename fec to executor_cache
naoyam Nov 5, 2024
f064e07
forgot to rename
naoyam Nov 5, 2024
629b774
rename fusion_executor_cache to executor_cache
naoyam Nov 5, 2024
14993fa
rename executor to ke
naoyam Nov 5, 2024
ba2a03b
Merge branch 'main' into rename
naoyam Nov 7, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
8 changes: 4 additions & 4 deletions benchmarks/cpp/gelu_backward.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -162,7 +162,7 @@ static void NvFuserScheduler_GeluBackward_Compile(
&fusion, SchedulerType::PointWise, c10::ArrayRef<c10::IValue>(inputs));

for (auto _ : benchmark_state) {
FusionExecutor executor;
KernelExecutor executor;
executor.compileFusion(&fusion, inputs, heuristic_params->lparams);
}
}
Expand All @@ -187,7 +187,7 @@ static void NvFuserScheduler_GeluBackward_RunFusion(
auto heuristic_params = SchedulerEntry::scheduleWith(
&fusion, SchedulerType::PointWise, c10::ArrayRef<c10::IValue>(inputs));

FusionExecutor executor;
KernelExecutor executor;
executor.compileFusion(&fusion, inputs, heuristic_params->lparams);

C10_CUDA_CHECK(cudaDeviceSynchronize());
Expand Down Expand Up @@ -218,7 +218,7 @@ static void NvFuserScheduler_GeluBackward_RunFusion_GpuOnly(
auto heuristic_params = SchedulerEntry::scheduleWith(
&fusion, SchedulerType::PointWise, c10::ArrayRef<c10::IValue>(inputs));

FusionExecutor executor;
KernelExecutor executor;
executor.compileFusion(&fusion, inputs, heuristic_params->lparams);

runBenchmarkIterations(
Expand Down Expand Up @@ -247,7 +247,7 @@ static void NvFuserScheduler_GeluBackward_RunFusion_CpuOnly(
auto heuristic_params = SchedulerEntry::scheduleWith(
&fusion, SchedulerType::PointWise, c10::ArrayRef<c10::IValue>(inputs));

FusionExecutor executor;
KernelExecutor executor;
executor.setExecuteKernelFlag(false);
executor.compileFusion(&fusion, inputs, heuristic_params->lparams);

Expand Down
24 changes: 12 additions & 12 deletions benchmarks/cpp/heuristic_cache.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -26,7 +26,7 @@ using namespace nvfuser;

static auto getLayerBackwardNormRuntime(
std::unique_ptr<Fusion> fusion_ptr,
std::unique_ptr<FusionExecutorCache>& fec,
std::unique_ptr<FusionExecutorCache>& executor_cache,
std::vector<c10::IValue>& aten_inputs,
std::vector<int64_t>& shape,
std::vector<int64_t>& norm_shape) {
Expand Down Expand Up @@ -84,12 +84,12 @@ static auto getLayerBackwardNormRuntime(
auto aten_mean = std::get<1>(aten_results);
auto aten_rstd = std::get<2>(aten_results);

fec = std::make_unique<FusionExecutorCache>(std::move(fusion_ptr));
executor_cache = std::make_unique<FusionExecutorCache>(std::move(fusion_ptr));
aten_inputs = {
aten_grad_out, aten_input, aten_mean, aten_rstd, aten_weight, aten_bias};
auto cg_outputs = fec->runFusionWithInputs(aten_inputs);
auto cg_outputs = executor_cache->runFusionWithInputs(aten_inputs);

return fec->getMostRecentKernelRuntime();
return executor_cache->getMostRecentKernelRuntime();
}

static void NvFuserScheduler_LayerNormBackward_HeuristicCache(
Expand All @@ -98,14 +98,14 @@ static void NvFuserScheduler_LayerNormBackward_HeuristicCache(
FusionGuard fg(fusion_ptr.get());

// PreAllocate
std::unique_ptr<FusionExecutorCache> fec;
std::unique_ptr<FusionExecutorCache> executor_cache;
std::vector<c10::IValue> aten_inputs;

std::vector<int64_t> shape{20, 100, 35, 67};
std::vector<int64_t> norm_shape{67};

auto runtime = getLayerBackwardNormRuntime(
std::move(fusion_ptr), fec, aten_inputs, shape, norm_shape);
std::move(fusion_ptr), executor_cache, aten_inputs, shape, norm_shape);

KernelArgumentHolder args =
KernelArgumentHolder::createKernelArgumentHolder(aten_inputs);
Expand All @@ -120,7 +120,7 @@ static void NvFuserScheduler_LayerNormBackward_HeuristicCache(

static auto getLayerForwardNormRuntime(
std::unique_ptr<Fusion> fusion_ptr,
std::unique_ptr<FusionExecutorCache>& fec,
std::unique_ptr<FusionExecutorCache>& executor_cache,
std::vector<c10::IValue>& aten_inputs,
std::vector<int64_t>& shape,
std::vector<int64_t>& norm_shape) {
Expand All @@ -141,11 +141,11 @@ static auto getLayerForwardNormRuntime(
auto options = at::TensorOptions().dtype(at::kFloat).device(at::kCUDA, 0);
at::Tensor aten_input = at::randn(shape, options);

fec = std::make_unique<FusionExecutorCache>(std::move(fusion_ptr));
executor_cache = std::make_unique<FusionExecutorCache>(std::move(fusion_ptr));
aten_inputs = {aten_input};
auto cg_outputs = fec->runFusionWithInputs(aten_inputs);
auto cg_outputs = executor_cache->runFusionWithInputs(aten_inputs);

return fec->getMostRecentKernelRuntime();
return executor_cache->getMostRecentKernelRuntime();
}

static void NvFuserScheduler_LayerNormForward_HeuristicCache(
Expand All @@ -154,14 +154,14 @@ static void NvFuserScheduler_LayerNormForward_HeuristicCache(
FusionGuard fg(fusion_ptr.get());

// PreAllocate
std::unique_ptr<FusionExecutorCache> fec;
std::unique_ptr<FusionExecutorCache> executor_cache;
std::vector<c10::IValue> aten_inputs;

std::vector<int64_t> shape{20, 100, 35, 67};
std::vector<int64_t> norm_shape{67};

auto runtime = getLayerForwardNormRuntime(
std::move(fusion_ptr), fec, aten_inputs, shape, norm_shape);
std::move(fusion_ptr), executor_cache, aten_inputs, shape, norm_shape);

KernelArgumentHolder args =
KernelArgumentHolder::createKernelArgumentHolder(aten_inputs);
Expand Down
24 changes: 12 additions & 12 deletions benchmarks/cpp/heuristic_lookup.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -26,7 +26,7 @@ using namespace nvfuser;

static auto getLayerBackwardNormRuntime(
std::unique_ptr<Fusion> fusion_ptr,
std::unique_ptr<FusionExecutorCache>& fec,
std::unique_ptr<FusionExecutorCache>& executor_cache,
std::vector<c10::IValue>& aten_inputs,
std::vector<int64_t>& shape,
std::vector<int64_t>& norm_shape) {
Expand Down Expand Up @@ -86,12 +86,12 @@ static auto getLayerBackwardNormRuntime(
auto aten_mean = std::get<1>(aten_results);
auto aten_rstd = std::get<2>(aten_results);

fec = std::make_unique<FusionExecutorCache>(std::move(fusion_ptr));
executor_cache = std::make_unique<FusionExecutorCache>(std::move(fusion_ptr));
aten_inputs = {
aten_grad_out, aten_input, aten_mean, aten_rstd, aten_weight, aten_bias};
auto cg_outputs = fec->runFusionWithInputs(aten_inputs);
auto cg_outputs = executor_cache->runFusionWithInputs(aten_inputs);

return fec->getMostRecentKernelRuntime();
return executor_cache->getMostRecentKernelRuntime();
}

static void NvFuserScheduler_LayerNormBackward_HeuristicLookup(
Expand All @@ -100,14 +100,14 @@ static void NvFuserScheduler_LayerNormBackward_HeuristicLookup(
FusionGuard fg(fusion_ptr.get());

// PreAllocate
std::unique_ptr<FusionExecutorCache> fec;
std::unique_ptr<FusionExecutorCache> executor_cache;
std::vector<c10::IValue> aten_inputs;

std::vector<int64_t> shape{20, 100, 35, 67};
std::vector<int64_t> norm_shape{67};

auto runtime = getLayerBackwardNormRuntime(
std::move(fusion_ptr), fec, aten_inputs, shape, norm_shape);
std::move(fusion_ptr), executor_cache, aten_inputs, shape, norm_shape);

KernelArgumentHolder args =
KernelArgumentHolder::createKernelArgumentHolder(aten_inputs);
Expand All @@ -122,7 +122,7 @@ static void NvFuserScheduler_LayerNormBackward_HeuristicLookup(

static auto getLayerForwardNormRuntime(
std::unique_ptr<Fusion> fusion_ptr,
std::unique_ptr<FusionExecutorCache>& fec,
std::unique_ptr<FusionExecutorCache>& executor_cache,
std::vector<c10::IValue>& aten_inputs,
std::vector<int64_t>& shape,
std::vector<int64_t>& norm_shape) {
Expand All @@ -143,11 +143,11 @@ static auto getLayerForwardNormRuntime(
auto options = at::TensorOptions().dtype(at::kFloat).device(at::kCUDA, 0);
at::Tensor aten_input = at::randn(shape, options);

fec = std::make_unique<FusionExecutorCache>(std::move(fusion_ptr));
executor_cache = std::make_unique<FusionExecutorCache>(std::move(fusion_ptr));
aten_inputs = {aten_input};
auto cg_outputs = fec->runFusionWithInputs(aten_inputs);
auto cg_outputs = executor_cache->runFusionWithInputs(aten_inputs);

return fec->getMostRecentKernelRuntime();
return executor_cache->getMostRecentKernelRuntime();
}

static void NvFuserScheduler_LayerNormForward_HeuristicLookup(
Expand All @@ -156,14 +156,14 @@ static void NvFuserScheduler_LayerNormForward_HeuristicLookup(
FusionGuard fg(fusion_ptr.get());

// PreAllocate
std::unique_ptr<FusionExecutorCache> fec;
std::unique_ptr<FusionExecutorCache> executor_cache;
std::vector<c10::IValue> aten_inputs;

std::vector<int64_t> shape{20, 100, 35, 67};
std::vector<int64_t> norm_shape{67};

auto runtime = getLayerForwardNormRuntime(
std::move(fusion_ptr), fec, aten_inputs, shape, norm_shape);
std::move(fusion_ptr), executor_cache, aten_inputs, shape, norm_shape);

KernelArgumentHolder args =
KernelArgumentHolder::createKernelArgumentHolder(aten_inputs);
Expand Down
4 changes: 2 additions & 2 deletions benchmarks/cpp/indexselect.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -132,7 +132,7 @@ static void NvFuserScheduler_IndexSelect_Compile(
&fusion, SchedulerType::PointWise, c10::ArrayRef<c10::IValue>(inputs));

for (auto _ : benchmark_state) {
FusionExecutor executor;
KernelExecutor executor;
executor.compileFusion(
&fusion, c10::ArrayRef<c10::IValue>(inputs), heuristic_params->lparams);
}
Expand All @@ -155,7 +155,7 @@ static void NvFuserScheduler_IndexSelect_RunFusion(
auto heuristic_params = SchedulerEntry::scheduleWith(
&fusion, SchedulerType::PointWise, c10::ArrayRef<c10::IValue>(inputs));

FusionExecutor executor;
KernelExecutor executor;
executor.compileFusion(
&fusion, c10::ArrayRef<c10::IValue>(inputs), heuristic_params->lparams);

Expand Down
8 changes: 4 additions & 4 deletions benchmarks/cpp/lstm_cell.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -155,7 +155,7 @@ static void NvFuserScheduler_LstmCell_Compile(
&fusion, SchedulerType::PointWise, c10::ArrayRef<c10::IValue>(inputs));

for (auto _ : benchmark_state) {
FusionExecutor executor;
KernelExecutor executor;
executor.compileFusion(&fusion, inputs);
}
}
Expand All @@ -182,7 +182,7 @@ static void NvFuserScheduler_LstmCell_RunFusion(
auto heuristic_params = SchedulerEntry::scheduleWith(
&fusion, SchedulerType::PointWise, c10::ArrayRef<c10::IValue>(inputs));

FusionExecutor executor;
KernelExecutor executor;
executor.compileFusion(&fusion, inputs);

C10_CUDA_CHECK(cudaDeviceSynchronize());
Expand Down Expand Up @@ -220,7 +220,7 @@ static void NvFuserScheduler_LstmCell_RunFusion_GpuOnly(
auto heuristic_params = SchedulerEntry::scheduleWith(
&fusion, SchedulerType::PointWise, c10::ArrayRef<c10::IValue>(inputs));

FusionExecutor executor;
KernelExecutor executor;
executor.compileFusion(&fusion, inputs);

runBenchmarkIterations(
Expand Down Expand Up @@ -259,7 +259,7 @@ static void NvFuserScheduler_LstmCell_RunFusion_CpuOnly(
auto heuristic_params = SchedulerEntry::scheduleWith(
&fusion, SchedulerType::PointWise, c10::ArrayRef<c10::IValue>(inputs));

FusionExecutor executor;
KernelExecutor executor;
executor.setExecuteKernelFlag(false);
executor.compileFusion(&fusion, inputs);

Expand Down
30 changes: 15 additions & 15 deletions benchmarks/cpp/matmul.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -175,19 +175,19 @@ static void SingleMatmulBase(

// Compile kernel
auto launch_constraints = LaunchParams();
FusionExecutor fe;
fe.compileFusion(fusion, args, launch_constraints, cparams);
KernelExecutor ke;
ke.compileFusion(fusion, args, launch_constraints, cparams);
NVF_CHECK(
getBankConflictInfo(fe.kernel(), launch_constraints).empty(),
getBankConflictInfo(ke.kernel(), launch_constraints).empty(),
"Shared memory bank conflict not removed.");

std::vector<c10::IValue> aten_inputs({inputs.first, inputs.second});

// Warm up run
auto outputs = fe.runFusion(aten_inputs);
auto outputs = ke.runFusion(aten_inputs);
checkMatch(expected_output, outputs.at(0).to(at::kDouble), k);

runBenchmarkIterations(benchmark_state, &fe, aten_inputs);
runBenchmarkIterations(benchmark_state, &ke, aten_inputs);

// TODO: FLOPS calculation
}
Expand Down Expand Up @@ -355,19 +355,19 @@ static void SingleMatmulPartitionedK(
cparams.index_type = computeIndexType(M, N, K);

// Compile kernel
FusionExecutor fe;
KernelExecutor ke;
auto lparams = LaunchParams();
fe.compileFusion(fusion, args, lparams, cparams);
ke.compileFusion(fusion, args, lparams, cparams);
NVF_CHECK(
getBankConflictInfo(fe.kernel(), lparams).empty(),
getBankConflictInfo(ke.kernel(), lparams).empty(),
"Shared memory bank conflict not removed.");

// Warm up run
auto outputs = fe.runFusion(aten_inputs);
auto outputs = ke.runFusion(aten_inputs);

checkMatch(expected_output, outputs.at(0).to(at::kDouble), Ki);

runBenchmarkIterations(benchmark_state, &fe, aten_inputs);
runBenchmarkIterations(benchmark_state, &ke, aten_inputs);

// TODO: FLOPS calculation
}
Expand Down Expand Up @@ -461,21 +461,21 @@ static void NvFuserScheduler_MatmulSplitKReduction(
KernelArgumentHolder::createKernelArgumentHolder(aten_inputs);

// Compile kernel
FusionExecutor fe;
fe.compileFusion(
KernelExecutor ke;
ke.compileFusion(
fusion, args, heuristic_params->lparams, heuristic_params->cparams);

NVF_CHECK(
getBankConflictInfo(fe.kernel(), heuristic_params->lparams).empty(),
getBankConflictInfo(ke.kernel(), heuristic_params->lparams).empty(),
"Shared memory bank conflict not removed.");

// Warm up run
auto outputs = fe.runFusion(aten_inputs, heuristic_params->lparams);
auto outputs = ke.runFusion(aten_inputs, heuristic_params->lparams);

checkMatch(expected_output, outputs.at(0).to(at::kDouble), splitk_factor);

runBenchmarkIterations(
benchmark_state, &fe, aten_inputs, heuristic_params->lparams);
benchmark_state, &ke, aten_inputs, heuristic_params->lparams);

// TODO: FLOPS calculation
}
Expand Down
Loading
Loading