#include #include #include "src/turbomind/kernels/gpt_kernels.h" #include "src/turbomind/utils/memory_utils.h" #include "unittest_utils.h" int test_find_context_dups(); int test_compact(); int test_uncompact(); int main(int argc, char* argv[]) { bool all_passed = true; bool passed; passed = test_find_context_dups() == EXIT_SUCCESS; all_passed |= passed; printf("%s", passed ? "." : "X"); if (!passed) { puts("\ntest_find_context_dups: FAILED"); } passed = test_compact() == EXIT_SUCCESS; all_passed |= passed; printf("%s", passed ? "." : "X"); if (!passed) { puts("\ntest_compact: FAILED"); } passed = test_uncompact() == EXIT_SUCCESS; all_passed |= passed; printf("%s", passed ? "." : "X"); if (!passed) { puts("\ntest_uncompact: FAILED"); } puts(""); return all_passed ? EXIT_SUCCESS : EXIT_FAILURE; } int test_find_context_dups() { const size_t vec_size = 1234; const size_t batch_size = 8; // Reference to the first unique vector const std::vector shared_contexts_ref {0, 0, 2, 3, 4, 4, 3, 3}; // Which compact index belong to what vector const std::vector batch_idx_to_compact_idx {0, 0, 1, 2, 3, 3, 2, 2}; std::vector batch_idx_to_compact_idx_test(batch_size); // Reverse map of batch_idx_to_compact_idx const std::vector compact_idx_to_batch_idx {0, 2, 3, 4, -1, -1, -1, -1}; std::vector compact_idx_to_batch_idx_test(batch_size, -1); std::vector input_ids; std::vector default_vector(vec_size, 0); for (size_t i = 0; i < batch_size; ++i) { default_vector[vec_size - 1] = shared_contexts_ref[i]; input_ids.insert(input_ids.end(), default_vector.begin(), default_vector.end()); } std::vector shared_contexts_test(batch_size); int* d_input_ids; int* d_shared_contexts_test; int* d_batch_idx_to_compact_idx; int* d_compact_to_batch; int* d_compact_size; cudaMalloc(&d_input_ids, batch_size * vec_size * sizeof(int)); cudaMalloc(&d_shared_contexts_test, batch_size * sizeof(int)); cudaMalloc(&d_batch_idx_to_compact_idx, batch_size * sizeof(int)); cudaMalloc(&d_compact_to_batch, batch_size * sizeof(int)); cudaMalloc(&d_compact_size, sizeof(int)); cudaH2Dcpy(d_input_ids, input_ids.data(), batch_size * vec_size); cudaH2Dcpy(d_compact_to_batch, compact_idx_to_batch_idx_test.data(), batch_size); invokeFindContextDups(d_shared_contexts_test, d_batch_idx_to_compact_idx, d_compact_to_batch, d_compact_size, d_input_ids, batch_size, vec_size); int compact_size; cudaD2Hcpy(shared_contexts_test.data(), d_shared_contexts_test, batch_size); cudaD2Hcpy(batch_idx_to_compact_idx_test.data(), d_batch_idx_to_compact_idx, batch_size); cudaD2Hcpy(compact_idx_to_batch_idx_test.data(), d_compact_to_batch, batch_size); cudaD2Hcpy(&compact_size, d_compact_size, 1); cudaFree(d_input_ids); cudaFree(d_shared_contexts_test); EXPECT_TRUE(shared_contexts_test == shared_contexts_ref); EXPECT_TRUE(batch_idx_to_compact_idx == batch_idx_to_compact_idx_test); EXPECT_TRUE(compact_idx_to_batch_idx_test == compact_idx_to_batch_idx); EXPECT_TRUE(compact_size == 4); return EXIT_SUCCESS; } int test_compact() { size_t batch_size = 128; size_t compact_size = 5; size_t seq_len = 40; size_t hidden_dimension = 8; auto generator_f = std::bind(std::uniform_real_distribution(-1.0, 1.0), std::mt19937()); auto generator_i = std::bind(std::uniform_int_distribution(0, 128), std::mt19937()); // decoder_input [batch_size, seq_len, hidden_dimension] -> // compact_decoder_input [compact_size, seq_len, hidden_dimension] std::vector decoder_input(batch_size * seq_len * hidden_dimension); std::vector compact_decoder_input(compact_size * seq_len * hidden_dimension); std::generate(decoder_input.begin(), decoder_input.end(), generator_f); float *d_decoder_input, *d_compact_decoder_input; cudaMalloc(&d_decoder_input, decoder_input.size() * sizeof(float)); cudaMalloc(&d_compact_decoder_input, compact_decoder_input.size() * sizeof(float)); cudaH2Dcpy(d_decoder_input, decoder_input.data(), decoder_input.size()); // attention_mask [batch_size, seq_len, seq_len] -> // compact_attention_mask [compact_size, seq_len, seq_len] std::vector attention_mask(batch_size * seq_len * seq_len); std::vector compact_attention_mask(compact_size * seq_len * seq_len); std::generate(attention_mask.begin(), attention_mask.end(), generator_f); float *d_attention_mask, *d_compact_attention_mask; cudaMalloc(&d_attention_mask, attention_mask.size() * sizeof(float)); cudaMalloc(&d_compact_attention_mask, compact_attention_mask.size() * sizeof(float)); cudaH2Dcpy(d_attention_mask, attention_mask.data(), attention_mask.size()); // input_lengths [batch_size] -> compact_input_lengths [compact_size] std::vector input_lengths(batch_size); std::vector compact_input_lengths(compact_size); std::generate(input_lengths.begin(), input_lengths.end(), generator_i); int *d_input_lengths, *d_compact_input_lengths; cudaMalloc(&d_input_lengths, input_lengths.size() * sizeof(int)); cudaMalloc(&d_compact_input_lengths, compact_input_lengths.size() * sizeof(int)); cudaH2Dcpy(d_input_lengths, input_lengths.data(), input_lengths.size()); // compact_idx [compact_size] /* std::vector compact_idx {0, 3}; */ std::vector compact_idx {0, 29, 42, 44, 100}; int *d_compact_idx; cudaMalloc(&d_compact_idx, compact_idx.size() * sizeof(int)); cudaH2Dcpy(d_compact_idx, compact_idx.data(), compact_idx.size()); invokeCompactInputs(d_compact_decoder_input, d_compact_attention_mask, d_compact_input_lengths, d_decoder_input, d_attention_mask, d_input_lengths, d_compact_idx, compact_size, seq_len, hidden_dimension); cudaD2Hcpy(compact_decoder_input.data(), d_compact_decoder_input, compact_decoder_input.size()); cudaD2Hcpy(compact_attention_mask.data(), d_compact_attention_mask, compact_attention_mask.size()); cudaD2Hcpy(compact_input_lengths.data(), d_compact_input_lengths, compact_input_lengths.size()); for (size_t i = 0; i < compact_size; i++) { for (size_t t = 0; t < seq_len; t++) { for (size_t h = 0; h < hidden_dimension; h++) { EXPECT_TRUE(compact_decoder_input[(i * seq_len + t) * hidden_dimension + h] == decoder_input[(compact_idx[i] * seq_len + t) * hidden_dimension + h]); } } } for (size_t i = 0; i < compact_size; i++) { for (size_t t1 = 0; t1 < seq_len; t1++) { for (size_t t2 = 0; t2 < seq_len; t2++) { EXPECT_TRUE(compact_attention_mask[(i * seq_len + t1) * seq_len + t2] == attention_mask[(compact_idx[i] * seq_len + t1) * seq_len + t2]); } } } for (size_t i = 0; i < compact_size; i++) { EXPECT_TRUE(compact_input_lengths[i] == input_lengths[compact_idx[i]]); } cudaFree(d_decoder_input); cudaFree(d_compact_decoder_input); cudaFree(d_attention_mask); cudaFree(d_compact_attention_mask); cudaFree(d_input_lengths); cudaFree(d_compact_input_lengths); cudaFree(d_compact_idx); return EXIT_SUCCESS; } int test_uncompact() { // compact_decoder_outputs [compact_size, seq_len, hidden_dimension] -> // decoder_outputs [batch_size, seq_len, hidden_dimension] size_t batch_size = 128; size_t compact_size = 6; size_t local_batch_size = compact_size / 2; size_t seq_len = 40; size_t max_seq_len = 60; size_t hidden_dimension = 8; size_t num_layer = 2; size_t num_head = 2; size_t size_per_head = 4; auto generator_f = std::bind(std::uniform_real_distribution(-1.0, 1.0), std::mt19937()); auto generator_i = std::bind(std::uniform_int_distribution(0, compact_size - 1), std::mt19937()); std::vector compact_decoder_outputs(compact_size * seq_len * hidden_dimension); std::vector decoder_outputs(batch_size * seq_len * hidden_dimension); std::vector k_cache_compact(num_layer * compact_size * num_head * size_per_head * seq_len); std::vector v_cache_compact(num_layer * compact_size * num_head * seq_len * size_per_head); std::vector k_cache_out(num_layer * batch_size * num_head * size_per_head * max_seq_len); std::vector v_cache_out(num_layer * batch_size * num_head * max_seq_len * size_per_head); std::generate(compact_decoder_outputs.begin(), compact_decoder_outputs.end(), generator_f); std::generate(k_cache_compact.begin(), k_cache_compact.end(), generator_f); std::generate(v_cache_compact.begin(), v_cache_compact.end(), generator_f); std::vector batch_to_compact_idx(batch_size); std::generate(batch_to_compact_idx.begin(), batch_to_compact_idx.end(), generator_i); float *d_compact_decoder_outputs, *d_decoder_outputs, *d_k_cache, *d_v_cache; float *d_k_cache_compact, *d_v_cache_compact; cudaMalloc(&d_compact_decoder_outputs, compact_decoder_outputs.size() * sizeof(float)); cudaH2Dcpy(d_compact_decoder_outputs, compact_decoder_outputs.data(), compact_decoder_outputs.size()); cudaMalloc(&d_k_cache_compact, k_cache_compact.size() * sizeof(float)); cudaMalloc(&d_v_cache_compact, v_cache_compact.size() * sizeof(float)); cudaH2Dcpy(d_k_cache_compact, k_cache_compact.data(), k_cache_compact.size()); cudaH2Dcpy(d_v_cache_compact, v_cache_compact.data(), v_cache_compact.size()); cudaMalloc(&d_k_cache, k_cache_out.size() * sizeof(float)); cudaMalloc(&d_v_cache, v_cache_out.size() * sizeof(float)); cudaMemset(d_k_cache, 0, k_cache_out.size() * sizeof(float)); cudaMemset(d_v_cache, 0, v_cache_out.size() * sizeof(float)); cudaMalloc(&d_decoder_outputs, decoder_outputs.size() * sizeof(float)); int *d_batch_to_compact_idx; cudaMalloc(&d_batch_to_compact_idx, batch_to_compact_idx.size() * sizeof(int)); cudaH2Dcpy(d_batch_to_compact_idx, batch_to_compact_idx.data(), batch_to_compact_idx.size()); const size_t cache_stride_dst = max_seq_len * hidden_dimension; const size_t cache_stride_src = seq_len * hidden_dimension; for (size_t ite = 0; ite < (batch_size / local_batch_size); ite++) { for (size_t l = 0; l < num_layer; l++) { const float *k_cache_offset = d_k_cache_compact + (l * compact_size + ite * local_batch_size) * cache_stride_src; const float *v_cache_offset = d_v_cache_compact + (l * compact_size + ite * local_batch_size) * cache_stride_src; invokeUnCompactCaches(d_k_cache + l * batch_size * cache_stride_dst, d_v_cache + l * batch_size * cache_stride_dst, k_cache_offset, v_cache_offset, d_batch_to_compact_idx, batch_size, num_head, max_seq_len, seq_len, size_per_head, local_batch_size, ite); } } invokeUnCompactOutputs(d_decoder_outputs, d_compact_decoder_outputs, d_batch_to_compact_idx, batch_size, cache_stride_src); cudaD2Hcpy(decoder_outputs.data(), d_decoder_outputs, decoder_outputs.size()); cudaD2Hcpy(k_cache_out.data(), d_k_cache, k_cache_out.size()); cudaD2Hcpy(v_cache_out.data(), d_v_cache, v_cache_out.size()); for (size_t i = 0; i < batch_size; i++) { for (size_t t = 0; t < seq_len; t++) { for (size_t h = 0; h < hidden_dimension; h++) { EXPECT_TRUE(decoder_outputs[(i * seq_len + t) * hidden_dimension] == compact_decoder_outputs[(batch_to_compact_idx[i] * seq_len + t) * hidden_dimension]); } } } size_t x_size = (16 / sizeof(float)); for (size_t l = 0; l < num_layer; l++) { for (size_t i = 0; i < batch_size; i++) { for (size_t h = 0; h < num_head; h++) { for (size_t dh = 0; dh < size_per_head / x_size; dh++) { for (size_t t = 0; t < seq_len; t++) { for (size_t x = 0; x < x_size; x++) { auto src = batch_to_compact_idx[i]; EXPECT_TRUE( k_cache_out[((((l * batch_size + i ) * num_head + h) * (size_per_head / x_size) + dh) * max_seq_len + t) * x_size + x] == k_cache_compact[((((l * compact_size + src) * num_head + h) * (size_per_head / x_size) + dh) * seq_len + t) * x_size + x]); } } } } } } for (size_t l = 0; l < num_layer; l++) { for (size_t i = 0; i < batch_size; i++) { for (size_t h = 0; h < num_head; h++) { for (size_t t = 0; t < seq_len; t++) { for (size_t dh = 0; dh < size_per_head; dh++) { auto src = batch_to_compact_idx[i]; EXPECT_TRUE( v_cache_out[(((l * batch_size + i ) * num_head + h) * max_seq_len + t) * size_per_head + dh] == v_cache_compact[(((l * compact_size + src) * num_head + h) * seq_len + t) * size_per_head + dh]); } } } } } cudaFree(d_compact_decoder_outputs); cudaFree(d_k_cache_compact); cudaFree(d_v_cache_compact); cudaFree(d_k_cache); cudaFree(d_v_cache); cudaFree(d_decoder_outputs); cudaFree(d_batch_to_compact_idx); return EXIT_SUCCESS; }