| 103 | // *********************** compresse transformer input *********************** |
| 104 | template <typename T> |
| 105 | __global__ void compress_bert_input(const T *from_tensor, T *to_tensor, const int *batch_idx, |
| 106 | const int *word_idx, int hidden_dim) { |
| 107 | int offset = __ldg(&word_idx[blockIdx.x]); |
| 108 | int dst_idx = blockIdx.x * hidden_dim + threadIdx.x; |
| 109 | int src_idx = offset * hidden_dim + threadIdx.x; |
| 110 | ((float4 *)to_tensor)[dst_idx] = ((const float4 *)from_tensor)[src_idx]; |
| 111 | } |
| 112 | |
| 113 | template <typename T> |
| 114 | void compressBertInput_kernelLauncher(const T *from_tensor, T *to_tensor, int *batch_idx, |
nothing calls this directly
no outgoing calls
no test coverage detected