In this blog, I will guide you through how to code the cuda kernel for 2D convolution.

Libs Required:

#include <stdio.h>
#include <cuda_runtime.h>

Kernel:

#define IS 5
#define KS 3

__global__
void convolution2DKernel(const float *input, const float *kernel, float *output,
                          int inputWidth, int inputHeight,
                          int kernelWidth, int kernelHeight) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;

    if (col < inputWidth && row < inputHeight) {
        int halfKernelWidth = kernelWidth / 2;
        int halfKernelHeight = kernelHeight / 2;

        float result = 0.0f;

        for (int i = 0; i < kernelHeight; ++i) {
            for (int j = 0; j < kernelWidth; ++j) {
                int inputRow = row - halfKernelHeight + i;
                int inputCol = col - halfKernelWidth + j;

                if (inputRow >= 0 && inputRow < inputHeight && inputCol >= 0 && inputCol < inputWidth) {
                    result += input[inputRow * inputWidth + inputCol] * kernel[i * kernelWidth + j];
                }
            }
        }

        output[row * inputWidth + col] = result;
    }
}

In this program, we have a kernel function called “convolution2DKernel”, which takes four arguments: two float arrays “input” and “kernal”, an float array “output”, and an integer “kernelSize”. The kernel function apply “kernel ” on “input ” array and stores the result in array “output”.

Running Kernel:

int main() {
    // Example input data
    const int inputWidth = IS;
    const int inputHeight = IS;
    const int kernelWidth = KS;
    const int kernelHeight = KS;

    float input[inputHeight * inputWidth];

    for (int i = 0; i < inputHeight; i++)
        for (int j = 0; j < inputWidth; j++)
            input[i * inputHeight + j] = i + j;

    float kernel[kernelHeight * kernelWidth] = {
        0.5, 1.0, 0.5,
        1.0, 2.0, 1.0,
        0.5, 1.0, 0.5
    };

    float output[inputHeight * inputWidth];


    printf("array:input \n");
    for (int i = 0; i < inputHeight ; i++) {
        for (int j = 0; j < inputWidth; j++)
            printf("%f ", input[i * inputHeight + j]);
        printf("\n");
    }

    printf("array:kernel \n");
    for (int i = 0; i < kernelHeight; i++) {
        for (int j = 0; j < kernelWidth; j++)
            printf("%f ", kernel[i * kernelHeight + j]);
        printf("\n");
    }

    float *d_input, *d_kernel, *d_output;
    cudaMalloc((void**)&d_input, inputHeight * inputWidth * sizeof(float));
    cudaMalloc((void**)&d_kernel, kernelHeight * kernelWidth * sizeof(float));
    cudaMalloc((void**)&d_output, inputHeight * inputWidth * sizeof(float));

    cudaMemcpy(d_input, input, inputHeight * inputWidth * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_kernel, kernel, kernelHeight * kernelWidth * sizeof(float), cudaMemcpyHostToDevice);

    dim3 blockSize(16, 16);
    dim3 gridSize((inputWidth + blockSize.x - 1) / blockSize.x, (inputHeight + blockSize.y - 1) / blockSize.y);

    
    convolution2DKernel<<<gridSize, blockSize>>>(d_input, d_kernel, d_output,
                                                 inputWidth, inputHeight,
                                                 kernelWidth, kernelHeight);

    
    cudaMemcpy(output, d_output, inputHeight * inputWidth * sizeof(float), cudaMemcpyDeviceToHost);

    cudaFree(d_input);
    cudaFree(d_kernel);
    cudaFree(d_output);

    printf("array:output \n");

    for (int i = 0; i < inputHeight; i++) {
        for (int j = 0; j < inputWidth; j++)
            printf("%f ", output[i * inputHeight + j]);
        printf("\n");
    }

    return 0;
}

In the main function, we declare and initialize the necessary variables. We allocate memory on both the host and the device using the “malloc” and “cudaMalloc” functions, respectively. We then initialize the matrix “input” and “kernel” with values, and copy them to the device using the “cudaMemcpy” function. We launch the kernel function using the “<<< >>>” syntax, and then copy the result from the device to the host using “cudaMemcpy” again. Finally, we print the result of the vector output and free the memory.

Compile and Running:

To compile the program, we need to use the “nvcc” compiler provided by the CUDA Toolkit. We can compile the program with the following command:

nvcc 2d_convolution_code.cu -o 2d_convolution_code

To run the program, we simply execute the binary file generated by the compiler:

./2d_convolution_code

Conclusion:

I hope this blog has given you a good introduction to CUDA programming with C, and that you’re excited to explore more advanced topics in CUDA programming. Happy coding!

In this blog, I will guide you through how to code the cuda kernel for 1D convolution.

Libs Required:

#include <stdio.h>
#include <cuda_runtime.h>

Kernel:

#define KS 3
#define IS 10

__global__
void convolutionKernel(const float *input, const float *kernel, float *output, int inputSize, int kernelSize) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;

    if (tid < inputSize) {
        int halfKernelSize = kernelSize / 2;
        float result = 0.0f;

        for (int i = 0; i < kernelSize; ++i) {
            int inputIndex = tid - halfKernelSize + i;
            if (inputIndex >= 0 && inputIndex < inputSize) {
                result += input[inputIndex] * kernel[i];
            }
        }

        output[tid] = result;
    }
}

In this program, we have a kernel function called “convolutionKernel”, which takes four arguments: two float arrays “input” and “kernal”, an float array “output”, and an integer “kernelSize”. The kernel function apply “kernel ” on “input ” array and stores the result in array “output”.

Running Kernel:

