37 lines
999 B
Plaintext
37 lines
999 B
Plaintext
#ifndef __SOFTMAX_KERNEL_CUH__
|
|
#define __SOFTMAX_KERNEL_CUH__
|
|
|
|
#include <cuda_runtime.h>
|
|
#include <math.h>
|
|
#include <float.h>
|
|
#include "../utils/cuda_utils.cuh"
|
|
|
|
__global__ void softmax_forward_kernel(float* __restrict__ out, float inv_temperature, const float* __restrict__ inp, int N, int T) {
|
|
// Implement this
|
|
int x = blockDim.x * blockIdx.x + threadIdx.x;
|
|
|
|
if(x < N*T){
|
|
int own_pos = x % T;
|
|
const float* __restrict__ a = inp + x * T;
|
|
float* __restrict__ o = out + x * T;
|
|
|
|
float maxval = -FLT_MAX;
|
|
for (int i = 0; i <= own_pos; i++) {
|
|
maxval = fmaxf(maxval,a[i]);
|
|
}
|
|
|
|
float sumval = 0.0f;
|
|
for (int i = 0; i <= own_pos; ++i) {
|
|
float ev = expf(inv_temperature * (a[i] - maxval));
|
|
sumval += ev;
|
|
o[i] = ev;
|
|
}
|
|
|
|
float norm = 1.0f / sumval;
|
|
for (int i = 0; i <= own_pos; ++i) {
|
|
o[i] *= norm;
|
|
}
|
|
}
|
|
}
|
|
|
|
#endif // __SOFTMAX_KERNEL_CUH__ |