batched matmul 2D and 3D tensors

This commit is contained in:
lucasdelimanogueira 2024-05-02 19:13:05 -03:00
parent a538ab4ecd
commit 31be051028
14 changed files with 134 additions and 18 deletions

Binary file not shown.

Binary file not shown.

Binary file not shown.

Binary file not shown.

Binary file not shown.

Binary file not shown.

View file

@ -45,6 +45,26 @@ void matmul_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data) {
}
}
void batched_matmul_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data) {
int tensor2_offset = tensor2->shape[1] * tensor2->shape[2];
int result_data_offset = tensor1->shape[0] * tensor2->shape[2];
for (int batch = 0; batch < tensor2->shape[0]; batch++) {
for (int i = 0; i < tensor1->shape[0]; i++) {
for (int j = 0; j < tensor2->shape[2]; j++) {
float sum = 0.0;
for (int k = 0; k < tensor1->shape[1]; k++) {
sum += tensor1->data[i * tensor1->shape[1] + k] * tensor2->data[batch*tensor2_offset + (k * tensor2->shape[2] + j)];
}
result_data[(batch * result_data_offset) + (i * tensor2->shape[2] + j)] = sum;
}
}
}
}
void pow_tensor_cpu(Tensor* tensor, float power, float* result_data) {
for (int i = 0; i < tensor->size; i++) {

View file

@ -8,6 +8,7 @@ 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 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_mul_tensor_cpu(Tensor* tensor, float scalar, float* result_data);
void ones_like_tensor_cpu(Tensor* tensor, float* result_data);

View file

@ -5,7 +5,6 @@
#define THREADS_PER_BLOCK 128
#define TILE_SIZE 32
#define SHMEM_SIZE THREADS_PER_BLOCK * sizeof(float)
__host__ void cpu_to_cuda(Tensor* tensor) {
@ -61,7 +60,7 @@ __host__ void add_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_da
__global__ void sum_tensor_cuda_kernel(float* data, float* result_data, int size) {
__shared__ float partial_sum[SHMEM_SIZE];
__shared__ float partial_sum[THREADS_PER_BLOCK * sizeof(float)];
int tid = threadIdx.x;
int i = blockIdx.x * blockDim.x + threadIdx.x;

View file

@ -327,9 +327,10 @@ extern "C" {
}
Tensor* matmul_tensor(Tensor* tensor1, Tensor* tensor2) {
//MxN @ NxP = MxP
// Check if tensors have compatible shapes for matrix multiplication
if (tensor1->shape[1] != tensor2->shape[0]) {
fprintf(stderr, "Incompatible shapes for matrix multiplication\n");
fprintf(stderr, "Incompatible shapes for matrix multiplication %dx%d and %dx%d\n", tensor1->shape[0], tensor1->shape[1], tensor2->shape[0], tensor2->shape[1]);
exit(1);
}
@ -387,6 +388,69 @@ extern "C" {
}
}
Tensor* batched_matmul_tensor(Tensor* tensor1, Tensor* tensor2) {
//MxN @ BATCHxNxP = BATCHxMxP
// Check if tensors have compatible shapes for matrix multiplication
if (tensor1->shape[1] != tensor2->shape[1]) {
fprintf(stderr, "Incompatible shapes for matrix multiplication %dx%d and %dx%d\n", tensor1->shape[0], tensor1->shape[1], tensor2->shape[0], tensor2->shape[1]);
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 = 3;
int* shape = (int*)malloc(ndim * sizeof(int));
if (shape == NULL) {
fprintf(stderr, "Memory allocation failed\n");
exit(1);
}
shape[0] = tensor2->shape[0];;
shape[1] = tensor1->shape[0];
shape[2] = tensor2->shape[2];
int size = 1;
for (int i = 0; i < ndim; i++) {
size *= shape[i];
}
float* result_data = (float*)malloc(size * sizeof(float));
if (result_data == NULL) {
fprintf(stderr, "Memory allocation failed\n");
exit(1);
}
if (strcmp(tensor1->device, "cuda") == 0) {
float* result_data;
cudaMalloc((void **)&result_data, size * sizeof(float));
matmul_tensor_cuda(tensor1, tensor2, result_data);
return create_tensor(result_data, shape, ndim, device);
}
else {
float* result_data = (float*)malloc(size * sizeof(float));
if (result_data == NULL) {
fprintf(stderr, "Memory allocation failed\n");
exit(1);
}
batched_matmul_tensor_cpu(tensor1, tensor2, result_data);
return create_tensor(result_data, shape, ndim, device);
}
}
Tensor* pow_tensor(Tensor* tensor, float power) {
char* device = (char*)malloc(strlen(tensor->device) + 1);
if (device != NULL) {

View file

@ -291,22 +291,38 @@ class Tensor:
return self
def __matmul__(self, other):
if self.ndim != 2 or other.ndim != 2:
raise ValueError("Matrix multiplication requires 2D tensors")
if other.ndim == 3:
#batched 3D matmul
if self.shape[1] != other.shape[0]:
raise ValueError("Incompatible shapes for matrix multiplication")
Tensor._C.batched_matmul_tensor.argtypes = [ctypes.POINTER(CTensor), ctypes.POINTER(CTensor)]
Tensor._C.batched_matmul_tensor.restype = ctypes.POINTER(CTensor)
result_tensor_ptr = Tensor._C.batched_matmul_tensor(self.tensor, other.tensor)
Tensor._C.matmul_tensor.argtypes = [ctypes.POINTER(CTensor), ctypes.POINTER(CTensor)]
Tensor._C.matmul_tensor.restype = ctypes.POINTER(CTensor)
result_data = Tensor()
result_data.tensor = result_tensor_ptr
result_data.shape = [other.shape[0], self.shape[0], other.shape[2]]
result_data.ndim = 3
result_data.device = self.device
else:
#2D matmul
if self.ndim != 2 or other.ndim != 2:
raise ValueError("Matrix multiplication requires 2D tensors")
result_tensor_ptr = Tensor._C.matmul_tensor(self.tensor, other.tensor)
if self.shape[1] != other.shape[0]:
raise ValueError("Incompatible shapes for matrix multiplication")
result_data = Tensor()
result_data.tensor = result_tensor_ptr
result_data.shape = [self.shape[0], other.shape[1]]
result_data.ndim = 2
result_data.device = self.device
Tensor._C.matmul_tensor.argtypes = [ctypes.POINTER(CTensor), ctypes.POINTER(CTensor)]
Tensor._C.matmul_tensor.restype = ctypes.POINTER(CTensor)
result_tensor_ptr = Tensor._C.matmul_tensor(self.tensor, other.tensor)
result_data = Tensor()
result_data.tensor = result_tensor_ptr
result_data.shape = [self.shape[0], other.shape[1]]
result_data.ndim = 2
result_data.device = self.device
result_data.requires_grad = self.requires_grad or other.requires_grad
if result_data.requires_grad:

22
test.py
View file

@ -20,8 +20,25 @@ if __name__ == "__main__":
import random
import numpy as np
a = norch.Tensor([
[[1.234, 2.123], [3.635, 4.456], [5.678, 6.789]],
[[7.890, 8.901], [9.012, 1.234], [2.345, 3.456]],
[[4.567, 5.678], [6.789, 7.890], [8.901, 9.012]],
[[1.234, 2.345], [3.456, 4.567], [5.678, 6.789]],
[[7.890, 8.901], [9.012, 1.234], [2.345, 3.456]]
])
b = norch.Tensor([
[1.234, 2.123, 1.5],
[5.678, 6.789, 1.293],
[3.635, 4.456, 1.0202],
[7.890, 8.901, 1.91],
])
result = b @ a
print(result)
#a = Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]])#.to("cuda")
#a = norch.Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]])#.to("cuda")
#b = Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]])#.to("cuda")
#c = Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]])#.to("cuda")
@ -37,7 +54,6 @@ if __name__ == "__main__":
print(a.grad)"""
"""#print(a)
"""
N = 10
a = norch.Tensor([[1 for _ in range(N)] for _ in range(N)])
#b = norch.Tensor([[random.uniform(0, 1) for _ in range(N)] for _ in range(N)])
@ -56,7 +72,7 @@ if __name__ == "__main__":
print("\n\n")
"""
"""
a = [[random.uniform(0, 1) for _ in range(N)] for _ in range(N)]
b = [[random.uniform(0, 1) for _ in range(N)] for _ in range(N)]
ini = time.time()