int main() {
    const int inputSize = IS;
    const int kernelSize = KS;

    float input[inputSize];
    float kernel[kernelSize] = {0.5, 1.0, 0.5};
    float output[inputSize];

    for (int i = 0; i < inputSize; i++)
        input[i] = i;

    printf("input: ");
    for (int i = 0; i < inputSize; i++)
        printf("%f ", input[i]);
    printf("\n");

    printf("kernel: ");
    for (int i = 0; i < kernelSize; i++)
        printf("%f ", kernel[i]);
    printf("\n");

    float *d_input, *d_kernel, *d_output;
    cudaMalloc((void**)&d_input, inputSize * sizeof(float));
    cudaMalloc((void**)&d_kernel, kernelSize * sizeof(float));
    cudaMalloc((void**)&d_output, inputSize * sizeof(float));

    cudaMemcpy(d_input, input, inputSize * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_kernel, kernel, kernelSize * sizeof(float), cudaMemcpyHostToDevice);

    int blockSize = 256;
    int gridSize = (inputSize + blockSize - 1) / blockSize;

    convolutionKernel<<<gridSize, blockSize>>>(d_input, d_kernel, d_output, inputSize, kernelSize);

    cudaMemcpy(output, d_output, inputSize * sizeof(float), cudaMemcpyDeviceToHost);

    cudaFree(d_input);
    cudaFree(d_kernel);
    cudaFree(d_output);

    printf("output: ");
    for (int i = 0; i < inputSize; ++i) {
        printf("%f ", output[i]);
    }
    printf("\n");

    return 0;
}

In the main function, we declare and initialize the necessary variables. We allocate memory on both the host and the device using the “malloc” and “cudaMalloc” functions, respectively. We then initialize the matrix “input” and “kernel” with values, and copy them to the device using the “cudaMemcpy” function. We launch the kernel function using the “<<< >>>” syntax, and then copy the result from the device to the host using “cudaMemcpy” again. Finally, we print the result of the vector output and free the memory.

Compile and Running:

To compile the program, we need to use the “nvcc” compiler provided by the CUDA Toolkit. We can compile the program with the following command:

nvcc 1d_convolution_code.cu -o 1d_convolution_code

To run the program, we simply execute the binary file generated by the compiler:

./1d_convolution_code

Conclusion:

I hope this blog has given you a good introduction to CUDA programming with C, and that you’re excited to explore more advanced topics in CUDA programming. Happy coding!

In this blog, I will guide you through how to code the cuda kernel for MNIST MLP inference.

Introduction:

In the ever-evolving landscape of artificial intelligence and deep learning, neural networks have emerged as a dominant force, propelling groundbreaking advancements across various domains. Among these neural network architectures, the Multilayer Perceptron (MLP) stands out as one of the most fundamental and widely-used models for its simplicity and effectiveness in solving various tasks. One such task is digit recognition, where the classic MNIST dataset has become the de facto benchmark for evaluating the performance of machine learning algorithms.

In this blog post, we embark on an exhilarating journey into the realm of high-performance computing and CUDA programming to unleash the full potential of custom MNIST MLP inference. NVIDIA’s CUDA (Compute Unified Device Architecture) has revolutionized the way we harness the power of Graphics Processing Units (GPUs) to accelerate computationally intensive tasks. By leveraging the parallel processing capabilities of GPUs, we can achieve significant speedups in training and inference times, making them indispensable tools for modern deep learning applications.

Our objective throughout this article is to demystify the complexities of CUDA programming and empower you to build a custom MLP inference engine tailored to the MNIST dataset. Whether you are a seasoned deep learning practitioner or a curious enthusiast eager to dive into the world of GPU acceleration, this guide will equip you with the essential knowledge to wield CUDA effectively.

The knowledge gained from this blog will not only empower you to optimize MNIST digit recognition but also serve as a stepping stone to explore more complex deep learning tasks that demand high computational power.

So, buckle up and get ready to harness the parallel processing capabilities of CUDA as we embark on an exhilarating journey to supercharge our custom MNIST MLP inference engine!

Training the Model and Saving Weights in NumPy:

In this section of the blog, we covered the initial steps of implementing custom MNIST MLP inference using CUDA. We began by setting up the MNIST dataset and then designing the architecture of our Multilayer Perceptron (MLP) model and training it using PyTorch.

To preserve the trained model’s learned parameters for future use and compatibility with other environments, we saved the optimized weights and biases in NumPy format.

Data reading and Model training:

import torch
import torch.nn as nn
import torch.nn.functional as F
from torchvision import datasets, transforms

import numpy as np
import matplotlib.pyplot as plt

import os


if torch.cuda.is_available():
    device = torch.device("cuda")
else:
    device = torch.device("cpu")

print("Using PyTorch version:", torch.__version__, " Device:", device)

batch_size = 32

train_dataset = datasets.MNIST(
    "./data", train=True, download=True, transform=transforms.ToTensor()
)

validation_dataset = datasets.MNIST(
    "./data", train=False, transform=transforms.ToTensor()
)

train_loader = torch.utils.data.DataLoader(
    dataset=train_dataset, batch_size=batch_size, shuffle=True
)

validation_loader = torch.utils.data.DataLoader(
    dataset=validation_dataset, batch_size=batch_size, shuffle=False
)

for X_train, y_train in train_loader:
    print("X_train:", X_train.size(), "type:", X_train.type())
    print("y_train:", y_train.size(), "type:", y_train.type())
    break


class Net(nn.Module):
    def __init__(self):
        super(Net, self).__init__()
        self.fc1 = nn.Linear(28 * 28, 50)
        self.fc1_drop = nn.Dropout(0.2)
        self.fc2 = nn.Linear(50, 50)
        self.fc2_drop = nn.Dropout(0.2)
        self.fc3 = nn.Linear(50, 10)

    def forward(self, x):
        x = x.view(-1, 28 * 28)
        x = F.relu(self.fc1(x))
        x = self.fc1_drop(x)
        x = F.relu(self.fc2(x))
        x = self.fc2_drop(x)
        return F.log_softmax(self.fc3(x), dim=1)


model = Net().to(device)
optimizer = torch.optim.SGD(model.parameters(), lr=0.01, momentum=0.5)
criterion = nn.CrossEntropyLoss()

print(model)


