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

refactoring & remove unused functions to reduce LOC (+wip profile.py improvements)

Browse source
This commit is contained in:
ademeure 2024-05-04 23:40:15 +01:00
parent 66b9755ce9
commit bfb9c51446
4 changed files with 325 additions and 549 deletions
Download patch file Download diff file Expand all files Collapse all files

51
profile_gpt2.cu
View file

@ -28,44 +28,7 @@ the profile.ncu-rep from a cloud box to local to pretty view.
#include "train_gpt2.cu"
int main() {
// set up the device
int deviceIdx = 0;
cudaCheck(cudaSetDevice(deviceIdx));
cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp, deviceIdx);
printf("[System]\n");
printf("Device %d: %s\n", deviceIdx, deviceProp.name);
cuda_num_SMs = deviceProp.multiProcessorCount;
cuda_threads_per_SM = deviceProp.maxThreadsPerMultiProcessor;
cuda_arch_major = deviceProp.major;
cuda_arch_minor = deviceProp.minor;
cudaCheck(cudaStreamCreate(&main_stream));
cudaEventCreateWithFlags(&main_event, cudaEventDisableTiming);
cudaEventCreateWithFlags(&loss_event, cudaEventDisableTiming);
for (int i = 0; i < num_parallel_streams; i++) {
cudaCheck(cudaStreamCreate(&parallel_streams[i]));
cudaEventCreateWithFlags(&parallel_events[i], cudaEventDisableTiming);
}
// setup cuBLAS and cuBLASLt
cublasCheck(cublasCreate(&cublas_handle));
cublasCheck(cublasSetStream(cublas_handle, main_stream));
cublasCheck(cublasLtCreate(&cublaslt_handle));
// TF32 precision is equivalent to torch.set_float32_matmul_precision('high')
int enable_tf32 = deviceProp.major >= 8 ? 1 : 0;
printf("enable_tf32: %d\n", enable_tf32);
cublas_compute_type = enable_tf32 ? CUBLAS_COMPUTE_32F_FAST_TF32 : CUBLAS_COMPUTE_32F;
cublasMath_t cublas_math_mode = enable_tf32 ? CUBLAS_TF32_TENSOR_OP_MATH : CUBLAS_DEFAULT_MATH;
cublasCheck(cublasSetMathMode(cublas_handle, cublas_math_mode));
// setup the (global) cuBLASLt workspace
cudaCheck(cudaMalloc(&cublaslt_workspace, cublaslt_workspace_size));
#ifdef ENABLE_CUDNN
checkCudnnErr(cudnnCreate(&cudnn_handle));
#endif
common_start();
// build the GPT-2 model from a checkpoint
GPT2 model;
@ -91,16 +54,8 @@ int main() {
gpt2_backward(&model);
gpt2_update(&model, 1e-4f, 0.9f, 0.999f, 1e-8f, 0.0f, 1);
cudaCheck(cudaDeviceSynchronize()); // finish all CUDA work to get correct precise timings
// free
gpt2_free(&model);
#ifdef ENABLE_CUDNN
if (cudnn_workspace != NULL) { cudaCheck(cudaFree(cudnn_workspace)); }
checkCudnnErr(cudnnDestroy(cudnn_handle));
#endif
cudaCheck(cudaFree(cublaslt_workspace));
cublasCheck(cublasDestroy(cublas_handle));
cublasCheck(cublasLtDestroy(cublaslt_handle));
common_free(model);
return 0;
}

19
profile_gpt2cu.py
View file

