@@ -36,12 +36,12 @@ __global__ void KeGruForwardResetOutput(OpResetOutput op_reset_output,
36
36
T *prev_output_value, int frame_size,
37
37
int batch_size,
38
38
activation_mode_t active_gate) {
39
- const int frame_idx = block_idx .x * block_dim .x + thread_idx .x ;
39
+ const int frame_idx = blockIdx .x * blockDim .x + threadIdx .x ;
40
40
if (frame_idx >= frame_size) return ;
41
41
42
42
int batch_idx = 0 ;
43
43
if (is_batch) {
44
- batch_idx = block_idx .y * block_dim .y + thread_idx .y ;
44
+ batch_idx = blockIdx .y * blockDim .y + threadIdx .y ;
45
45
if (batch_idx >= batch_size) return ;
46
46
gate_value += batch_idx * 3 * frame_size;
47
47
reset_output_value += batch_idx * frame_size;
@@ -75,11 +75,11 @@ __global__ void KeGruForwardFinalOutput(OpFinalOutput op_final_output,
75
75
T *output_value, int frame_size,
76
76
int batch_size,
77
77
activation_mode_t active_node) {
78
- const int frame_idx = block_idx .x * block_dim .x + thread_idx .x ;
78
+ const int frame_idx = blockIdx .x * blockDim .x + threadIdx .x ;
79
79
if (frame_idx >= frame_size) return ;
80
80
int batch_idx = 0 ;
81
81
if (is_batch) {
82
- batch_idx = block_idx .y * block_dim .y + thread_idx .y ;
82
+ batch_idx = blockIdx .y * blockDim .y + threadIdx .y ;
83
83
if (batch_idx >= batch_size) return ;
84
84
gate_value += batch_idx * 3 * frame_size;
85
85
output_value += batch_idx * frame_size;
@@ -112,11 +112,11 @@ __global__ void KeGruBackwardStateGrad(OpStateGrad op_state_grad, T *gate_value,
112
112
T *prev_out_grad, T *output_grad,
113
113
int frame_size, int batch_size,
114
114
activation_mode_t active_node) {
115
- const int frame_idx = block_idx .x * block_dim .x + thread_idx .x ;
115
+ const int frame_idx = blockIdx .x * blockDim .x + threadIdx .x ;
116
116
if (frame_idx >= frame_size) return ;
117
117
int batch_idx = 0 ;
118
118
if (is_batch) {
119
- batch_idx = block_idx .y * block_dim .y + thread_idx .y ;
119
+ batch_idx = blockIdx .y * blockDim .y + threadIdx .y ;
120
120
if (batch_idx >= batch_size) return ;
121
121
gate_value += batch_idx * 3 * frame_size;
122
122
gate_grad += batch_idx * 3 * frame_size;
@@ -160,11 +160,11 @@ __global__ void KeGruBackwardResetGrad(OpResetGrad op_reset_grad, T *gate_value,
160
160
T *prev_out_grad, T *reset_output_grad,
161
161
int frame_size, int batch_size,
162
162
activation_mode_t active_gate) {
163
- const int frame_idx = block_idx .x * block_dim .x + thread_idx .x ;
163
+ const int frame_idx = blockIdx .x * blockDim .x + threadIdx .x ;
164
164
if (frame_idx >= frame_size) return ;
165
165
int batch_idx = 0 ;
166
166
if (is_batch) {
167
- batch_idx = block_idx .y * block_dim .y + thread_idx .y ;
167
+ batch_idx = blockIdx .y * blockDim .y + threadIdx .y ;
168
168
if (batch_idx >= batch_size) return ;
169
169
gate_value += batch_idx * 3 * frame_size;
170
170
gate_grad += batch_idx * 3 * frame_size;
0 commit comments