def train(epoch, log_interval=200):
    # Set model to training mode
    model.train()

    # Loop over each batch from the training set
    for batch_idx, (data, target) in enumerate(train_loader):
        data = data.to(device)
        target = target.to(device)
        optimizer.zero_grad()
        output = model(data)
        loss = criterion(output, target)
        loss.backward()
        optimizer.step()

        if batch_idx % log_interval == 0:
            print(
                "Train Epoch: {} [{}/{} ({:.0f}%)]\tLoss: {:.6f}".format(
                    epoch,
                    batch_idx * len(data),
                    len(train_loader.dataset),
                    100.0 * batch_idx / len(train_loader),
                    loss.data.item(),
                )
            )

def validate(loss_vector, accuracy_vector):
    model.eval()
    val_loss, correct = 0, 0
    for data, target in validation_loader:
        data = data.to(device)
        target = target.to(device)
        output = model(data)
        val_loss += criterion(output, target).data.item()
        pred = output.data.max(1)[1]  # get the index of the max log-probability
        correct += pred.eq(target.data).cpu().sum()

    val_loss /= len(validation_loader)
    loss_vector.append(val_loss)

    accuracy = 100.0 * correct.to(torch.float32) / len(validation_loader.dataset)
    accuracy_vector.append(accuracy)

    print(
        "\nValidation set: Average loss: {:.4f}, Accuracy: {}/{} ({:.0f}%)\n".format(
            val_loss, correct, len(validation_loader.dataset), accuracy
        )
    )

epochs = 5

lossv, accv = [], []
for epoch in range(1, epochs + 1):
    train(epoch)
    validate(lossv, accv)

Weights preserving:

weights = [
    (name, param.detach().cpu().numpy()) for name, param in model.named_parameters()
]

for w in weights:
    with open("./model/" + str(w[0]) + ".npy", "wb") as f:
        print("shape")

        if "bias" in w[0]:
            weights = np.expand_dims(w[1], axis=(0))
            print(weights.shape)

            weights = np.ascontiguousarray(weights)
            np.save(f, weights)
        else:
            weights = w[1].T
            print(weights.shape)

            weights = np.ascontiguousarray(weights)
            np.save(f, weights)

Loading NumPy Weights in CUDA in C++:

In this section we will read the numpy weight in C++ using python interpeter. As we ready using Python interpeter we get PyArrayObject , which we will convert it into float. And we will move the float array to GPU using CUDA.

#include <stdio.h>
#include <stdlib.h>
#include <Python.h>
#include <cuda_runtime.h>
#include <numpy/arrayobject.h>

#define NPY_NO_DEPRECATED_API NPY_1_7_API_VERSION

struct Weights {
    float* matrix;
    int ndims;
    int *shape;
    long int size;
};

PyObject* read_numpy_file(const char* file_path) {
    // Import the Python module containing the function
    PyObject* numpy_module = PyImport_ImportModule("numpy");
    if (numpy_module == nullptr) {
        PyErr_Print();
        return nullptr;
    }

    // Get the reference to the function
    PyObject* numpy_function = PyObject_GetAttrString(numpy_module, "load");
    if (numpy_function == nullptr) {
        PyErr_Print();
        Py_DECREF(numpy_module);
        return nullptr;
    }

    // Create the arguments tuple
    PyObject* args = PyTuple_New(1);
    PyTuple_SetItem(args, 0, PyUnicode_FromString(file_path));

    // Call the Python function with the arguments
    PyObject* result = PyObject_CallObject(numpy_function, args);
    if (result == nullptr) {
        PyErr_Print();
        Py_DECREF(numpy_module);
        Py_DECREF(numpy_function);
        Py_DECREF(args);
        return nullptr;
    }

    // Print the shape of the NumPy array
    PyObject* shape = PyObject_GetAttrString(result, "shape");
    if (shape != nullptr) {
        PyObject* repr = PyObject_Repr(shape);
        const char* str = PyUnicode_AsUTF8(repr);
        printf("Shape: %s\n", str);
        Py_DECREF(repr);
        Py_DECREF(shape);
    } else {
        printf("Failed to get shape.\n");
    }

    // Clean up references
    Py_DECREF(numpy_module);
    Py_DECREF(numpy_function);
    Py_DECREF(args);
    return result;
}

PyArrayObject* read_weights_from_numpy(const char* file_path, int print) {
    // Call the Python function and get the PyObject reference to the NumPy array
    PyObject* numpy_array = read_numpy_file(file_path);
    if (numpy_array == nullptr) {
        // Handle error
        return nullptr;
    }

    // Use the PyObject reference to the NumPy array as needed
    // Example: Print the array
    if (print == 1){
        PyObject* repr = PyObject_Repr(numpy_array);
        const char* str = PyUnicode_AsUTF8(repr);
        printf("%s\n", str);
    }

    // Convert the result to a NumPy array
    PyArrayObject* array = reinterpret_cast<PyArrayObject*>(numpy_array);

    return array;
}

int get_numpy_ndims(PyArrayObject* array){
    // Get number of dimensions
    int ndim = PyArray_NDIM(array);
    printf("%d \n", ndim);
    return ndim;
}

long int get_numpy_size(PyArrayObject* array){
    // Get the total size of the array
    npy_intp total_size_intp = PyArray_SIZE(array);
    printf("%" NPY_INTP_FMT "\n", total_size_intp);
    
    long int total_size = static_cast<long int>(total_size_intp);
    return total_size;
}

