diff --git a/build/cpu.o b/build/cpu.o index 6289686..9b48aff 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 6d729b8..3ddee86 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 2bb1564..dacb48e 100755 Binary files a/build/libtensor.so and b/build/libtensor.so differ diff --git a/build/tensor.o b/build/tensor.o index 49c51f6..549b305 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 343e53b..e230121 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 e5d728b..3ce2caf 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 3cdf8df..0761d0d 100644 --- a/norch/autograd/functions.py +++ b/norch/autograd/functions.py @@ -64,13 +64,12 @@ class PowBackward: class LogBackward: - def __init__(self, input_tensor): - self.input_tensor = input_tensor + def __init__(self, x): + self.input = [x] def backward(self, gradient): - input_tensor = self.input_tensor - grad_input = gradient / input_tensor + grad_input = gradient / self.input[0] return [grad_input] diff --git a/norch/csrc/cpu.cpp b/norch/csrc/cpu.cpp index 13a658f..ff5995d 100644 --- a/norch/csrc/cpu.cpp +++ b/norch/csrc/cpu.cpp @@ -47,6 +47,13 @@ void tensor_div_scalar_cpu(Tensor* tensor, float scalar, float* result_data) { } } +void tensor_div_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data) { + + for (int i = 0; i < tensor1->size; i++) { + result_data[i] = tensor1->data[i] / tensor2->data[i]; + } +} + void matmul_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data) { for (int i = 0; i < tensor1->shape[0]; i++) { diff --git a/norch/csrc/cpu.h b/norch/csrc/cpu.h index 1cc52fb..6e0e2a3 100644 --- a/norch/csrc/cpu.h +++ b/norch/csrc/cpu.h @@ -9,6 +9,7 @@ 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 tensor_div_tensor_cpu(Tensor* tensor1, Tensor* tensor2, 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); diff --git a/norch/csrc/cuda.cu b/norch/csrc/cuda.cu index 67dc8a0..73efdd2 100644 --- a/norch/csrc/cuda.cu +++ b/norch/csrc/cuda.cu @@ -217,6 +217,28 @@ __host__ void tensor_div_scalar_cuda(Tensor* tensor, float scalar, float* result cudaDeviceSynchronize(); } +__global__ void tensor_div_tensor_cuda_kernel(float* data1, float* data2, float* result_data, int size) { + + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < size) { + result_data[i] = data1[i] / data2[i]; + } +} + +__host__ void tensor_div_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_data) { + + int number_of_blocks = (tensor1->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; + tensor_div_tensor_cuda_kernel<<>>(tensor1->data, tensor2->data, result_data, tensor1->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; diff --git a/norch/csrc/cuda.h b/norch/csrc/cuda.h index 87ef884..d21beef 100644 --- a/norch/csrc/cuda.h +++ b/norch/csrc/cuda.h @@ -25,6 +25,9 @@ __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 tensor_div_tensor_cuda_kernel(float* data1, float* data2, float* result_data, int size); + __host__ void tensor_div_tensor_cuda(Tensor* tensor1, Tensor* tensor2, 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); diff --git a/norch/csrc/tensor.cpp b/norch/csrc/tensor.cpp index 05232d0..741662e 100644 --- a/norch/csrc/tensor.cpp +++ b/norch/csrc/tensor.cpp @@ -411,6 +411,57 @@ extern "C" { } } + Tensor* tensor_div_tensor(Tensor* tensor1, Tensor* tensor2) { + if (tensor1->ndim != tensor2->ndim) { + fprintf(stderr, "Tensors must have the same number of dimensions %d and %d for element-wise multiplication\n", tensor1->ndim, tensor2->ndim); + exit(1); + } + + if (strcmp(tensor1->device, tensor2->device) != 0) { + fprintf(stderr, "Tensors must be on the same device: %s and %s\n", tensor1->device, tensor2->device); + exit(1); + } + + char* device = (char*)malloc(strlen(tensor1->device) + 1); + if (device != NULL) { + strcpy(device, tensor1->device); + } else { + fprintf(stderr, "Memory allocation failed\n"); + exit(-1); + } + int ndim = tensor1->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++) { + if (tensor1->shape[i] != tensor2->shape[i]) { + fprintf(stderr, "Tensors must have the same shape %d and %d at index %d for subtraction\n", tensor1->shape[i], tensor2->shape[i], i); + exit(1); + } + shape[i] = tensor1->shape[i]; + } + + if (strcmp(tensor1->device, "cuda") == 0) { + + float* result_data; + cudaMalloc((void **)&result_data, tensor1->size * sizeof(float)); + tensor_div_tensor_cuda(tensor1, tensor2, result_data); + return create_tensor(result_data, shape, ndim, device); + } + else { + float* result_data = (float*)malloc(tensor1->size * sizeof(float)); + if (result_data == NULL) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + tensor_div_tensor_cpu(tensor1, tensor2, 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 diff --git a/norch/csrc/tensor.h b/norch/csrc/tensor.h index a9b8f8e..90b4d0d 100644 --- a/norch/csrc/tensor.h +++ b/norch/csrc/tensor.h @@ -22,6 +22,7 @@ extern "C" { 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* tensor_div_tensor(Tensor* tensor1, Tensor* tensor2); Tensor* reshape_tensor(Tensor* tensor, int* new_shape, int new_ndim); Tensor* matmul_tensor(Tensor* tensor1, Tensor* tensor2); Tensor* tensor_pow_scalar(Tensor* tensor, float exponent); diff --git a/norch/tensor.py b/norch/tensor.py index 31c5258..633b8ae 100644 --- a/norch/tensor.py +++ b/norch/tensor.py @@ -153,7 +153,8 @@ class Tensor: if self.grad_fn is not None: # not a leaf grads = self.grad_fn.backward(gradient) for tensor, grad in zip(self.grad_fn.input, grads): - tensor.backward(grad) + if isinstance(tensor, Tensor): + tensor.backward(grad) def zero_grad(self): self.grad = None @@ -200,6 +201,9 @@ class Tensor: return self.__str__() def __add__(self, other): + if isinstance(other, (int, float)): + other = other * self.ones_like() + if self.shape != other.shape: raise ValueError("Tensors must have the same shape for addition") @@ -221,6 +225,9 @@ class Tensor: return result_data def __sub__(self, other): + if isinstance(other, (int, float)): + other = other * self.ones_like() + if self.shape != other.shape: raise ValueError("Tensors must have the same shape for subtraction") @@ -345,51 +352,11 @@ class Tensor: return result_data 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) - - else: - raise ValueError("Tensor to tensor power is not supported") - - return result_data - - 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) + 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_div_scalar(other.tensor, ctypes.c_float(self)) + result_tensor_ptr = Tensor._C.tensor_pow_scalar(self.tensor, ctypes.c_float(other)) result_data = Tensor() result_data.tensor = result_tensor_ptr @@ -399,7 +366,62 @@ class Tensor: result_data.requires_grad = self.requires_grad if result_data.requires_grad: - result_data.grad_fn = DivisionBackward(self, other) + result_data.grad_fn = PowBackward(self, other) + + return result_data + + def __rpow__(self, other): + other = float(other) + 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(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 = PowBackward(other, self) + + return result_data + + def __truediv__(self, other): + if isinstance(other, (int, float)): + 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.tensor_div_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 = DivisionBackward(self, other) + + elif isinstance(self, Tensor) and isinstance(other, Tensor): + Tensor._C.tensor_div_tensor.argtypes = [ctypes.POINTER(CTensor), ctypes.POINTER(CTensor)] + Tensor._C.tensor_div_tensor.restype = ctypes.POINTER(CTensor) + + result_tensor_ptr = Tensor._C.tensor_div_tensor(self.tensor, other.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 or other.requires_grad + if result_data.requires_grad: + result_data.grad_fn = DivisionBackward(self, other) return result_data @@ -418,7 +440,7 @@ class Tensor: 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) diff --git a/test.py b/test.py index b905f25..e473a71 100644 --- a/test.py +++ b/test.py @@ -75,7 +75,10 @@ if __name__ == "__main__": [[19, 20], [21, 22], [23, 24]], [[25, 26], [27, 28], [29, 30]]], requires_grad=True) - print(tensor1 / 1) + result = 2 ** tensor1 + result = result.sum() + result.backward() + print(tensor1.grad) exit() # Reshape tensor1 to 2x3x5