Skip to content
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.

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

Fix MFU printing #585

Closed
wants to merge 2 commits into from
Closed
Show file tree
Hide file tree
Changes from 1 commit
Commits
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
1 change: 1 addition & 0 deletions llmc/mfu.h
Original file line number Diff line number Diff line change
Expand Up @@ -73,6 +73,7 @@ static GPUEntry gpu_db[] = {
{"NVIDIA GeForce RTX 4070", &ADA, 184, 2475},
{"NVIDIA GeForce RTX 4060 Ti", &ADA, 136, 2535},
{"NVIDIA GeForce RTX 4060", &ADA, 96, 2460},
{"NVIDIA H100 PCIe", &HOPPER, 456, 1620},
{"NVIDIA H100 80GB HBM3", &HOPPER, 528, 1830}, // HBM3 = SXM5
};

Expand Down
19 changes: 11 additions & 8 deletions train_gpt2.cu
Original file line number Diff line number Diff line change
Expand Up @@ -1239,7 +1239,7 @@ void gpt2_multi_gpu_param_gather(GPT2 *model, MultiGpuConfig* multi_gpu_config)
cudaCheck(cudaDeviceSynchronize());
}

float gpt2_estimate_mfu(GPT2 *model, int num_tokens, float dt) {
void gpt2_estimate_mfu(GPT2 *model, int num_tokens, float dt, char* mfu_str) {
/*
Estimate model flops utilization (MFU)
ref: Section 2.1 of https://arxiv.org/pdf/2001.08361
Expand All @@ -1261,11 +1261,12 @@ float gpt2_estimate_mfu(GPT2 *model, int num_tokens, float dt) {
// express our flops throughput as ratio of A100 bfloat16 peak flops
float flops_achieved = (float)flops_per_step * (1.0f / dt); // per second
float flops_promised = get_flops_promised(deviceProp.name, PRECISION_MODE) * 1e12f;
if(flops_promised < 0) {
return -1.f; // don't know
if(flops_promised < 0) { // don't know
snprintf(mfu_str, sizeof(mfu_str), "n/a");
} else {
float mfu = flops_achieved / flops_promised;
snprintf(mfu_str, sizeof(mfu_str), "%.1f%%", 100 * mfu);
}
float mfu = flops_achieved / flops_promised;
return mfu;
}

void gpt2_free(GPT2 *model) {
Expand Down Expand Up @@ -1671,6 +1672,7 @@ int main(int argc, char *argv[]) {
}

// train
char* mfu_str = (char*)malloc(16);
Copy link
Owner

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I'm a bit nervous about introducing this, hardcoding 16 here, having to free it, etc.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I agree there is a slight overhead but this is done once and left. Plus even if we had a leak it's 16 bytes :D

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Lmk if you have a better proposal! I still think in net it's better than having -100% for non-supported GPUs.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

can't we just put these 16 bytes on the stack?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

definitely - that's cleaner!

cudaEvent_t start, end;
cudaCheck(cudaEventCreate(&start));
cudaCheck(cudaEventCreate(&end));
Expand Down Expand Up @@ -1851,10 +1853,10 @@ int main(int argc, char *argv[]) {
bias_corrected_ema_tokens_per_second = ema_tokens_per_second / (1.0f - powf(0.95f, step));
}
float accumulated_loss = multi_gpu_config.num_processes == 1 ? model.mean_loss : model.accumulated_mean_loss;
float mfu = gpt2_estimate_mfu(&model, B * T * grad_accum_steps, time_elapsed_ms / 1000.0f);
printf0("step %4d/%d | train loss %7.6f | norm %6.4f | lr %.2e | %.2f ms | %.1f%% bf16 MFU | %.0f tok/s\n",
gpt2_estimate_mfu(&model, B * T * grad_accum_steps, time_elapsed_ms / 1000.0f, mfu_str);
printf0("step %4d/%d | train loss %7.6f | norm %6.4f | lr %.2e | %.2f ms | %s bf16 MFU | %.0f tok/s\n",
step + 1, train_num_batches, accumulated_loss, grad_norm, step_learning_rate,
time_elapsed_ms, 100*mfu, bias_corrected_ema_tokens_per_second);
time_elapsed_ms, mfu_str, bias_corrected_ema_tokens_per_second);
logger_log_train(&logger, step, model.mean_loss, step_learning_rate, grad_norm);

// disable the profiler after 3 steps of optimization
Expand All @@ -1864,6 +1866,7 @@ int main(int argc, char *argv[]) {
printf0("total average iteration time: %f ms\n", total_sum_iteration_time_s / (train_num_batches-1) * 1000);

// free and destroy everything
free(mfu_str);
cudaCheck(cudaEventDestroy(end));
cudaCheck(cudaEventDestroy(start));
if (run_hellaswag) { evalloader_free(&eval_loader); }
Expand Down