Kernels
iamwyldecat commited on
Commit
32c2bde
·
1 Parent(s): e85ecc9

refactor(poly-norm): fix indentation

Browse files
Files changed (1) hide show
  1. activation/activation_kernels.cu +63 -62
activation/activation_kernels.cu CHANGED
@@ -130,60 +130,60 @@ __global__ void poly_norm_backward_kernel(
130
  acc_t sq_mean_4 = rsqrtf(mean_4) * mean_4;
131
  acc_t sq_mean_6 = rsqrtf(mean_6) * mean_6;
132
 
133
- acc_t denom_2 = mean_2 * sq_mean_2;
134
- acc_t denom_4 = mean_4 * sq_mean_4;
135
- acc_t denom_6 = mean_6 * sq_mean_6;
136
-
137
- acc_t sum_dw0 = 0;
138
- acc_t sum_dw1 = 0;
139
- acc_t sum_dw2 = 0;
140
- acc_t sum_db = 0;
141
-
142
- for (int64_t idx = threadIdx.x; idx < d; idx += blockDim.x) {
143
- acc_t dy = output_grad[token_idx * d + idx];
144
- acc_t x_1 = input[token_idx * d + idx];
145
- acc_t x_2 = x_1 * x_1;
146
- acc_t x_3 = x_2 * x_1;
147
-
148
- acc_t _dx_1 = w2 * dy;
149
- acc_t _dx_2 = w1 * dy;
150
- acc_t _dx_3 = w0 * dy;
151
-
152
- acc_t dx_3 =
153
- 3 * x_2 * (_dx_3 / sq_mean_6 - x_3 * sum_dx_3 / (d * denom_6));
154
- acc_t dx_2 =
155
- 2 * x_1 * (_dx_2 / sq_mean_4 - x_2 * sum_dx_2 / (d * denom_4));
156
- acc_t dx_1 =
157
- _dx_1 / sq_mean_2 - x_1 * sum_dx_1 / (d * denom_2);
158
-
159
- if (input_grad) {
160
- input_grad[token_idx * d + idx] = dx_1 + dx_2 + dx_3;
161
- }
162
-
163
- sum_dw0 += dy * (x_3 / sq_mean_6);
164
- sum_dw1 += dy * (x_2 / sq_mean_4);
165
- sum_dw2 += dy * (x_1 / sq_mean_2);
166
- sum_db += dy;
167
- }
168
-
169
- if (weight_grad) {
170
- sum_dw0 = _block_reduce_sum<acc_t, BLOCK_SIZE>(shared, sum_dw0, d);
171
- sum_dw1 = _block_reduce_sum<acc_t, BLOCK_SIZE>(shared, sum_dw1, d);
172
- sum_dw2 = _block_reduce_sum<acc_t, BLOCK_SIZE>(shared, sum_dw2, d);
173
-
174
- if (threadIdx.x == 0) {
175
- atomic_add(&weight_grad[0], sum_dw0);
176
- atomic_add(&weight_grad[1], sum_dw1);
177
- atomic_add(&weight_grad[2], sum_dw2);
178
- }
179
- }
180
-
181
- if (bias_grad) {
182
- sum_db = _block_reduce_sum<acc_t, BLOCK_SIZE>(shared, sum_db, d);
183
- if (threadIdx.x == 0) {
184
- atomic_add(&bias_grad[0], sum_db);
185
- }
186
- }
187
  }
188
 
189
  } // namespace motif
