48 lines
1.7 KiB
Common Lisp
48 lines
1.7 KiB
Common Lisp
typedef struct __attribute__((packed)) {
|
|
uint rows;
|
|
uint cols;
|
|
char transposed;
|
|
} cl_GPUMat;
|
|
|
|
#define matrixAt(mat, mat_values, r, c) mat_values[(!mat.transposed ? r * mat.cols + c : c * mat.rows + r)]
|
|
|
|
void mat_multiply(cl_GPUMat matA, __global float *matA_values, cl_GPUMat matB, __global float *matB_values, cl_GPUMat matOut, __global float *matOut_values) {
|
|
uint idx = get_global_id(0);
|
|
if(idx >= matOut.rows * matOut.cols) return;
|
|
|
|
uint i = idx / matOut.cols;
|
|
uint j = idx % matOut.cols;
|
|
|
|
float sum = 0;
|
|
for(unsigned int k = 0; k < matA.cols; k++) {
|
|
sum += matrixAt(matA, matA_values, i, k) * matrixAt(matB, matB_values, k, j);
|
|
}
|
|
matrixAt(matOut, matOut_values, i, j) = sum;
|
|
}
|
|
|
|
void mat_add(cl_GPUMat matA, __global float *matA_values, cl_GPUMat matB, __global float *matB_values) {
|
|
uint idx = get_global_id(0);
|
|
if(idx >= matA.rows * matA.cols) return;
|
|
|
|
matA_values[idx] += matB_values[idx];
|
|
}
|
|
|
|
void mat_sigmoid(cl_GPUMat mat, __global float *mat_values) {
|
|
uint idx = get_global_id(0);
|
|
if(idx >= mat.rows * mat.cols) return;
|
|
|
|
mat_values[idx] = 1 / (1 + exp(-mat_values[idx]));
|
|
}
|
|
|
|
// clm_Matrix input;
|
|
// clm_Matrix weights;
|
|
// clm_Matrix bias;
|
|
// clm_Matrix output;
|
|
|
|
__kernel void linear_forward(cl_GPUMat input, __global float *input_values, cl_GPUMat weights, __global float *weights_values, cl_GPUMat bias, __global float *bias_values, cl_GPUMat out, __global float *out_values) {
|
|
// FIXME mat_multiply possibly doesn't index at idx when out is transposed, which is important to ensure it always accesses the same index for all operations!
|
|
mat_multiply(weights, weights_values, input, input_values, out, out_values);
|
|
mat_add(out, out_values, bias, bias_values);
|
|
mat_sigmoid(out, out_values);
|
|
}
|