| 74 | |
| 75 | template <typename T, typename U, int kUnroll> |
| 76 | __global__ void GenerateNormalizedProb(const T* logits, const U* sum_probs, |
| 77 | const T* max_logits, T* output, |
| 78 | const int num_rows, const int num_cols, |
| 79 | const bool in_log_space) { |
| 80 | int tid = blockIdx.x * blockDim.x + threadIdx.x; |
| 81 | int row, col; |
| 82 | |
| 83 | // TODO(jamesqin): change to half2 load when inputs are Eigen::half. |
| 84 | U input[kUnroll]; |
| 85 | U max_val[kUnroll]; |
| 86 | U result[kUnroll]; |
| 87 | for (int i = 0; i < kUnroll; i++) { |
| 88 | row = tid / num_cols; |
| 89 | col = tid % num_cols; |
| 90 | if (row < num_rows && col < num_cols) { |
| 91 | input[i] = strict_cast<U>(logits[tid]); |
| 92 | max_val[i] = strict_cast<U>(ldg(max_logits + row)); |
| 93 | } |
| 94 | tid += gridDim.x * blockDim.x; |
| 95 | } |
| 96 | |
| 97 | tid = blockIdx.x * blockDim.x + threadIdx.x; |
| 98 | for (int i = 0; i < kUnroll; i++) { |
| 99 | row = tid / num_cols; |
| 100 | col = tid % num_cols; |
| 101 | if (row < num_rows && col < num_cols) { |
| 102 | if (in_log_space) { |
| 103 | result[i] = input[i] - max_val[i] - log(ldg(sum_probs + row)); |
| 104 | } else { |
| 105 | result[i] = exp(input[i] - max_val[i]) / ldg(sum_probs + row); |
| 106 | } |
| 107 | output[tid] = strict_cast<T>(result[i]); |
| 108 | } |
| 109 | tid += gridDim.x * blockDim.x; |
| 110 | } |
| 111 | } |
| 112 | |
| 113 | template <> |
| 114 | __global__ void GenerateNormalizedProb<Eigen::half, float, 8>( |