@ -31,7 +31,7 @@ metrics = [
"dram__bytes_write.sum", # DRAM writes
"lts__t_sectors_srcunit_tex_op_read.sum", # L2 reads (sectors -- 32B)
"lts__t_sectors_srcunit_tex_op_write.sum", # L2 reads (sectors -- 32B)
"smsp__inst_executed.sum", # instructions
"sm__pipe_tensor_op_hmma_cycles_active.avg.pct_of_peak_sustained_active", # todo - tensor core %
]
cmd = [NCU, "-i", "profile.ncu-rep", "--csv", "--page", "raw", "--metrics", ",".join(metrics)]
result = subprocess.check_output(cmd, text=True).strip()
@ -55,11 +55,11 @@ print("Kernel calls:")
for rid, row in enumerate(reader):
if rid == 0:
# headings
print(f"id pass {'name':<40} {'time':>8} {'RAM rd':>8} {'RAM wt':>8} {'L2 rd':>8} {'L2 wt':>8} {'inst':>8}")
print(f"id pass {'name':<70} {'time':>8} {'RAM BW':>8} {'RAM rd':>8} {'RAM wt':>8} {'L2 rd':>8} {'L2 wt':>8} {'inst':>8}")
continue
if rid == 1:
# units
units = f" {'':<40} {'ms':>8} {'GiB':>8} {'GiB':>8} {'GiB':>8} {'GiB':>8} {'MInst':>8}"
units = f" {'':<70} {'ms':>8} {'GB/s':>8} {'GiB':>8} {'GiB':>8} {'GiB':>8} {'GiB':>8} {'MInst':>8}"
print(units)
print("." * len(units))
continue
@ -74,7 +74,7 @@ for rid, row in enumerate(reader):
write = float(row[12])
l2_read = float(row[14])
l2_write = float(row[15])
inst = float(row[16]) / 1e6
inst = float(row[16])
kid = rid - 2
@ -118,18 +118,21 @@ for rid, row in enumerate(reader):
total['l2_write'] += l2_write
total['inst'] += inst
print(f"{kid:02} {pass_name:4} {fn_name:<40} {time:8.2f} {read:8.2f} {write:8.2f} {l2_read:8.2f} {l2_write:8.2f} {inst:8.2f}")
dram_bw = (read + write) / (time / 1000.0);
print(f"{kid:02} {pass_name:4} {fn_name:<70} {time:8.2f} {dram_bw:8.1f} {read:8.2f} {write:8.2f} {l2_read:8.2f} {l2_write:8.2f} {inst:8.2f}")
total_time = total['time']
total_dram_bw = (total['read'] + total['write']) / (total_time / 1000.0);
print("." * len(units))
print(f" {'Total':<40} {total['time']:8.2f} {total['read']:8.2f} {total['write']:8.2f} {total['l2_read']:8.2f} {total['l2_write']:8.2f} {total['inst']:8.2f}")
print(f" {'Total':<70} {total['time']:8.2f} {total_dram_bw:8.1f} {total['read']:8.2f} {total['write']:8.2f} {total['l2_read']:8.2f} {total['l2_write']:8.2f} {total['inst']:8.2f}")
print()
print("Kernel type summaries:")
print(f" {'name':<40} {'time':>6} {'frac':>6}")
print(f" {'name':<70} {'time':>6} {'frac':>6}")
ordered = sorted(summaries.items(), key=lambda x: x[1], reverse=True)
for entry, value in ordered:
print(f" {entry:<40} {value:6.2f} {100*value / total_time:6.2f}%")
print(f" {entry:<70} {value:6.2f} {100*value / total_time:6.2f}%")
ts = total_time / 1000

49
test_gpt2.cu
View file

@ -83,44 +83,7 @@ float* float_cpu_malloc_and_point_parameters(FloatParameterTensors* params, size
}
int main(int argc, char *argv[]) {
// set up the device
int deviceIdx = 0;
cudaCheck(cudaSetDevice(deviceIdx));
cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp, deviceIdx);
printf("[System]\n");
printf("Device %d: %s\n", deviceIdx, deviceProp.name);
cuda_num_SMs = deviceProp.multiProcessorCount;
cuda_threads_per_SM = deviceProp.maxThreadsPerMultiProcessor;
cuda_arch_major = deviceProp.major;
cuda_arch_minor = deviceProp.minor;
cudaCheck(cudaStreamCreate(&main_stream));
cudaEventCreateWithFlags(&main_event, cudaEventDisableTiming);
cudaEventCreateWithFlags(&loss_event, cudaEventDisableTiming);
for (int i = 0; i < num_parallel_streams; i++) {
cudaCheck(cudaStreamCreate(&parallel_streams[i]));
cudaEventCreateWithFlags(&parallel_events[i], cudaEventDisableTiming);
}
// setup cuBLAS and cuBLASLt
cublasCheck(cublasCreate(&cublas_handle));
cublasCheck(cublasSetStream(cublas_handle, main_stream));
cublasCheck(cublasLtCreate(&cublaslt_handle));
// TF32 precision is equivalent to torch.set_float32_matmul_precision('high')
int enable_tf32 = cuda_arch_major >= 8 ? 1 : 0;
enable_tf32 = 0; // NOTE: disable TF32 for testing!!!
printf("enable_tf32: %d\n", enable_tf32);
cublas_compute_type = enable_tf32 ? CUBLAS_COMPUTE_32F_FAST_TF32 : CUBLAS_COMPUTE_32F;
cublasMath_t cublas_math_mode = enable_tf32 ? CUBLAS_TF32_TENSOR_OP_MATH : CUBLAS_DEFAULT_MATH;
cublasCheck(cublasSetMathMode(cublas_handle, cublas_math_mode));
cudaCheck(cudaMalloc(&cublaslt_workspace, cublaslt_workspace_size));
#ifdef ENABLE_CUDNN
checkCudnnErr(cudnnCreate(&cudnn_handle));
#endif
common_start(false);
// build the GPT-2 model from a checkpoint
GPT2 model;
@ -327,6 +290,7 @@ int main(int argc, char *argv[]) {
printf("overall okay: %d\n", allok);
// free everything
common_free(model);
free(x);
free(y);
free(logits_cpu_raw);
@ -336,14 +300,5 @@ int main(int argc, char *argv[]) {
free(expected_grads_memory);
free(grads_memory_cpu);
free(grads_memory_cpu_float);
gpt2_free(&model);
#ifdef ENABLE_CUDNN
if (cudnn_workspace != NULL) { cudaCheck(cudaFree(cudnn_workspace)); }
checkCudnnErr(cudnnDestroy(cudnn_handle));
#endif
cudaCheck(cudaFree(cublaslt_workspace));
cublasCheck(cublasDestroy(cublas_handle));
cublasCheck(cublasLtDestroy(cublaslt_handle));
return 0;
}

755
train_gpt2.cu
View file

File diff suppressed because it is too large Load diff