@@ -191,9 +191,9 @@ __global__ void poly_norm_backward_kernel(
191
 
192
  void poly_norm(torch::Tensor& out, // [..., d]
193
  const torch::Tensor& input, // [..., d]
194
- const torch::Tensor& weight, // [3]
195
- const torch::Tensor& bias, // [1]
196
- double eps)
197
  {
198
  AssertTensorShapeEqual(input, out, "input", "out");
199
  AssertTensorNotNull(weight, "weight");
@@ -244,11 +244,12 @@ void poly_norm_backward(
244
 
245
  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
246
 
247
- if (weight_grad.defined())
248
- cudaMemsetAsync(weight_grad.data_ptr(), 0, weight_grad.numel() * weight_grad.element_size(), stream);
249
- if (bias_grad.defined()) {
250
- cudaMemsetAsync(bias_grad.data_ptr(), 0, bias_grad.numel() * bias_grad.element_size(), stream);
251
- }
 
252
 
253
  const at::cuda::OptionalCUDAGuard device_guard(device_of(input));
254
  MOTIF_DISPATCH_FLOATING_TYPES(
 
130
  acc_t sq_mean_4 = rsqrtf(mean_4) * mean_4;
131
  acc_t sq_mean_6 = rsqrtf(mean_6) * mean_6;
132
 
133
+ acc_t denom_2 = mean_2 * sq_mean_2;
134
+ acc_t denom_4 = mean_4 * sq_mean_4;
135
+ acc_t denom_6 = mean_6 * sq_mean_6;
136
+
137
+ acc_t sum_dw0 = 0;
138
+ acc_t sum_dw1 = 0;
139
+ acc_t sum_dw2 = 0;
140
+ acc_t sum_db = 0;
141
+
142
+ for (int64_t idx = threadIdx.x; idx < d; idx += blockDim.x) {
143
+ acc_t dy = output_grad[token_idx * d + idx];
144
+ acc_t x_1 = input[token_idx * d + idx];
145
+ acc_t x_2 = x_1 * x_1;
146
+ acc_t x_3 = x_2 * x_1;
147
+
148
+ acc_t _dx_1 = w2 * dy;
149
+ acc_t _dx_2 = w1 * dy;
150
+ acc_t _dx_3 = w0 * dy;
151
+
152
+ acc_t dx_3 =
153
+ 3 * x_2 * (_dx_3 / sq_mean_6 - x_3 * sum_dx_3 / (d * denom_6));
154
+ acc_t dx_2 =
155
+ 2 * x_1 * (_dx_2 / sq_mean_4 - x_2 * sum_dx_2 / (d * denom_4));
156
+ acc_t dx_1 =
157
+ _dx_1 / sq_mean_2 - x_1 * sum_dx_1 / (d * denom_2);
158
+
159
+ if (input_grad) {
160
+ input_grad[token_idx * d + idx] = dx_1 + dx_2 + dx_3;
161
+ }
162
+
163
+ sum_dw0 += dy * (x_3 / sq_mean_6);
164
+ sum_dw1 += dy * (x_2 / sq_mean_4);
165
+ sum_dw2 += dy * (x_1 / sq_mean_2);
166
+ sum_db += dy;
167
+ }
168
+
169
+ if (weight_grad) {
170
+ sum_dw0 = _block_reduce_sum<acc_t, BLOCK_SIZE>(shared, sum_dw0, d);
171
+ sum_dw1 = _block_reduce_sum<acc_t, BLOCK_SIZE>(shared, sum_dw1, d);
172
+ sum_dw2 = _block_reduce_sum<acc_t, BLOCK_SIZE>(shared, sum_dw2, d);
173
+
174
+ if (threadIdx.x == 0) {
175
+ atomic_add(&weight_grad[0], sum_dw0);
176
+ atomic_add(&weight_grad[1], sum_dw1);
177
+ atomic_add(&weight_grad[2], sum_dw2);
178
+ }
179
+ }
180
+
181
+ if (bias_grad) {
182
+ sum_db = _block_reduce_sum<acc_t, BLOCK_SIZE>(shared, sum_db, d);
183
+ if (threadIdx.x == 0) {
184
+ atomic_add(&bias_grad[0], sum_db);
185
+ }
186
+ }
187
  }
188
 
189
  } // namespace motif
 
191
 
192
  void poly_norm(torch::Tensor& out, // [..., d]
193
  const torch::Tensor& input, // [..., d]
194
+ const torch::Tensor& weight, // [3]
195
+ const torch::Tensor& bias, // [1]
196
+ double eps)
197
  {
198
  AssertTensorShapeEqual(input, out, "input", "out");
199
  AssertTensorNotNull(weight, "weight");
 
244
 
245
  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
246
 
247
+ if (weight_grad.defined()) {
248
+ cudaMemsetAsync(weight_grad.data_ptr(), 0, weight_grad.numel() * weight_grad.element_size(), stream);
249
+ }
250
+ if (bias_grad.defined()) {
251
+ cudaMemsetAsync(bias_grad.data_ptr(), 0, bias_grad.numel() * bias_grad.element_size(), stream);
252
+ }
253
 
254
  const at::cuda::OptionalCUDAGuard device_guard(device_of(input));
255
  MOTIF_DISPATCH_FLOATING_TYPES(