Wmma::store_matrix_sync( d_C, acc_frag, 16, wmma::mem_row_major ) įor ( auto& a : A ) a = float( rand() ) / RAND_MAX įor ( auto& b : B ) b = float( rand() ) / RAND_MAX ĬudaMalloc( &d_A, 16 * 16 * sizeof( dtype ) ) ĬudaMalloc( &d_B, 16 * 16 * sizeof( dtype ) ) ĬudaMalloc( &d_C, 16 * 16 * sizeof( dtype ) ) ĬudaMemcpy( d_A, A.data(), 16 * 16 * sizeof( dtype ), cudaMemcpyHostToDevice ) ĬudaMemcpy( d_B, B.data(), 16 * 16 * sizeof( dtype ), cudaMemcpyHostToDevice ) ĬudaMemcpy( d_C, C.data(), 16 * 16 * sizeof( dtype ), cudaMemcpyHostToDevice ) Wmma::mma_sync( acc_frag, a_frag, b_frag, acc_frag ) Wmma::load_matrix_sync( a_frag, d_A, 16 ) _global_ void test_wmma( _half* d_A, _half* d_B, _half* d_C ) In the hope that someone here can help me understand what I am doing wrong, I will post a small repro-case here. I am attempting to use the tensor cores efficiently in a custom DL inference kernel, but I get very poor performance.
0 Comments
Leave a Reply. |
AuthorWrite something about yourself. No need to be fancy, just an overview. ArchivesCategories |