Streaming data ke kernel CUDA tanpa banyak peluncuran

Saya mencoba mempercepat algoritma GPU di mana saya menerima aliran partikel asinkron dalam ruang 3D $p=[x,y,t]$. Setiap vektor $p_n$ perlu dikalikan dengan sekumpulan matriks transformasi. Karena transformasi ini tidak bergantung satu sama lain, transformasi ini dapat terjadi secara paralel, jadi saya telah menulis kernel CUDA untuk melakukan itu. Ini berfungsi dengan baik, tetapi tentu saja untuk setiap $p_n$ yang masuk saya akhirnya meluncurkan kernel CUDA lagi. Meluncurkan kernel CUDA memerlukan penalti waktu yang besar, sehingga saya kehilangan keuntungan dari akselerasi GPU. Jadi pertanyaan saya adalah, bisakah saya membiarkan kernel tetap terbuka dan mengalirkan partikel ke dalamnya?

Jika ada bantuan, inilah kernel saya saat ini:

__global__
void project(float *projection_matrix, float *vector, float *output_matrix) {
    int col_index = blockIdx.x * blockDim.x + threadIdx.x;
    int row_index = blockIdx.y * blockDim.x + threadIdx.y;
    int output_index = (col_index*3 + threadIdx.y);
    int transform_first_element = col_index * 9 + threadIdx.y * 3;
    int stride = blockDim.x*blockDim.y*gridDim.x;

    while (output_index < (NUMBER_OF_TRANSFORMS * 3)) {
        output_matrix[output_index] = projection_matrix[transform_first_element]*vector[0]+ projection_matrix[(transform_first_element+1)]*vector[1] + projection_matrix[(transform_first_element+2)]*vector[2];
        output_index += stride;
    }
}

dan di sinilah saya menyebutnya:

...
project <<<num_blocks_dim, block_dim >>> (transformationList, inputVector, outputMatrix);
cudaDeviceSynchronize();
...