aparler/c-llm
1
0
Fork 0
mirror of https://github.com/karpathy/llm.c.git synced 2026-07-28 20:35:09 -04:00
Code Issues Projects Releases Packages Wiki Activity Actions

Merge branch 'ngc92-check-determinism'

Browse source
This commit is contained in:
Andrej Karpathy 2024-06-18 03:03:06 +00:00
commit 4fd2df8870
5 changed files with 102 additions and 29 deletions
Download patch file Download diff file Expand all files Collapse all files

14
llmc/cuda_common.h
View file

@ -42,7 +42,7 @@ extern cudaDeviceProp deviceProp;
// Error checking
// CUDA error checking
void inline cudaCheck(cudaError_t error, const char *file, int line) {
inline void cudaCheck(cudaError_t error, const char *file, int line) {
if (error != cudaSuccess) {
printf("[CUDA ERROR] at file %s:%d:\n%s\n", file, line, cudaGetErrorString(error));
exit(EXIT_FAILURE);
@ -50,6 +50,18 @@ void inline cudaCheck(cudaError_t error, const char *file, int line) {
};
#define cudaCheck(err) (cudaCheck(err, __FILE__, __LINE__))
// like cudaFree, but checks for errors _and_ resets the pointer.
template<class T>
inline void cudaFreeCheck(T** ptr, const char *file, int line) {
cudaError_t error = cudaFree(*ptr);
if (error != cudaSuccess) {
printf("[CUDA ERROR] at file %s:%d:\n%s\n", file, line, cudaGetErrorString(error));
exit(EXIT_FAILURE);
}
*ptr = nullptr;
}
#define cudaFreeCheck(ptr) (cudaFreeCheck(ptr, __FILE__, __LINE__))
// ----------------------------------------------------------------------------
// CUDA Precision settings and defines

25
llmc/dataloader.h
View file

@ -96,13 +96,6 @@ int64_t dataloader_load_shard_(DataLoader *loader, int shard_index) {
return ntok;
}
void dataloader_resume(DataLoader *loader, size_t current_shard_idx, size_t current_sample_idx) {
// used during model resumption (-y 1) flag
loader->current_shard_idx = current_shard_idx;
loader->current_sample_idx = current_sample_idx;
dataloader_load_shard_(loader, loader->current_shard_idx);
}
void prepare_intra_shard_indices_(DataLoader *loader) {
// shuffle the examples inside the shards
if (loader->intra_shard_indices != NULL) {
@ -206,12 +199,14 @@ void dataloader_init(DataLoader *loader,
// reset the loader, to initialize it
dataloader_reset(loader);
// we haven't drawn a current sample yet
// note that this line will underflow but that's ok, as +1 later will just overflow to 0
loader->current_sample_idx = (size_t) -1;
}
void dataloader_next_batch(DataLoader *loader) {
void dataloader_load_batch(DataLoader* loader) {
assert(!loader->should_shuffle || (loader->should_shuffle && loader->intra_shard_indices != NULL));
assert(loader->current_sample_idx < loader->shard_num_samples);
size_t idx = loader->should_shuffle ? loader->intra_shard_indices[loader->current_sample_idx] : loader->current_sample_idx;
size_t global_batch_offset_bytes = idx * loader->total_batch_size_bytes;
int64_t current_offset = loader->header_bytes + global_batch_offset_bytes + loader->local_batch_offset_bytes;
@ -226,12 +221,24 @@ void dataloader_next_batch(DataLoader *loader) {
loader->inputs[i] = (int)loader->buffer[i];
loader->targets[i] = (int)loader->buffer[i+1];
}
}
void dataloader_next_batch(DataLoader *loader) {
loader->current_sample_idx += 1;
// if the next batch would go past the end of the file, advance the loader
if (loader->current_sample_idx >= loader->shard_num_samples) {
dataloader_advance_(loader);
}
dataloader_load_batch(loader);
}
void dataloader_resume(DataLoader *loader, size_t current_shard_idx, size_t current_sample_idx) {
// used during model resumption (-y 1) flag
loader->current_shard_idx = current_shard_idx;
loader->current_sample_idx = current_sample_idx;
dataloader_load_shard_(loader, loader->current_shard_idx);
dataloader_load_batch(loader);
}
void dataloader_free(DataLoader *loader) {

1
profile_gpt2.cu
View file

@ -33,6 +33,7 @@ int main(int argc, char *argv[]) {
// build the GPT-2 model from a checkpoint
GPT2 model;
gpt2_init_common(&model);
gpt2_build_from_checkpoint(&model, "gpt2_124M_bf16.bin");
int B = 24; // if program OOMs decrease this number, e.g. all the way down to 4 or etc

53
test_gpt2.cu
View file

@ -101,6 +101,7 @@ int main(int argc, char *argv[]) {
// build the GPT-2 model from a checkpoint
GPT2 model;
gpt2_init_common(&model);
gpt2_build_from_checkpoint(&model, load_filename);
size_t V = model.config.vocab_size;
@ -226,7 +227,7 @@ int main(int argc, char *argv[]) {
}
// move the (mixed precision) grads from GPU to CPU
cudaMemcpy(grads_memory_cpu, model.grads_memory, model.num_parameters_bytes, cudaMemcpyDeviceToHost);
cudaCheck(cudaMemcpy(grads_memory_cpu, model.grads_memory, model.num_parameters_bytes, cudaMemcpyDeviceToHost));
// convert all gradients to float on the CPU
char* src_iterator = (char*)grads_memory_cpu; // can be lower precision, so we use char*
@ -312,10 +313,60 @@ int main(int argc, char *argv[]) {
}
}
// Finally, let's check determinism
gpt2_write_to_checkpoint(&model, "test_gpt2cu_model.ckpt");
DataLoader loader;
dataloader_init(&loader, "dev/data/tinyshakespeare/tiny_shakespeare_val.bin", B, T, multi_gpu_config.process_rank, multi_gpu_config.num_processes, 1);
dataloader_next_batch(&loader);
save_state("test_gpt2cu_state.ckpt", 10, &model, &loader);
int tokens[10];
for (int step = 0; step < 10; step++) {
gpt2_forward(&model, loader.inputs, loader.targets, B, T);
gpt2_zero_grad(&model);
gpt2_backward(&model, loader.inputs, true);
gpt2_update(&model, 1e-4f, 0.9f, 0.95f, 1e-8f, 0.0f, 1.0f, step+11, &multi_gpu_config);
losses[step] = model.mean_loss;
tokens[step] = loader.inputs[0];
dataloader_next_batch(&loader);
}
// reload
gpt2_free(&model);
gpt2_build_from_checkpoint(&model, "test_gpt2cu_model.ckpt");
int ld_step;
load_state(&ld_step, &model, &loader, "test_gpt2cu_state.ckpt");
for (int step = 0; step < 10; step++) {
gpt2_forward(&model, loader.inputs, loader.targets, B, T);
gpt2_zero_grad(&model);
gpt2_backward(&model, loader.inputs, true);
gpt2_update(&model, 1e-4f, 0.9f, 0.95f, 1e-8f, 0.0f, 1.0f, step+11, &multi_gpu_config);
if(loader.inputs[0] != tokens[step]) {
printf("Nondeterminism! Token mismatch at step %d: %d vs %d\n", step, tokens[step], loader.inputs[0]);
allok = false;
break;
}
if(losses[step] != model.mean_loss) {
printf("Nondeterminism! Loss mismatch at step %d: %.15f vs %.15f\n", step, losses[step], model.mean_loss);
allok = false;
break;
} else {
printf("loss ok at step %d: %f %f\n", step, losses[step], model.mean_loss);
}
dataloader_next_batch(&loader);
}
// final approval
printf("overall okay: %d\n", allok);
// delete intermediate test files
remove("test_gpt2cu_model.ckpt");
remove("test_gpt2cu_state.ckpt");
// free everything
dataloader_free(&loader);
gpt2_free(&model);
common_free(model);
free(x);

38
train_gpt2.cu
View file

@ -350,6 +350,7 @@ void gpt2_init_common(GPT2 *model) {
model->batch_size = 0;
model->seq_len = 0;
model->mean_loss = -1.0f; // -1.0f designates no loss, set at end of forward()
model->params_memory = NULL;
// memory lazily initialized in backward()
model->grads_memory = NULL;
model->workload_indices = NULL; // on cpu, for encoder_backward
@ -444,6 +445,7 @@ void gpt2_build_from_checkpoint(GPT2 *model, const char* checkpoint_path) {
}
// create memory for model parameters on the device
assert(model->params_memory == nullptr && "Old model needs to be freed before loading from checkpoint again");
model->params_memory = malloc_and_point_parameters(&model->params, model->param_elements, model->param_sizeof);
// read in all the parameters from file and copy them to device
@ -453,7 +455,6 @@ void gpt2_build_from_checkpoint(GPT2 *model, const char* checkpoint_path) {
free(params_memory_cpu);
fcloseCheck(model_file);
gpt2_init_common(model);
// only return from this function once we are certain the params are ready on the GPU
cudaCheck(cudaDeviceSynchronize());
}
@ -543,8 +544,6 @@ void gpt2_build_from_random(GPT2 *model, int depth) {
// copy them to GPU
cudaCheck(cudaMemcpy(model->params_memory, params_memory_cpu, model->num_parameters_bytes, cudaMemcpyHostToDevice));
free(params_memory_cpu);
gpt2_init_common(model);
}
void gpt2_forward(GPT2 *model, const int* inputs, const int* targets, size_t B, size_t T, int grad_accum_steps=1) {
@ -1105,14 +1104,14 @@ float gpt2_estimate_mfu(GPT2 *model, int num_tokens, float dt) {
}
void gpt2_free(GPT2 *model) {
cudaCheck(cudaFree(model->params_memory));
cudaCheck(cudaFree(model->grads_memory));
cudaCheck(cudaFree(model->m_memory));
cudaCheck(cudaFree(model->v_memory));
cudaCheck(cudaFree(model->master_weights));
cudaCheck(cudaFree(model->acts_memory));
cudaCheck(cudaFree(model->inputs));
cudaCheck(cudaFree(model->targets));
cudaFreeCheck(&model->params_memory);
cudaFreeCheck(&model->grads_memory);
cudaFreeCheck(&model->m_memory);
cudaFreeCheck(&model->v_memory);
cudaFreeCheck(&model->master_weights);
cudaFreeCheck(&model->acts_memory);
cudaFreeCheck(&model->inputs);
cudaFreeCheck(&model->targets);
cudaCheck(cudaFreeHost(model->cpu_losses));
cudaCheck(cudaFreeHost(model->cpu_losses_fp32));
free(model->workload_indices);
@ -1160,13 +1159,6 @@ void common_free(GPT2 &model) {
#endif
}
#ifndef TESTING
// if we are TESTING (see test_gpt2.cu), we'll skip everything below this point
// ----------------------------------------------------------------------------
// training resumption logic, very useful when jobs crash once in a while
// the goal is that we can resume optimization from any checkpoint, bit-perfect
// note that "state" refers to things not already saved in the model checkpoint file
void save_state(const char* filename, int step, GPT2* model, DataLoader* loader) {
printf("Writing state to %s\n", filename);
@ -1280,6 +1272,15 @@ void load_state(int* step, GPT2* model, DataLoader* loader, const char* filename
fclose(state_file);
}
#ifndef TESTING
// if we are TESTING (see test_gpt2.cu), we'll skip everything below this point
// ----------------------------------------------------------------------------
// training resumption logic, very useful when jobs crash once in a while
// the goal is that we can resume optimization from any checkpoint, bit-perfect
// note that "state" refers to things not already saved in the model checkpoint file
// ----------------------------------------------------------------------------
// CLI, poor man's argparse
// unclaimed flags lol: k,p
@ -1453,6 +1454,7 @@ int main(int argc, char *argv[]) {
// build the GPT-2 model
GPT2 model;
gpt2_init_common(&model);
// if load_filename is of the form "dX" where X is an integer (e.g. d12), then we build
// a random model with the depth of the model specified by X (e.g. 12). otherwise interpret
// this variable as a checkpoint filename, and load that checkpoint