float* convert_PyArrayObject_to_float(PyArrayObject* array, int print, int *shape, int ndim) {

    printf("values %d %d \n", PyArray_TYPE(array), NPY_DOUBLE);
    // Check the data type of the numpy array
    if (PyArray_TYPE(array) != NPY_FLOAT32) {
        printf("Input numpy array is not of type float.\n");
    }

    // Convert the weight into float
    float* matrix = static_cast<float*>(PyArray_DATA(array));

    // Printing weights for checking it after conversion 
    if (print == 1){
        if (ndim == 2){
            for (int i = 0; i < shape[0]; ++i) {
                if (i ==0){
                    for (int j = 0; j < shape[1]; ++j) {
                        // matrix[j+(i*shape[1])] = 0.5f;
                        if(j<3 || j>shape[1]-3){
                            printf("%d index %.5f ", j+(i*shape[1]), matrix[j+(i*shape[1])]);
                        }
                    }
                    printf("\n");
                }
            }
        }
        else if (ndim == 1){
            for (int i = 0; i < shape[0]; ++i) {
                // matrix[i] = 0.5f;
                if(i<5 || i>shape[0]-5){
                    printf("%d index %.5f ", i, matrix[i]);
                }
                
            }
            printf("\n");
        }
       
    }

    return matrix; 
}


float* move_weight_to_cuda(float* weights ,long int total_size){
    // Allocate CUDA device memory
    float* d_data;
    printf("total size %d \n", total_size);

    cudaError_t cudaStatus;
    cudaStatus = cudaMalloc((void**)&d_data, total_size * sizeof(float));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed: %s\n", cudaGetErrorString(cudaStatus));
        // Handle the error or return an error code
    }
    // Copy the array data from host (CPU) to device (CUDA)
    cudaStatus = cudaMemcpy(d_data, weights, total_size * sizeof(float), cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed: %s\n", cudaGetErrorString(cudaStatus));
        // Handle the error or return an error code
    }
    return d_data;
}

Weights read_weights(const char* file_path, int print){
    Weights weight;

    // reading numpy weights
    PyArrayObject* array = read_weights_from_numpy(file_path, print);
    if (array == nullptr) {
        // Handle error
        weight.matrix = nullptr;
        return weight;
    }

    int ndims = get_numpy_ndims(array);
    get_numpy_shape(array, weight, ndims);
    
    long int size = get_numpy_size(array);
    float* matrix = convert_PyArrayObject_to_float(array, print, weight.shape, ndims);
    
    // Release the PyObject reference

    float* cuda_weights = move_weight_to_cuda(matrix, size);
    Py_DECREF(array);

    weight.ndims = ndims;
    weight.size = size;
    weight.matrix = cuda_weights;

    return weight;
}

Defining CUDA kernels for MLP in C++:

We have to define 3 kernel for make MLP to work matrixMulKernel, matrixAddKernel and softmaxKernel. For more details please check links in references.

__global__ void matrixMulKernel(float* matrixA, float* matrixB, float* matrixC, int rowsA, int colsA, int colsB) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    if (row < rowsA && col < colsB) {
        float sum = 0.0f;
        for (int k = 0; k < colsA; k++) {
            sum += matrixA[row * colsA + k] * matrixB[k * colsB + col];
        }
        matrixC[row * colsB + col] = sum;

        // Print intermediate result
        // printf("Intermediate result at rowsA %d, colsB %d  [%d][%d]: %.5f\n", rowsA, colsB, row, col, sum);
    }
}


__global__ void matrixAddKernel(float* matrixA, float* matrixB, float* matrixC, int rows, int cols) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    if (row < rows && col < cols) {
        int index = row * cols + col;
        matrixC[index] = matrixA[index] + matrixB[index];
    }
}

__global__ void softmaxKernel(float* input, float* output, int rows, int cols) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    if (row < rows && col < cols) {
        int index = row * cols + col;

        // Compute the exponential of each element
        float expVal = expf(input[index]);

        // Compute the sum of exponentials for the row
        float sumExp = 0.0f;
        for (int i = 0; i < cols; ++i) {
            sumExp += expf(input[row * cols + i]);
        }

        // Compute the softmax value for the element
        output[index] = expVal / sumExp;
    }
}


float* matrixMul(float* matrixA, float* matrixB, int rowsA, int colsA, int rowsB, int colsB){
    float* matrixC;

    printf(" ROW A %d  COLS A %d \n",rowsA, colsA);
    printf(" ROW B %d  COLS B %d \n",rowsB, colsB);

    cudaMalloc((void **)&matrixC, rowsA * colsB * sizeof(float));

    dim3 blockSize(16, 16);
    dim3 gridSize((colsB + blockSize.x - 1) / blockSize.x, (rowsA + blockSize.y - 1) / blockSize.y);

    matrixMulKernel<<<gridSize, blockSize>>>(matrixA, matrixB, matrixC, rowsA, colsA, colsB);

    return matrixC;
}

float* matrixAdd(float* matrixA, float* matrixB, int rows, int cols){
    float* matrixC;

    printf(" ROW %d  COLS %d \n",rows, cols);

    cudaMalloc((void **)&matrixC, rows * cols * sizeof(float));

    dim3 blockSize(16, 16);
    dim3 gridSize((cols + blockSize.x - 1) / blockSize.x, (rows + blockSize.y - 1) / blockSize.y);

    matrixAddKernel<<<gridSize, blockSize>>>(matrixA, matrixB, matrixC, rows, cols);

    return matrixC;
}


float* softmax(float* input, int rows, int cols){
    float* matrixC;

    printf(" ROW %d  COLS %d \n",rows, cols);

    cudaMalloc((void **)&matrixC, rows * cols * sizeof(float));

    dim3 blockSize(16, 16);
    dim3 gridSize((cols + blockSize.x - 1) / blockSize.x, (rows + blockSize.y - 1) / blockSize.y);

    softmaxKernel<<<gridSize, blockSize>>>(input, matrixC, rows, cols);

    return matrixC;
}

Defining main function for inference in C++:

In this section, we will combine all the reading and kernel defination for MLP inference on CUDA.

