diff --git a/build/Makefile b/build/Makefile index 81c04e7..7e003df 100644 --- a/build/Makefile +++ b/build/Makefile @@ -3,7 +3,7 @@ CC = g++ NVCC = nvcc # Compiler flags -CFLAGS = -Wall -Wextra -std=c++11 -I/usr/local/cuda/include +CFLAGS = -Wall -Wextra -std=c++11 NVCCFLAGS = -std=c++11 # Directories @@ -19,7 +19,7 @@ CU_OBJS = $(patsubst $(SRCDIR)/%.cu, $(BUILDDIR)/%.cu.o, $(CU_SRCS)) # CUDA flags and libraries CUDAFLAGS = -arch=sm_75 -CUDALIBS = -L/usr/local/cuda -lcudart -lcuda +CUDALIBS = -I/usr/local/cuda/include -L/usr/local/cuda/lib64 -lcudart -lcuda # Rule to build the target $(TARGET): $(OBJS) $(CU_OBJS) @@ -27,7 +27,7 @@ $(TARGET): $(OBJS) $(CU_OBJS) # Rule to compile C++ source files $(BUILDDIR)/%.o: $(SRCDIR)/%.cpp - $(CC) $(CFLAGS) -fPIC -c $< -o $@ + $(CC) $(CFLAGS) -fPIC -c $< -o $@ $(CUDALIBS) # Rule to compile CUDA source files $(BUILDDIR)/%.cu.o: $(SRCDIR)/%.cu diff --git a/build/cpu.o b/build/cpu.o new file mode 100644 index 0000000..1910307 Binary files /dev/null and b/build/cpu.o differ diff --git a/build/cuda.cu.o b/build/cuda.cu.o index c04ec68..bd53e6b 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 69817f2..b90fbf6 100755 Binary files a/build/libtensor.so and b/build/libtensor.so differ diff --git a/build/tensor.o b/build/tensor.o index 3b6337d..1e213bd 100644 Binary files a/build/tensor.o and b/build/tensor.o differ diff --git a/compiled_example b/compiled_example new file mode 100755 index 0000000..af71cff Binary files /dev/null and b/compiled_example differ diff --git a/csrc/cpu.cpp b/csrc/cpu.cpp new file mode 100644 index 0000000..2392876 --- /dev/null +++ b/csrc/cpu.cpp @@ -0,0 +1,17 @@ +#include "tensor.h" +#include +#include +#include + +void add_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data) { + result_data = (float*)malloc(tensor1->size * sizeof(float)); + + if (result_data == NULL) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + + for (int i = 0; i < tensor1->size; i++) { + result_data[i] = tensor1->data[i] + tensor2->data[i]; + } +} diff --git a/csrc/cpu.h b/csrc/cpu.h new file mode 100644 index 0000000..d5ff611 --- /dev/null +++ b/csrc/cpu.h @@ -0,0 +1 @@ +void add_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data); \ No newline at end of file diff --git a/csrc/cuda.cu b/csrc/cuda.cu index dc64303..3bbdf5b 100644 --- a/csrc/cuda.cu +++ b/csrc/cuda.cu @@ -1,7 +1,57 @@ -#ifdef __CUDACC__ - __global__ void add_tensor_cuda(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]; +#include "tensor.h" +#include +#include + +#define THREADS_PER_BLOCK 128 + +__host__ void cpu_to_cuda(Tensor* tensor) { + float* data_tmp; + cudaMalloc((void **)&data_tmp, tensor->size * sizeof(float)); + cudaMemcpy(data_tmp, tensor->data, tensor->size * sizeof(float), cudaMemcpyHostToDevice); + + tensor->data = data_tmp; + + const char* device_str = "cuda"; + tensor->device = (char*)malloc(strlen(device_str) + 1); + strcpy(tensor->device, device_str); + + printf("Successfully sent tensor to: %s\n", tensor->device); +} + +__host__ void cuda_to_cpu(Tensor* tensor) { + float* data_tmp = (float*)malloc(tensor->size * sizeof(float)); + + cudaMemcpy(data_tmp, tensor->data, tensor->size * sizeof(float), cudaMemcpyDeviceToHost); + cudaFree(tensor->data); + + tensor->data = data_tmp; + + const char* device_str = "cpu"; + tensor->device = (char*)malloc(strlen(device_str) + 1); + strcpy(tensor->device, device_str); + + printf("Successfully sent tensor to: %s\n", tensor->device); +} + +__global__ void add_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]; } -#endif \ No newline at end of file +} + +__host__ void add_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_data) { + cudaMalloc((void **)&result_data, tensor1->size * sizeof(float)); + + int number_of_blocks = (tensor1->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; + add_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(); +} + diff --git a/csrc/cuda.h b/csrc/cuda.h index a747fcb..11cd5bf 100644 --- a/csrc/cuda.h +++ b/csrc/cuda.h @@ -1,8 +1,10 @@ #ifndef CUDA_KERNEL_H_ #define CUDA_KERNEL_H_ -#ifdef __CUDACC__ - __global__ void add_tensor_cuda(float* data1, float* data2, float* result_data, int size); -#endif + __host__ void cpu_to_cuda(Tensor* tensor); + __host__ void cuda_to_cpu(Tensor* tensor); + __global__ void add_tensor_cuda_kernel(float* data1, float* data2, float* result_data, int size); + __host__ void add_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_data); + #endif /* CUDA_KERNEL_H_ */ diff --git a/csrc/tensor.cpp b/csrc/tensor.cpp index ad3f4f1..a8407e3 100644 --- a/csrc/tensor.cpp +++ b/csrc/tensor.cpp @@ -5,7 +5,7 @@ #include #include "tensor.h" #include "cuda.h" -#define THREADS_PER_BLOCK 128 +#include "cpu.h" extern "C" { @@ -74,39 +74,18 @@ extern "C" { return tensor->data[index]; } - void to_device(Tensor* tensor, char* device) { - printf("Sending tensor to device: %s\n", device); - #ifdef __CUDACC__ + void to_device(Tensor* tensor, char* target_device) { + printf("Sending tensor to device: %s\n", target_device); - if ((strcmp(device, "cuda") == 0) && (strcmp(tensor->device, "cpu") == 0)) { - - float* data_tmp; + if ((strcmp(target_device, "cuda") == 0) && (strcmp(tensor->device, "cpu") == 0)) { + cpu_to_cuda(tensor); + } - cudaMalloc((void **)&data_tmp, tensor->size * sizeof(float)); - cudaMemcpy(data_tmp, tensor->data, tensor->size * sizeof(float), cudaMemcpyHostToDevice); - - free(tensor->data); - tensor->data = data_tmp; - tensor->device = device; - } - - else if ((strcmp(device, "cpu") == 0) && (strcmp(tensor->device, "cuda") == 0)) { - float* data_tmp = (float*)malloc(tensor->size * sizeof(float)); - - cudaMemcpy(data_tmp, tensor->data, tensor->size * sizeof(float), cudaMemcpyDeviceToHost); - cudaFree(tensor->data); - - tensor->data = data_tmp; - tensor->device = device; - } - #else - fprintf(stderr, "ERROR: CUDA is not available. Cannot perform GPU operations.\n"); - exit(0); - #endif + else if ((strcmp(target_device, "cpu") == 0) && (strcmp(tensor->device, "cuda") == 0)) { + cuda_to_cpu(tensor); + } } - - Tensor* add_tensor(Tensor* tensor1, Tensor* tensor2) { printf("Adding tensor\n"); if (tensor1->ndim != tensor2->ndim) { @@ -173,44 +152,21 @@ extern "C" { shape[i] = tensor1->shape[i]; } - #ifdef __CUDACC__ - if (strcmp(tensor1->device, "cuda") != 0) { + if (strcmp(tensor1->device, "cuda") != 0) { - float* result_data; - - cudaMalloc((void **)&result_data, tensor1->size * sizeof(float)); - - int number_of_blocks = (tensor1->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; - add_tensor_cuda<<>>(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(); - - return create_tensor(result_data, shape, ndim, device); - } - #endif - - float* result_data = (float*)malloc(tensor1->size * sizeof(float)); - if (result_data == NULL) { - fprintf(stderr, "Memory allocation failed\n"); - exit(1); - } - - for (int i = 0; i < tensor1->size; i++) { - result_data[i] = tensor1->data[i] + tensor2->data[i]; - } - - - return create_tensor(result_data, shape, ndim, device); + float* result_data; + add_tensor_cuda(tensor1, tensor2, result_data); + return create_tensor(result_data, shape, ndim, device); + } + else { + float* result_data; + add_tensor_cpu(tensor1, tensor2, result_data); + return create_tensor(result_data, shape, ndim, device); + } } Tensor* sub_tensor(Tensor* tensor1, Tensor* tensor2) { - printf("Adding tensor\n"); + printf("Subtracting tensor\n"); if (tensor1->ndim != tensor2->ndim) { fprintf(stderr, "Tensors must have the same number of dimensions %d and %d for subtraction\n", tensor1->ndim, tensor2->ndim); exit(1); @@ -290,7 +246,7 @@ extern "C" { } Tensor* elementwise_mul_tensor(Tensor* tensor1, Tensor* tensor2) { - printf("Adding tensor\n"); + printf("Elementwise multiplying tensor\n"); 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); diff --git a/example.cu b/example.cu new file mode 100644 index 0000000..d4e8a9f --- /dev/null +++ b/example.cu @@ -0,0 +1,54 @@ +#include + +// Kernel definition +__global__ void AddTwoVectors(float A[], float B[], float C[]) { + int i = threadIdx.x; + C[i] = A[i] + B[i]; +} + +int main() { + printf("kkkkkkkkkkkkkkkkkkkk"); + int N = 1000; // Size of the vectors + float A[N], B[N], C[N]; // Arrays for vectors A, B, and C + + // Initialize vectors A and B + for (int i = 0; i < N; ++i) { + A[i] = 1; + B[i] = 3; + } + + float *d_A, *d_B, *d_C; // Device pointers for vectors A, B, and C + + // Allocate memory on the device for vectors A, B, and C + cudaMalloc((void **)&d_A, N * sizeof(float)); + cudaMalloc((void **)&d_B, N * sizeof(float)); + cudaMalloc((void **)&d_C, N * sizeof(float)); + printf("\nkk2kkkkkkkkkkkkkkkkkk"); + + // Copy vectors A and B from host to device + cudaMemcpy(d_A, A, N * sizeof(float), cudaMemcpyHostToDevice); + cudaMemcpy(d_B, B, N * sizeof(float), cudaMemcpyHostToDevice); + printf("\nkk3kkkkkkkkkkkkkkkkkk"); + + // Kernel invocation with N threads + AddTwoVectors<<<1, N>>>(d_A, d_B, d_C); + printf("\nkk4kkkkkkkkkkkkkkkkkk"); + + // Check for error + cudaError_t error = cudaGetLastError(); + if(error != cudaSuccess) { + printf("CUDA error: %s\n", cudaGetErrorString(error)); + exit(-1); + } + + // Waits untill all CUDA threads are executed + cudaDeviceSynchronize(); + + // Copy vector C from device to host + cudaMemcpy(C, d_C, N * sizeof(float), cudaMemcpyDeviceToHost); + + // Free device memory + cudaFree(d_A); + cudaFree(d_B); + cudaFree(d_C); +} \ No newline at end of file diff --git a/norch/__pycache__/tensor.cpython-38.pyc b/norch/__pycache__/tensor.cpython-38.pyc index 84e4a97..8c72345 100644 Binary files a/norch/__pycache__/tensor.cpython-38.pyc and b/norch/__pycache__/tensor.cpython-38.pyc differ diff --git a/norch/tensor.py b/norch/tensor.py index 66a1a59..f7ae02c 100644 --- a/norch/tensor.py +++ b/norch/tensor.py @@ -78,7 +78,7 @@ class Tensor: def to(self, device): self.device = device - self.device_ctype = device.encode('utf-8') + self.device_ctype = self.device.encode('utf-8') Tensor._C.to_device.argtypes = [ctypes.POINTER(CTensor), ctypes.c_char_p] Tensor._C.to_device.restype = None @@ -234,8 +234,13 @@ if __name__ == "__main__": import numpy as np - a = Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]]).to("cuda") - print(a ** 2) + a = 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") + + d = b+c + print(d) + #print(a ** 2) """#print(a) N = 1000