diff --git a/build/cpu.o b/build/cpu.o index 7a91d37..6289686 100644 Binary files a/build/cpu.o and b/build/cpu.o differ diff --git a/build/cuda.cu.o b/build/cuda.cu.o index 069d281..6d729b8 100644 Binary files a/build/cuda.cu.o and b/build/cuda.cu.o differ diff --git a/build/libtensor.so b/build/libtensor.so index 5722336..2bb1564 100755 Binary files a/build/libtensor.so and b/build/libtensor.so differ diff --git a/build/tensor.o b/build/tensor.o index fa22bdf..49c51f6 100644 Binary files a/build/tensor.o and b/build/tensor.o differ diff --git a/norch/__pycache__/tensor.cpython-38.pyc b/norch/__pycache__/tensor.cpython-38.pyc index 8ee1ec7..343e53b 100644 Binary files a/norch/__pycache__/tensor.cpython-38.pyc and b/norch/__pycache__/tensor.cpython-38.pyc differ diff --git a/norch/autograd/__pycache__/functions.cpython-38.pyc b/norch/autograd/__pycache__/functions.cpython-38.pyc index 6dc5d32..e5d728b 100644 Binary files a/norch/autograd/__pycache__/functions.cpython-38.pyc and b/norch/autograd/__pycache__/functions.cpython-38.pyc differ diff --git a/norch/autograd/functions.py b/norch/autograd/functions.py index 78d992d..3cdf8df 100644 --- a/norch/autograd/functions.py +++ b/norch/autograd/functions.py @@ -1,3 +1,5 @@ +import math + class AddBackward: def __init__(self, x, y): self.input = [x, y] @@ -35,13 +37,42 @@ class MatmulBackward: x, y = self.input return [gradient @ y.transpose(-1,-2), x.transpose(-1,-2) @ gradient] -class PowBackward: +"""class PowBackward: def __init__(self, x, power): self.input = [x] self.power = power def backward(self, gradient): - return [(gradient * self.power) * (self.input[0]) ** (self.power - 1)] + return [(gradient * self.power) * (self.input[0]) ** (self.power - 1)]""" + +class PowBackward: + def __init__(self, base, exponent): + self.input = [base, exponent] + + def backward(self, gradient): + base, exponent = self.input[0], self.input[1] + + if isinstance(base, (int, float)): + grad_base = gradient * (base ** (exponent - 1)) + grad_exponent = (gradient * base ** exponent) * math.log(base) + + else: + grad_base = gradient * exponent * (base ** (exponent - 1)) + grad_exponent = (gradient * base ** exponent) * (base.log()) + + return [grad_base, grad_exponent] + + +class LogBackward: + def __init__(self, input_tensor): + self.input_tensor = input_tensor + + def backward(self, gradient): + input_tensor = self.input_tensor + + grad_input = gradient / input_tensor + + return [grad_input] class SumBackward: def __init__(self, x): @@ -66,3 +97,14 @@ class TransposeBackward: def backward(self, gradient): return [gradient.transpose(self.axis2, self.axis1)] + +class DivisionBackward: + def __init__(self, x, y): + self.input = [x, y] + + def backward(self, gradient): + x, y = self.input + grad_x = gradient / y + grad_y = -1 * gradient * (x / (y * y)) + return [grad_x, grad_y] + diff --git a/norch/csrc/cpu.cpp b/norch/csrc/cpu.cpp index c45a8e9..13a658f 100644 --- a/norch/csrc/cpu.cpp +++ b/norch/csrc/cpu.cpp @@ -18,6 +18,7 @@ void sub_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data) { } } + void elementwise_mul_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data) { for (int i = 0; i < tensor1->size; i++) { @@ -32,6 +33,20 @@ void scalar_mul_tensor_cpu(Tensor* tensor, float scalar, float* result_data) { } } +void scalar_div_tensor_cpu(float scalar, Tensor* tensor, float* result_data) { + + for (int i = 0; i < tensor->size; i++) { + result_data[i] = scalar / tensor->data[i]; + } +} + +void tensor_div_scalar_cpu(Tensor* tensor, float scalar, float* result_data) { + + for (int i = 0; i < tensor->size; i++) { + result_data[i] = tensor->data[i] / scalar; + } +} + void matmul_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data) { for (int i = 0; i < tensor1->shape[0]; i++) { @@ -85,10 +100,23 @@ void batched_matmul_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_d } } -void pow_tensor_cpu(Tensor* tensor, float power, float* result_data) { +void scalar_pow_tensor_cpu(float base, Tensor* tensor, float* result_data) { for (int i = 0; i < tensor->size; i++) { - result_data[i] = powf(tensor->data[i], power); + result_data[i] = powf(base, tensor->data[i]); + } +} + +void tensor_pow_scalar_cpu(Tensor* tensor, float exponent, float* result_data) { + + for (int i = 0; i < tensor->size; i++) { + result_data[i] = powf(tensor->data[i], exponent); + } +} + +void log_tensor_cpu(Tensor* tensor, float* result_data) { + for (int i = 0; i < tensor->size; i++) { + result_data[i] = logf(tensor->data[i]); } } diff --git a/norch/csrc/cpu.h b/norch/csrc/cpu.h index a3c5879..1cc52fb 100644 --- a/norch/csrc/cpu.h +++ b/norch/csrc/cpu.h @@ -7,10 +7,14 @@ void add_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data); void sum_tensor_cpu(Tensor* tensor1, float* result_data); void sub_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data); void elementwise_mul_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data); +void scalar_div_tensor_cpu(float scalar, Tensor* tensor, float* result_data); +void tensor_div_scalar_cpu(Tensor* tensor, float scalar, float* result_data); void matmul_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data); void broadcasted_batched_matmul_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data); void batched_matmul_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data); -void pow_tensor_cpu(Tensor* tensor, float power, float* result_data); +void scalar_pow_tensor_cpu(float base, Tensor* tensor, float* result_data); +void tensor_pow_scalar_cpu(Tensor* tensor, float exponent, float* result_data); +void log_tensor_cpu(Tensor* tensor, float* result_data); void scalar_mul_tensor_cpu(Tensor* tensor, float scalar, float* result_data); void ones_like_tensor_cpu(Tensor* tensor, float* result_data); void zeros_like_tensor_cpu(Tensor* tensor, float* result_data); diff --git a/norch/csrc/cuda.cu b/norch/csrc/cuda.cu index afa9820..67dc8a0 100644 --- a/norch/csrc/cuda.cu +++ b/norch/csrc/cuda.cu @@ -107,9 +107,6 @@ __host__ void sum_tensor_cuda(Tensor* tensor, float* result_data) { cudaDeviceSynchronize(); } - - - __global__ void sub_tensor_cuda_kernel(float* data1, float* data2, float* result_data, int size) { int i = blockIdx.x * blockDim.x + threadIdx.x; @@ -176,6 +173,50 @@ __host__ void scalar_mul_tensor_cuda(Tensor* tensor, float scalar, float* result cudaDeviceSynchronize(); } +__global__ void scalar_div_tensor_cuda_kernel(float scalar, float* data, float* result_data, int size) { + + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < size) { + result_data[i] = scalar / data[i]; + } +} + +__host__ void scalar_div_tensor_cuda(float scalar, Tensor* tensor, float* result_data) { + + int number_of_blocks = (tensor->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; + scalar_div_tensor_cuda_kernel<<>>(scalar, tensor->data, result_data, tensor->size); + + cudaError_t error = cudaGetLastError(); + if (error != cudaSuccess) { + printf("CUDA error: %s\n", cudaGetErrorString(error)); + exit(-1); + } + + cudaDeviceSynchronize(); +} + +__global__ void tensor_div_scalar_cuda_kernel(float* data, float scalar, float* result_data, int size) { + + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < size) { + result_data[i] = data[i] / scalar; + } +} + +__host__ void tensor_div_scalar_cuda(Tensor* tensor, float scalar, float* result_data) { + + int number_of_blocks = (tensor->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; + tensor_div_scalar_cuda_kernel<<>>(tensor->data, scalar, result_data, tensor->size); + + cudaError_t error = cudaGetLastError(); + if (error != cudaSuccess) { + printf("CUDA error: %s\n", cudaGetErrorString(error)); + exit(-1); + } + + cudaDeviceSynchronize(); +} + /*__global__ void matmul_tensor_cuda_kernel(float* data1, float* data2, float* result_data, int rows1, int cols1, int cols2) { int row = blockIdx.y * blockDim.y + threadIdx.y; @@ -258,18 +299,18 @@ __host__ void matmul_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result cudaDeviceSynchronize(); } -__global__ void pow_tensor_cuda_kernel(float* data, float power, float* result_data, int size) { +__global__ void tensor_pow_scalar_cuda_kernel(float* data, float exponent, float* result_data, int size) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < size) { - result_data[i] = powf(data[i], power); + result_data[i] = powf(data[i], exponent); } } -__host__ void pow_tensor_cuda(Tensor* tensor, float power, float* result_data) { +__host__ void tensor_pow_scalar_cuda(Tensor* tensor, float exponent, float* result_data) { int number_of_blocks = (tensor->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; - pow_tensor_cuda_kernel<<>>(tensor->data, power, result_data, tensor->size); + tensor_pow_scalar_cuda_kernel<<>>(tensor->data, exponent, result_data, tensor->size); cudaError_t error = cudaGetLastError(); if (error != cudaSuccess) { @@ -280,6 +321,51 @@ __host__ void pow_tensor_cuda(Tensor* tensor, float power, float* result_data) { cudaDeviceSynchronize(); } +__global__ void scalar_pow_tensor_cuda_kernel(float base, float* data, float* result_data, int size) { + + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < size) { + result_data[i] = powf(base, data[i]); + } +} + +__host__ void scalar_pow_tensor_cuda(float base, Tensor* tensor, float* result_data) { + + int number_of_blocks = (tensor->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; + scalar_pow_tensor_cuda_kernel<<>>(base, tensor->data, result_data, tensor->size); + + cudaError_t error = cudaGetLastError(); + if (error != cudaSuccess) { + printf("CUDA error: %s\n", cudaGetErrorString(error)); + exit(-1); + } + + cudaDeviceSynchronize(); +} + +__global__ void log_tensor_cuda_kernel(float* data, float* result_data, int size) { + + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < size) { + result_data[i] = logf(data[i]); + } +} + +__host__ void log_tensor_cuda(Tensor* tensor, float* result_data) { + + int number_of_blocks = (tensor->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; + log_tensor_cuda_kernel<<>>(tensor->data, result_data, tensor->size); + + cudaError_t error = cudaGetLastError(); + if (error != cudaSuccess) { + printf("CUDA error: %s\n", cudaGetErrorString(error)); + exit(-1); + } + + cudaDeviceSynchronize(); +} + + __global__ void ones_like_tensor_cuda_kernel(float* data, float* result_data, int size) { int i = blockIdx.x * blockDim.x + threadIdx.x; diff --git a/norch/csrc/cuda.h b/norch/csrc/cuda.h index 48c87b8..87ef884 100644 --- a/norch/csrc/cuda.h +++ b/norch/csrc/cuda.h @@ -19,11 +19,23 @@ __global__ void scalar_mul_tensor_cuda_kernel(float* data, float scalar, float* result_data, int size); __host__ void scalar_mul_tensor_cuda(Tensor* tensor, float scalar, float* result_data); + __global__ void scalar_div_tensor_cuda_kernel(float scalar, float* data, float* result_data, int size); + __host__ void scalar_div_tensor_cuda(float scalar, Tensor* tensor, float* result_data); + + __global__ void tensor_div_scalar_cuda_kernel(float* data, float scalar, float* result_data, int size); + __host__ void tensor_div_scalar_cuda(Tensor* tensor, float scalar, float* result_data); + __global__ void matmul_tensor_cuda_kernel(float* data1, float* data2, float* result_data, int rows1, int cols1, int cols2); __host__ void matmul_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_data); - __global__ void pow_tensor_cuda_kernel(float* data, float power, float* result_data, int size); - __host__ void pow_tensor_cuda(Tensor* tensor, float power, float* result_data); + __global__ void tensor_pow_scalar_cuda_kernel(float* data, float exponent, float* result_data, int size); + __host__ void tensor_pow_scalar_cuda(Tensor* tensor, float exponent, float* result_data); + + __global__ void scalar_pow_tensor_cuda_kernel(float base, float* data, float* result_data, int size); + __host__ void scalar_pow_tensor_cuda(float base, Tensor* tensor, float* result_data); + + __global__ void log_tensor_cuda_kernel(float* data, float* result_data, int size); + __host__ void log_tensor_cuda(Tensor* tensor, float* result_data); __global__ void ones_like_tensor_cuda_kernel(float* data, float* result_data, int size); __host__ void ones_like_tensor_cuda(Tensor* tensor, float* result_data); diff --git a/norch/csrc/tensor.cpp b/norch/csrc/tensor.cpp index 03cc657..05232d0 100644 --- a/norch/csrc/tensor.cpp +++ b/norch/csrc/tensor.cpp @@ -335,6 +335,82 @@ extern "C" { } } + Tensor* scalar_div_tensor(float scalar, Tensor* tensor) { + + char* device = (char*)malloc(strlen(tensor->device) + 1); + if (device != NULL) { + strcpy(device, tensor->device); + } else { + fprintf(stderr, "Memory allocation failed\n"); + exit(-1); + } + int ndim = tensor->ndim; + int* shape = (int*)malloc(ndim * sizeof(int)); + if (shape == NULL) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + + for (int i = 0; i < ndim; i++) { + shape[i] = tensor->shape[i]; + } + + if (strcmp(tensor->device, "cuda") == 0) { + + float* result_data; + cudaMalloc((void **)&result_data, tensor->size * sizeof(float)); + scalar_div_tensor_cuda(scalar, tensor, result_data); + return create_tensor(result_data, shape, ndim, device); + } + else { + float* result_data = (float*)malloc(tensor->size * sizeof(float)); + if (result_data == NULL) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + scalar_div_tensor_cpu(scalar, tensor, result_data); + return create_tensor(result_data, shape, ndim, device); + } + } + + Tensor* tensor_div_scalar(Tensor* tensor, float scalar) { + + char* device = (char*)malloc(strlen(tensor->device) + 1); + if (device != NULL) { + strcpy(device, tensor->device); + } else { + fprintf(stderr, "Memory allocation failed\n"); + exit(-1); + } + int ndim = tensor->ndim; + int* shape = (int*)malloc(ndim * sizeof(int)); + if (shape == NULL) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + + for (int i = 0; i < ndim; i++) { + shape[i] = tensor->shape[i]; + } + + if (strcmp(tensor->device, "cuda") == 0) { + + float* result_data; + cudaMalloc((void **)&result_data, tensor->size * sizeof(float)); + tensor_div_scalar_cuda(tensor, scalar, result_data); + return create_tensor(result_data, shape, ndim, device); + } + else { + float* result_data = (float*)malloc(tensor->size * sizeof(float)); + if (result_data == NULL) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + tensor_div_scalar_cpu(tensor, scalar, result_data); + return create_tensor(result_data, shape, ndim, device); + } + } + Tensor* matmul_tensor(Tensor* tensor1, Tensor* tensor2) { //MxN @ NxP = MxP // Check if tensors have compatible shapes for matrix multiplication @@ -528,7 +604,7 @@ extern "C" { } - Tensor* pow_tensor(Tensor* tensor, float power) { + Tensor* tensor_pow_scalar(Tensor* tensor, float exponent) { char* device = (char*)malloc(strlen(tensor->device) + 1); if (device != NULL) { strcpy(device, tensor->device); @@ -551,7 +627,7 @@ extern "C" { float* result_data; cudaMalloc((void **)&result_data, tensor->size * sizeof(float)); - pow_tensor_cuda(tensor, power, result_data); + tensor_pow_scalar_cuda(tensor, exponent, result_data); return create_tensor(result_data, shape, ndim, device); } else { @@ -560,7 +636,81 @@ extern "C" { fprintf(stderr, "Memory allocation failed\n"); exit(1); } - pow_tensor_cpu(tensor, power, result_data); + tensor_pow_scalar_cpu(tensor, exponent, result_data); + return create_tensor(result_data, shape, ndim, device); + } + } + + Tensor* scalar_pow_tensor(float base, Tensor* tensor) { + char* device = (char*)malloc(strlen(tensor->device) + 1); + if (device != NULL) { + strcpy(device, tensor->device); + } else { + fprintf(stderr, "Memory allocation failed\n"); + exit(-1); + } + int ndim = tensor->ndim; + int* shape = (int*)malloc(ndim * sizeof(int)); + if (shape == NULL) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + + for (int i = 0; i < ndim; i++) { + shape[i] = tensor->shape[i]; + } + + if (strcmp(tensor->device, "cuda") == 0) { + + float* result_data; + cudaMalloc((void **)&result_data, tensor->size * sizeof(float)); + scalar_pow_tensor_cuda(base, tensor, result_data); + return create_tensor(result_data, shape, ndim, device); + } + else { + float* result_data = (float*)malloc(tensor->size * sizeof(float)); + if (result_data == NULL) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + scalar_pow_tensor_cpu(base, tensor, result_data); + return create_tensor(result_data, shape, ndim, device); + } + } + + Tensor* log_tensor(Tensor* tensor) { + char* device = (char*)malloc(strlen(tensor->device) + 1); + if (device != NULL) { + strcpy(device, tensor->device); + } else { + fprintf(stderr, "Memory allocation failed\n"); + exit(-1); + } + int ndim = tensor->ndim; + int* shape = (int*)malloc(ndim * sizeof(int)); + if (shape == NULL) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + + for (int i = 0; i < ndim; i++) { + shape[i] = tensor->shape[i]; + } + + if (strcmp(tensor->device, "cuda") == 0) { + + float* result_data; + cudaMalloc((void **)&result_data, tensor->size * sizeof(float)); + log_tensor_cuda(tensor, result_data); + return create_tensor(result_data, shape, ndim, device); + } + else { + float* result_data = (float*)malloc(tensor->size * sizeof(float)); + if (result_data == NULL) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + log_tensor_cpu(tensor, result_data); return create_tensor(result_data, shape, ndim, device); } } diff --git a/norch/csrc/tensor.h b/norch/csrc/tensor.h index dab0b23..a9b8f8e 100644 --- a/norch/csrc/tensor.h +++ b/norch/csrc/tensor.h @@ -20,9 +20,13 @@ extern "C" { Tensor* sub_tensor(Tensor* tensor1, Tensor* tensor2); Tensor* elementwise_mul_tensor(Tensor* tensor1, Tensor* tensor2); Tensor* scalar_mul_tensor(Tensor* tensor, float scalar); + Tensor* scalar_div_tensor(float scalar, Tensor* tensor); + Tensor* tensor_div_scalar(Tensor* tensor, float scalar); Tensor* reshape_tensor(Tensor* tensor, int* new_shape, int new_ndim); Tensor* matmul_tensor(Tensor* tensor1, Tensor* tensor2); - Tensor* pow_tensor(Tensor* tensor, float power); + Tensor* tensor_pow_scalar(Tensor* tensor, float exponent); + Tensor* scalar_pow_tensor(float base, Tensor* tensor); + Tensor* log_tensor(Tensor* tensor); void to_device(Tensor* tensor, char* device); Tensor* ones_like_tensor(Tensor* tensor); Tensor* zeros_like_tensor(Tensor* tensor); diff --git a/norch/nn/__init__.py b/norch/nn/__init__.py new file mode 100644 index 0000000..fc6fa1a --- /dev/null +++ b/norch/nn/__init__.py @@ -0,0 +1,2 @@ +from .modules import * +from .activations import * \ No newline at end of file diff --git a/norch/nn/activation.py b/norch/nn/activation.py new file mode 100644 index 0000000..1ab272f --- /dev/null +++ b/norch/nn/activation.py @@ -0,0 +1,20 @@ +from module import Module +import math + +class Activation(Module): + """ + Abstract classes for activations + """ + def __init__(self): + super(Activation, self).__init__() + + def forward(self, x): + raise NotImplementedError + + +class Sigmoid(Activation): + def __init__(self): + super().__init__() + + def forward(self, x): + return 1.0 / (1.0 + (math.e) ** (-x)) \ No newline at end of file diff --git a/norch/nn/module.py b/norch/nn/module.py new file mode 100644 index 0000000..88b408a --- /dev/null +++ b/norch/nn/module.py @@ -0,0 +1,81 @@ +from parameter import Parameter +from collections import OrderedDict +from abc import ABC +import pickle +import json +import inspect +import warnings + +class Module(ABC): + """ + Abstract class for modules + """ + def __init__(self): + self._modules = OrderedDict() + self._params = OrderedDict() + self._grads = OrderedDict() + self.training = True + + def __call__(self, *inputs, **kwargs): + return self.forward(*inputs, **kwargs) + + def train(self): + self.training = True + for param in self.parameters(): + param.requires_grad = True + + def eval(self): + self.training = False + for param in self.parameters(): + param.requires_grad = False + + def parameters(self): + for name, value in inspect.getmembers(self): + if isinstance(value, Parameter): + yield value + elif isinstance(value, Module): + yield from value.parameters() + + def to(self, device): + for parameter in self.parameters(): + parameter.to(device) + + def state_dict(self): + state = OrderedDict() + for i, param in enumerate(self.parameters()): + state[f'param{i}'] = param.tolist() + return state + + def load_state(self, state_dict): + for i, param in self.parameters(): + data = state_dict[f'param{i}'] + if param.shape != data.shape: + warnings.warn(f"The 'state_dict' shape does not match model's parameter shape. " + f"Got {data.shape}, expected {param.shape}.") + param.data = Parameter(data=data) + + def save(self, filename='model.pickle'): + with open(filename, 'wb') as f: + pickle.dump(self, f) + + def save_dict(self, filename='state_dict.json'): + state = self.state_dict() + with open(filename, 'w') as f: + json.dump(state, f) + + def inner_repr(self): + return "" + + def __repr__(self): + string = f"{self.get_name()}(" + tab = " " + modules = self._modules + if modules == {}: + string += f'\n{tab}(parameters): {self.inner_repr()}' + else: + for key, module in modules.items(): + string += f"\n{tab}({key}): {module.get_name()}({module.inner_repr()})" + return f'{string}\n)' + + def get_name(self): + return self.__class__.__name__ \ No newline at end of file diff --git a/norch/nn/modules/linear.py b/norch/nn/modules/linear.py new file mode 100644 index 0000000..032670e --- /dev/null +++ b/norch/nn/modules/linear.py @@ -0,0 +1,18 @@ +from module import Module +from parameter import Parameter + +class Linear(Module): + def __init__(self, input_dim, output_dim): + super().__init__() + self.input_dim = input_dim + self.output_dim = output_dim + self.weight = Parameter(shape=[self.input_dim, self.output_dim]) + self.bias = Parameter(shape=[self.output_dim, 1]) + + def forward(self, x): + z = self.weight @ x + self.bias + return z + + def inner_repr(self): + return f"input_dim={self.input_dim}, output_dim={self.output_dim}, " \ + f"bias={True if self.bias is not None else False}" \ No newline at end of file diff --git a/norch/nn/parameter.py b/norch/nn/parameter.py new file mode 100644 index 0000000..152d951 --- /dev/null +++ b/norch/nn/parameter.py @@ -0,0 +1,14 @@ +from tensor import Tensor +import random + +class Parameter(Tensor): + """ + A parameter is a trainable tensor. + """ + def __init__(self, shape): + data = [] + for dim_size in reversed(shape): + random_dim = [random.random() for _ in range(dim_size)] + data.insert(0, random_dim) + + super().__init__(data, requires_grad=True) \ No newline at end of file diff --git a/norch/tensor.py b/norch/tensor.py index b2844e3..31c5258 100644 --- a/norch/tensor.py +++ b/norch/tensor.py @@ -344,14 +344,93 @@ class Tensor: return result_data - def __pow__(self, power): + def __pow__(self, other): + if isinstance(other, (int, float)): + other = float(other) + Tensor._C.tensor_pow_scalar.argtypes = [ctypes.POINTER(CTensor), ctypes.c_float] + Tensor._C.tensor_pow_scalar.restype = ctypes.POINTER(CTensor) + + result_tensor_ptr = Tensor._C.tensor_pow_scalar(self.tensor, ctypes.c_float(other)) + + result_data = Tensor() + result_data.tensor = result_tensor_ptr + result_data.shape = self.shape.copy() + result_data.ndim = self.ndim + result_data.device = self.device + + result_data.requires_grad = self.requires_grad + if result_data.requires_grad: + result_data.grad_fn = PowBackward(self, other) + + elif isinstance(self, (int, float)): + self = float(self) + Tensor._C.scalar_pow_tensor.argtypes = [ctypes.c_float, ctypes.POINTER(CTensor)] + Tensor._C.scalar_pow_tensor.restype = ctypes.POINTER(CTensor) + + result_tensor_ptr = Tensor._C.scalar_pow_tensor(ctypes.c_float(self), other.tensor) + + result_data = Tensor() + result_data.tensor = result_tensor_ptr + result_data.shape = other.shape.copy() + result_data.ndim = other.ndim + result_data.device = other.device + + result_data.requires_grad = other.requires_grad + if result_data.requires_grad: + result_data.grad_fn = PowBackward(other, self) - power_ctypes = ctypes.c_float(power) + else: + raise ValueError("Tensor to tensor power is not supported") + + return result_data - Tensor._C.pow_tensor.argtypes = [ctypes.POINTER(CTensor), ctypes.c_float] - Tensor._C.pow_tensor.restype = ctypes.POINTER(CTensor) + def __truediv__(self, other): + other = float(other) + Tensor._C.tensor_div_scalar.argtypes = [ctypes.POINTER(CTensor), ctypes.c_float] + Tensor._C.tensor_div_scalar.restype = ctypes.POINTER(CTensor) - result_tensor_ptr = Tensor._C.pow_tensor(self.tensor, power_ctypes) + result_tensor_ptr = Tensor._C.tensor_div_scalar(other.tensor, ctypes.c_float(self)) + + result_data = Tensor() + result_data.tensor = result_tensor_ptr + result_data.shape = self.shape.copy() + result_data.ndim = self.ndim + result_data.device = self.device + + result_data.requires_grad = self.requires_grad + if result_data.requires_grad: + result_data.grad_fn = DivisionBackward(self, other) + + return result_data + + + + def __rtruediv__(self, other): + other = float(other) + + Tensor._C.scalar_div_tensor.argtypes = [ctypes.c_float, ctypes.POINTER(CTensor)] + Tensor._C.scalar_div_tensor.restype = ctypes.POINTER(CTensor) + + result_tensor_ptr = Tensor._C.scalar_div_tensor(ctypes.c_float(other), self.tensor) + + result_data = Tensor() + result_data.tensor = result_tensor_ptr + result_data.shape = self.shape.copy() + result_data.ndim = self.ndim + result_data.device = self.device + + result_data.requires_grad = self.requires_grad + if result_data.requires_grad: + result_data.grad_fn = DivisionBackward(other, self) + + return result_data + + + def log(self): + Tensor._C.log_tensor.argtypes = [ctypes.POINTER(CTensor)] + Tensor._C.log_tensor.restype = ctypes.POINTER(CTensor) + + result_tensor_ptr = Tensor._C.log_tensor(self.tensor) result_data = Tensor() result_data.tensor = result_tensor_ptr @@ -361,7 +440,7 @@ class Tensor: result_data.requires_grad = self.requires_grad if result_data.requires_grad: - result_data.grad_fn = PowBackward(self, power) + result_data.grad_fn = LogBackward(self) return result_data diff --git a/test.py b/test.py index 185fc29..b905f25 100644 --- a/test.py +++ b/test.py @@ -74,6 +74,9 @@ if __name__ == "__main__": [[13, 14], [15, 16], [17, 18]], [[19, 20], [21, 22], [23, 24]], [[25, 26], [27, 28], [29, 30]]], requires_grad=True) + + print(tensor1 / 1) + exit() # Reshape tensor1 to 2x3x5 reshaped_tensor = tensor1.transpose(1, 0)