int main() {
    // Initialize the Python interpreter
    Py_Initialize();

    // Ensure that NumPy is available
    import_array();
    Weights fc1_w = read_weights("./mlp/mnist_mlp/model/fc1.weight.npy", 0);
    Weights fc2_w = read_weights("./mlp/mnist_mlp/model/fc2.weight.npy", 0);
    Weights fc3_w = read_weights("./mlp/mnist_mlp/model/fc3.weight.npy", 0);
    printf("#############################\n");
    Weights fc1_b = read_weights("./mlp/mnist_mlp/model/fc1.bias.npy", 0);
    Weights fc2_b = read_weights("./mlp/mnist_mlp/model/fc2.bias.npy", 0);
    Weights fc3_b = read_weights("./mlp/mnist_mlp/model/fc3.bias.npy", 0);

    //Read image
    Inputs image = read_image("./mlp/mnist_mlp/images/1.npy", 0);
    // Inputs image = read_image("./mlp/mnist_mlp/2.npy", 0);

    // Finalize the Python interpreter
    Py_Finalize();

    int output_row = 1;
    int output_col = 10;

    float* matrixC = nullptr;
    
    matrixC = matrixMul(image.matrix, fc1_w.matrix, image.shape[0], image.shape[1], fc1_w.shape[0], fc1_w.shape[1]);
    matrixC = matrixAdd(matrixC, fc1_b.matrix, fc1_b.shape[0], fc1_b.shape[1]);

    matrixC = matrixMul(matrixC, fc2_w.matrix, fc1_b.shape[0], fc1_b.shape[1], fc2_w.shape[0], fc2_w.shape[1]);
    matrixC = matrixAdd(matrixC, fc2_b.matrix, fc2_b.shape[0], fc2_b.shape[1]);

    matrixC = matrixMul(matrixC, fc3_w.matrix, fc2_b.shape[0], fc2_b.shape[1], fc3_w.shape[0], fc3_w.shape[1]);
    matrixC = matrixAdd(matrixC, fc3_b.matrix, fc3_b.shape[0], fc3_b.shape[1]);

    matrixC = softmax(matrixC, fc3_b.shape[0], fc3_b.shape[1]);

    printf("output shape 1: %d, %d", fc3_b.shape[0], fc3_b.shape[1]);

    float* C = (float *)malloc(output_row * output_col * sizeof(float));

    cudaMemcpy(C, matrixC,  output_row * output_col * sizeof(float), cudaMemcpyDeviceToHost);
    cudaDeviceSynchronize();

    printf("rowsA %d\n", output_row);
    printf("colsB %d\n", output_col);
    for (int i = 0; i < output_row; i++) {
        for (int j = 0; j < output_col; j++){
            printf("%f ", C[i * output_row + j]);
            
        }
        printf("\n");
    }


    sleep(5);
 
    // Clean up CUDA device memory
    cudaFree(fc1_w.matrix);
    cudaFree(fc2_w.matrix);
    cudaFree(fc3_w.matrix);
    
    cudaFree(fc1_b.matrix);
    cudaFree(fc2_b.matrix);
    cudaFree(fc3_b.matrix);

    return 0;
}

Compile and Running:

To compile the program, we need to use the “nvcc” compiler provided by the CUDA Toolkit. We can compile the program with the following command:

nvcc -o ml.o ml.cu -I/usr/include/python3.8/ -I/user/lib/python3/dist-packages/numpy/ -lpython3.8

To run the program, we simply execute the binary file generated by the compiler:

./ml.o

Conclusion:

I hope this blog has given you a good introduction to CUDA programming with C, and that you’re excited to explore more advanced topics in CUDA programming. Happy coding!

References:

1. https://github.com/harsh2011/cuda_learning (Github Repo)
2. https://medium.com/p/4d1f0b8535a3 (CUDA Programming: Array Addition)
3. https://medium.com/p/1aa774b8d703 (CUDA Programming: 2D Matrix Multiplication)

The evolution of computer architecture has witnessed a shift from single-core to multi-core processing. CPUs have been developed to increase their core counts, but there is a limit to the number of cores that can be added to a CPU. As a result, researchers have turned their attention to another type of processor: the GPU (Graphics Processing Unit).

GPUs are highly parallel and have many more cores than CPUs. GPUs were first made for processing graphics, but they have become popular in scientific and technical computing because they can do a lot of calculations at once. NVIDIA made CUDA (Compute Unified Device Architecture) in 2006 so that GPUs could be used to their fullest potential. CUDA is a platform and programming model for parallel computing that lets developers use the power of graphics processing units (GPUs) for general-purpose computing.

In this series of blogs, we will cover the basics of CUDA programming, starting with the installation of the CUDA toolkit and moving on to the development of a simple CUDA program.

Installing the CUDA Toolkit

Before we can start programming in CUDA, we need to install the CUDA toolkit. The CUDA toolkit includes the necessary compilers, libraries, and development tools for CUDA programming.

To install the CUDA toolkit, follow these steps:

1. Visit the NVIDIA CUDA website at https://developer.nvidia.com/cuda-toolkit.
2. Select the appropriate version of the CUDA toolkit for your system.
3. Download and run the installer.
4. Follow the instructions provided by the installer.
5. Once the installation is complete, restart your computer.
6. Verify that the CUDA toolkit is installed correctly by running the following command in a terminal window:

nvcc --version

Steps of Cuda Programming

To begin, we must understand the structure of a cuda-based code; there are a few simple steps to take.

• Data initialization on the CPU
• Data should be transferred from CPU to GPU.
• Kernel launch (instructions on GPU)
• Bring the results from GPU back to CPU.
• Reclaim the memory from both CPU and GPU

In such an environment, we will call host code the code that is going to run on the CPU, and device code the code that is going to run on the GPU.

CUDA Kernels

We now understand how the fundamentals of cuda work. I'll go over some basic kernel programming examples to help you learn and understand how the kernel works.

1. Array Addition
2. 2D Matrix Multiplication

Conclusion:

I hope this blog has given you a good introduction to CUDA programming with C, and that you’re excited to explore more advanced topics in CUDA programming. Happy coding!

In this blog, I will guide you through how to code the cuda kernel for 2D matrix multiplication.

Libs Required:

#include <stdio.h>
#include <cuda_runtime.h>

Kernel:

#define N 3

__global__ void matrix_mul(int *a, int *b, int *c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    int sum = 0;
    if (i < n && j < n) {
        for (int k = 0; k < n; k++)
            sum += a[i * n + k] * b[k * n + j];
        c[i * n + j] = sum;
    }
}

In this program, we have a kernel function called “add”, which takes four arguments: two integer arrays “a” and “b”, an integer array “c”, and an integer “n”. The kernel function adds the corresponding elements of arrays “a” and “b” and stores the result in array “c”. The if statement ensures that only valid elements of arrays “a”, “b”, and “c” are accessed.

Running Kernel:

int main() {
    int n = N;
    int *a, *b, *c;
    int *d_a, *d_b, *d_c;
    int size = n * n * sizeof(int);

    a = (int *)malloc(size);
    b = (int *)malloc(size);
    c = (int *)malloc(size);

    for (int i = 0; i < n; i++)
        for (int j = 0; j < n; j++) {
            a[i * n + j] = i + j;
            b[i * n + j] = i * j;
        }

    cudaMalloc((void **)&d_a, size);
    cudaMalloc((void **)&d_b, size);
    cudaMalloc((void **)&d_c, size);

    cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice);

    dim3 blockSize(N, N);
    dim3 gridSize((n + N - 1) / N, (n + N - 1) / N);
    matrix_mul<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);

    cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost);

    for (int i = 0; i < n; i++) {
        for (int j = 0; j < n; j++)
            printf("%d ", c[i * n + j]);
        printf("\n");
    }

    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(a);
    free(b);
    free(c);

    return 0;
}

In the main function, we declare and initialize the necessary variables. We allocate memory on both the host and the device using the “malloc” and “cudaMalloc” functions, respectively. We then initialize the matrix “a” and “b” with values, and copy them to the device using the “cudaMemcpy” function. We launch the kernel function using the “<<< >>>” syntax, and then copy the result from the device to the host using “cudaMemcpy” again. Finally, we print the result of the vector addition and free the memory.

Compile and Running:

To compile the program, we need to use the “nvcc” compiler provided by the CUDA Toolkit. We can compile the program with the following command:

nvcc matrix_multiplication.cu -o matrix_multiplication

To run the program, we simply execute the binary file generated by the compiler:

./matrix_multiplication

Conclusion:

I hope this blog has given you a good introduction to CUDA programming with C, and that you’re excited to explore more advanced topics in CUDA programming. Happy coding!

In this blog, I will guide you through how to code the cuda kernel for array addition.

Libs Reguired:

#include <stdio.h>
#include <cuda_runtime.h>

Kernel:

__global__ void add(int *a, int *b, int *c, int n)
{
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < n) 
        c[i] = a[i] + b[i];
}

In this program, we have a kernel function called “add”, which takes four arguments: two integer arrays “a” and “b”, an integer array “c”, and an integer “n”. The kernel function adds the corresponding elements of arrays “a” and “b” and stores the result in array “c”. The if statement ensures that only valid elements of arrays “a”, “b”, and “c” are accessed.

Running Kernel:

int main() {
    int n = 10;
    int *a, *b, *c;
    int *d_a, *d_b, *d_c;
    int size = n * sizeof(int);

    a = (int *)malloc(size);
    b = (int *)malloc(size);
    c = (int *)malloc(size);

    for (int i = 0; i < n; i++) {
        a[i] = i;
        b[i] = i * 2;
    }

    cudaMalloc((void **)&d_a, size);
    cudaMalloc((void **)&d_b, size);
    cudaMalloc((void **)&d_c, size);

    cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice);

    add<<<1, n>>>(d_a, d_b, d_c, n);

    cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost);

    for (int i = 0; i < n; i++)
        printf("%d + %d = %d\n", a[i], b[i], c[i]);

    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(a);
    free(b);
    free(c);

    return 0;
}

In the main function, we declare and initialize the necessary variables. We allocate memory on both the host and the device using the “malloc” and “cudaMalloc” functions, respectively. We then initialize the arrays “a” and “b” with values, and copy them to the device using the “cudaMemcpy” function. We launch the kernel function using the “<<< >>>” syntax, and then copy the result from the device to the host using “cudaMemcpy” again. Finally, we print the result of the vector addition and free the memory.

Compile and Running:

To compile the program, we need to use the “nvcc” compiler provided by the CUDA Toolkit. We can compile the program with the following command:

nvcc vector_addition.cu -o vector_addition

To run the program, we simply execute the binary file generated by the compiler:

./vector_addition

Conclusion:

I hope this blog has given you a good introduction to CUDA programming with C, and that you’re excited to explore more advanced topics in CUDA programming. Happy coding!

Generative Adversarial Network is a deep network that works for generating images from random data. It comprises two different networks: the generator and the discriminator. The discriminator network is a standard convolutional neural network that can classify the images fed to it, similar to how a binomial classifier classifies images as real or fake. The generator is an inverse convolutional network that upsamples a vector of random noise to produce an image.

As a discriminator, you assist the generator in creating images that resemble real ones. It aids in the training of the GAN because it has a min-max loss function. Whereas the Generator minimizes log(1- D(G(z))) (z is a random data vector), the Discriminator maximizes the output log(D(x). (x is a real image which you can call as ground truth to generate an image)

Training a GAN is a difficult task because both networks are fighting to decrease its loss. They are inversely proportional; if one decreases, the other increases. We have to train GAN properly, giving them equal iterations to train. Also, check the generated image's output to see if it's correct.

What is the use of GAN in tictactoe.ai?

Thinking of taking this game to a single player, there should be someone who draws for the computer side, so I thought of introducing two different GANs to generate X and O.

Generative Adversarial Network

We only need a Generator part of the network.

Defining Input and Output node and giving them specify names.

input_node_name = 'noise_input_x'
output_node_name = 'gen_x'

gen_inp = tf.placeholder(tf.float32,shape=[None, noise_dim],name=input_node_name)
disc_inp = tf.placeholder(tf.float32,shape= [None, img_dim])

def gen_fun(x,output_node_name):
    h1 = tf.matmul(x,W["w1"])
    h1 = tf.add(h1,b["b1"])
    h1 = tf.nn.relu(h1)
    
    h1 = tf.matmul(h1,W["w2"])
    h1 = tf.add(h1,b["b2"])
    h1 = tf.nn.sigmoid(h1, name=output_node_name)
    
    return h1

def disc_fun(x):
    h2 = tf.matmul(x,W["w3"])
    h2 = tf.add(h2,b["b3"])
    h2 = tf.nn.relu(h2)
    
    h2 = tf.matmul(h2,W["w4"])
    h2 = tf.add(h2,b["b4"])
    h2 = tf.nn.sigmoid(h2)
    
    return h2

Training output

Extracting Generator from the Graph

freeze_graph.freeze_graph('out/' + MODEL_NAME + '.pbtxt', None, False,
    'out/' + MODEL_NAME + '.chkp', output_node_name, "save/restore_all",
    "save/Const:0", 'out/frozen_' + MODEL_NAME + '.pb', True, "")

input_graph_def = tf.GraphDef()
with tf.gfile.Open('out/frozen_' + MODEL_NAME + '.pb', "rb") as f:
    input_graph_def.ParseFromString(f.read())

input_node_names = ['noise_input_x']
    
output_graph_def = optimize_for_inference_lib.optimize_for_inference(
        input_graph_def, input_node_names, [output_node_name],
        tf.float32.as_datatype_enum)

with tf.gfile.FastGFile('out/opt_' + MODEL_NAME + '.pb', "wb") as f:
    f.write(output_graph_def.SerializeToString())

For the above code, you will get an Optimized graph of the generator network

Putting it in Android

Firstly, copy your Optimized Graph of the generator in asset folder of Android

private String MODEL_GAN_O = "file:///android_asset/opt_gan_O.pb";
private String MODEL_GAN_X = "file:///android_asset/opt_gan_X.pb";

private void loadGANs(){
        inferenceInterface_GAN_O = new TensorFlowInferenceInterface(getAssets(),MODEL_GAN_O);
        inferenceInterface_GAN_X = new TensorFlowInferenceInterface(getAssets(),MODEL_GAN_X);
}

To Generate O, we have to pass the random noise vector to the Graph

public int[] getGAN_O(){

        // random noise input vector    
        float [] noise_input = new float [100];
        for(int i =0 ; i<noise_input.length; i++){
            float value = rand.nextFloat();
            noise_input[i] = (value*2)-1;
        }
        inferenceInterface_GAN_O.feed(INPUT_NODE_GAN_O, noise_input,1,100);
        inferenceInterface_GAN_O.run(new String[]{OUTPUT_NODE_GAN_O},false);

        float [] image = new float[784];
        int [] int_image = new int[784];
        //get the output
        inferenceInterface_GAN_O.fetch(OUTPUT_NODE_GAN_O,image);

        //output image
        for(int i =0 ; i<image.length; i++){
            int_image[i] = (int)(image[i]*255);
        }
        return int_image;
}

Same to Generate X, we have to pass the random noise vector to the Graph

public int[] getGAN_X(){

        // random noise input vector    
        float [] noise_input = new float [100];
        for(int i =0 ; i<noise_input.length; i++){
            float value = rand.nextFloat();
            noise_input[i] = (value*2)-1;
        }
        inferenceInterface_GAN_X.feed(INPUT_NODE_GAN_X, noise_input,1,100);
        inferenceInterface_GAN_X.run(new String[]{OUTPUT_NODE_GAN_X},false);

float [] image = new float[784];
        int [] int_image = new int[784];
        //get the output
        inferenceInterface_GAN_X.fetch(OUTPUT_NODE_GAN_X,image);

        //output image
        for(int i =0 ; i<image.length; i++){
            int_image[i] = (int)(image[i]*255);
        }
        return int_image;
}

This is the result after integrating the CNN and GAN. Completing all the logic of the application.

You can try out the application it is uploaded on the play store: https://play.google.com/store/apps/details?id=com.harsh.xo

Future works: Maybe working on Conditional GAN, which helps me to generate X and O from the same network.

• ConvolutionNeuralNetwork(CNN) in Android — tictactoe.ai (Part 1)
• tictactoe.ai - Apps on Google Play

GenerativeAdversarialNetwork(GAN) in Android — tictactoe.ai (Part 2) was originally published in TDS Archive on Medium, where people are continuing the conversation by highlighting and responding to this story.

Google Tensorflow has now been used in Android for integrating Machine Learning Model.

So I want to create an Android App with Tensorflow, as other people just use the pre-trained model of the Tensorflow and integrate it inside the app. But I want to try something new, so have trained my own model and learned how Tensorflow interacts with Android. To know how to save the model for mobile usage.

Tensorflow

I have implemented code for Convolutional Neural Network to classify X and O for TicTacToe Game.

Input image size is 28x28 and Output layer size is 2. As we have to classify the image in X and O.

While building the model you have to define which is your output node and which is your input node, give those nodes some specify name.

Then training of the model, I have trained my model on 6000 different images of X and O, run over 3 epochs

As the Trained model has trained weight which gives the perfect classification of the X and O. Now, we have to save the graph of the Tensorflow with the weights.

While freezing a graph we require a list of Input and Output node name, that why we have given a name to those to input and output nodes, which helpful while freezing the graph.

Frozen Graph consist of sequence of operation that has to perform to get the output. Operation consist of the matrix multiplication, addition and reshaping of the matrix with the frozen weight of trained model.

Below is the code of freezing a save model.

freeze_graph.freeze_graph('out/' + MODEL_NAME + '.pbtxt', None, False,'out/' + MODEL_NAME + '.chkp', output_node_name, "save/restore_all","save/Const:0", 'out/frozen_' + MODEL_NAME + '.pb', True, "")

input_graph_def = tf.GraphDef()
with tf.gfile.Open('out/frozen_' + MODEL_NAME + '.pb', "rb") as f:
    input_graph_def.ParseFromString(f.read())

input_node_names = ['input']
    
output_graph_def = optimize_for_inference_lib.optimize_for_inference(
        input_graph_def, input_node_names, [output_node_name],
        tf.float32.as_datatype_enum)

with tf.gfile.FastGFile('out/opt_' + MODEL_NAME + '.pb', "wb") as f:
    f.write(output_graph_def.SerializeToString())

Check out the frozen graph of my project here

Android with Tensorflow

Starting with Android, Firstly, we have integrated Tensorflow library just add it to dependencies.

dependencies {
.....
implementation ‘org.tensorflow:tensorflow-android:+’
}

As Tensorflow is coded in C++, so you have to integrate NDK in our Android project. Then copy the model in the asset folder of the android directory.

How to run the graph in the Android using Java?

// added the path of model
private static final String MODEL_FILE_CNN = "file:///android_asset/opt_xo_differ_v2.pb";
//declare the input and output
private static final String INPUT_NODE = "input";
private static final String OUTPUT_NODE = "output"private TensorFlowInferenceInterface inferenceInterface;

//load graph
inferenceInterface = new TensorFlowInferenceInterface(getAssets(),MODEL_FILE_CNN);

Loading the model .pd file and create a TensorFlowInferenceInterface object

For detecting, I have created a method which reads Image pixels in a float. As DrawView is a canvas which gives us the pixels value of canvas drawn image then converting them into float and pass it to TensorflowInferenceInterface object.

So first you have to feed the data in Input Node (input_node and keep_prob), then run the InferenceInterface with the output_node

public float[] getDetectofDV(DrawView d){

float pixels[] = d.getPixelData();

inferenceInterface.feed(INPUT_NODE, pixels,1,28,28,1);
inferenceInterface.feed("keep_prob", new float[] { 1 });

final Operation operation = inferenceInterface.graphOperation(OUTPUT_NODE);

final int numClasses = (int) operation.output(0).shape().size(1);

// this run the operation of the graph
inferenceInterface.run(new String[]{OUTPUT_NODE},false);

float [] outputarray = new float[numClasses];
//get the output
inferenceInterface.fetch(OUTPUT_NODE,outputarray);

return outputarray;
}

outputarray returns the answer of the detection which has 2 values if one index 0 value in greater than index 1. Then the image is detected to be ‘O’, else it is ‘X’.

This is the result after integrating the model with Android. As if work on touch listener when you pick up your finger from the canvas it pass the image for the classification and you will get the output in below matrix.

Check out the application on playstore. https://play.google.com/store/apps/details?id=com.harsh.xo

Future work

First I will try to reduce the size of the model. Model size of CNN is 17MB as it took a lot of space in the application. And the final size of the application 55 MB.

Working on integrating new model in the same application. A single user can play with the computer.

https://play.google.com/store/apps/details?id=com.harsh.xo

ConvolutionNeuralNetwork(CNN) in Android — tictactoe.ai (Part 1) was originally published in TDS Archive on Medium, where people are continuing the conversation by highlighting and responding to this story.

As it’s small one but , I think it will be helpful to developer. This library is to get Firebase Reference JSON value.

While using ValueChangeListener to get data, I found some difficulties because it ran the code every time when data got change. So I wanted to get rid of this ValueChangeListener to get the data of the given reference.

That’s the reason I developed FirebaseJSONReference.

Details need in Release 0.1.1

Firebase URL: “https://example.firebaseio.com"

Reference URL: “/data/value”

User Token: Auth token from Firebase

How to get Auth Token :-

FirebaseUser mUser = FirebaseAuth.getInstance().getCurrentUser();
mUser.getIdToken(true)
 .addOnCompleteListener(new OnCompleteListener<GetTokenResult>() {
public void onComplete(@NonNull Task<GetTokenResult> task) {
if (task.isSuccessful()) {
String token = task.getResult().getToken();
// Send token to your backend via HTTPS
// ...
} else {
// Handle error -> task.getException();
}
}
});

Add Dependencies :-

Add following to your root build.gradle at the end of repositories:

allprojects {
repositories {
 ...
maven { url 'https://jitpack.io' }
}
}

Add following dependency

dependencies {
compile 'com.github.harsh2011:FirebaseJSONReference:0.1.1'
}

Demo of Library:-

How to use it.

FirebaseJSONReference firebaseJSONReference = new FirebaseJSONReference(MainActivity.this);
firebaseJSONReference.setFirebaseDatabaseURL("https://example.firebaseio.com");
firebaseJSONReference.setFirebaseRef("/requestlist");
firebaseJSONReference.setAuth("auth token");
firebaseJSONReference.setResponseListener(new ResponseListener() {
@Override
public void onResponse(String s) {
System.out.println("response:"+s);
try {
//JSONobject response
JSONObject main = new JSONObject(s);

} catch (JSONException e) {
e.printStackTrace();
}

}
});
//build the request
firebaseJSONReference.build();
//execute the request and get response in Response Listener
firebaseJSONReference.execute();

So I start finding out Libraries to ready Audio data

Libraries used:-

1. ffmpy (python) For file conversion
2. soundfile (python) Read and Write the WAV file in python
3. numpy (pyhton) Handling N dimensional array
4. cv2 (python) Use for writing the PNG image

Reading of WAV file

read() method of soundfile returns two thing

1. samplerate
2. 2.data (frames)(data types are (float,int))(array of data)

As this Program reads int16 ,datatype. range of (-32768 to 32767)

One of the given data

data = 6000 //range of (-32768 to 32767)
data = data + 32768
// divide data bits(data is of 16 bits)
datafront = first8bitsof(data) // 8bits
datalast = last8bitsof(data) // 8bits

then,for the 2 channel in the data of 16 bits convert in 4 channel of 8 bits.which easy represent the image pixel value between 0 to 255

Using open write the data in image form. But image file must be PNG(lossless compression).

Check out whole code on my github profile:-

https://github.com/harsh2011/Steganography/tree/master/AudiotoImage(hidding%20audio%20in%20image)

Here is an ENCODE IMAGE:-