Organize testing directory

Author: virnarula

llvm/test/Transforms/NVPTXMemOpts/already_coalesced/driver.cpp

@@ -0,0 +1,32 @@
+#include <stdio.h>
+
+// CUDA kernel for mat mul
+__global__ void naiveMM(const float **A, const float **B, float **C, int w) {
+    int idx = blockDim.x * blockIdx.x + threadIdx.x;
+    int idy = blockDim.y * blockIdx.y + threadIdx.y;
+    float sum = 0;
+    for (int i = 0; i < w; i++) {
+        sum += A[idy][i] * B[i][idx];
+    }
+    C[idy][idx] = sum;
+
+}
+
+// CUDA kernel for mat mul with shared memory
+__global__ void sharedMM(const float **A, const float **B, float **C, int w) {
+    __shared__ float As[32][32];
+    __shared__ float Bs[32][32];
+    int idx = blockDim.x * blockIdx.x + threadIdx.x;
+    int idy = blockDim.y * blockIdx.y + threadIdx.y;
+    float sum = 0;
+    for (int i = 0; i < w; i++) {
+        As[threadIdx.y][i] = A[idy][i];
+        Bs[i][threadIdx.x] = B[i][idx];
+        __syncthreads();
+        for (int j = 0; j < w; j++) {
+            sum += As[threadIdx.y][j] * Bs[j][threadIdx.x];
+        }
+        __syncthreads();
+    }
+    C[idy][idx] = sum;
+}

llvm/test/Transforms/NVPTXMemOpts/already_coalesced/test.cu

@@ -0,0 +1,13 @@
+// CUDA kernel for constant access
+__global__ void naive_add(const float *A, float **C, int w) {
+    int idx = blockDim.x * blockIdx.x + threadIdx.x;
+    int idy = blockDim.y * blockIdx.y + threadIdx.y;
+    float sum = 0;
+    for (int i = 0; i < w; i++) {
+        for (int j = 0; j < w; j++) {
+            sum += A[10][15] * B[23][54];
+        }
+    }
+    C[idy][idx] = sum;
+
+}

llvm/test/Transforms/NVPTXMemOpts/basic/test.cu

@@ -0,0 +1,17 @@
+#include <stdio.h>
+
+// the reads in this kernel are of constant indexes and therefore cannot be coallesced
+
+// CUDA kernel for constant access
+__global__ void naiveMM(const float **A, const float **B, float **C, int w) {
+    int idx = blockDim.x * blockIdx.x + threadIdx.x;
+    int idy = blockDim.y * blockIdx.y + threadIdx.y;
+    float sum = 0;
+    for (int i = 0; i < w; i++) {
+        for (int j = 0; j < w; j++) {
+            sum += A[10][15] * B[23][54];
+        }
+    }
+    C[idy][idx] = sum;
+
+}

llvm/test/Transforms/NVPTXMemOpts/constant_access/test.cu

@@ -0,0 +1,29 @@
+#include <cuda_runtime.h>
+
+// CUDA kernel for element-wise addition of two matrices
+__global__ void matrixAdd(const int *A, const int *B, int *C, int numRows, int numCols) {
+    int row = blockIdx.y * blockDim.y + threadIdx.y;
+    int col = blockIdx.x * blockDim.x + threadIdx.x;
+    if (row < numRows && col < numCols) {
+        int idx = row * numCols + col;
+        C[idx] = A[idx] + B[idx];
+    }
+}
+
+
+// CUDA kernel for element-wise addition of two matrices with memory coalescing
+__global__ void matrixAdd_coalesced(const int *A, const int *B, int *C, int numRows, int numCols) {
+    int row = blockIdx.y * blockDim.y + threadIdx.y;
+    int col = blockIdx.x * blockDim.x + threadIdx.x;
+    __shared__ int A_shared[16][16];
+    __shared__ int B_shared[16][16];
+    A_shared[threadIdx.y][threadIdx.x] = A[row * numCols + col];
+    B_shared[threadIdx.y][threadIdx.x] = B[row * numCols + col];
+    __syncthreads();
+    if (row < numRows && col < numCols) {
+        int idx = row * numCols + col;
+        C[idx] = A_shared[threadIdx.y][threadIdx.x] + B_shared[threadIdx.y][threadIdx.x];
+    }
+}
+
+ 

llvm/test/Transforms/NVPTXMemOpts/driver.cpp renamed to llvm/test/Transforms/NVPTXMemOpts/matrix_add/driver.cpp

@@ -12,6 +12,22 @@ __global__ void transposeMatrix(const float *input, float *output, int width, in
     }
 }
 
+// CUDA kernel to transpose a matrix
+__global__ void transposeMatrix_coalesced(const float *input, float *output, int width, int height) {
+    int x = blockIdx.x * blockDim.x + threadIdx.x;
+    int y = blockIdx.y * blockDim.y + threadIdx.y;
+
+    __shared__ float input_shared[16][16];
+    input_shared[threadIdx.y][threadIdx.x] = input[y * width + x];
+    __syncthreads();
+
+    if (x < width && y < height) {
+        int pos = y * width + x;
+        int transPos = x * height + y;
+        output[transPos] = input[pos];
+    }
+}
+
 int main(void) {
     int width = 1024;
     int height = 1024;

llvm/test/Transforms/NVPTXMemOpts/matrix_add/matrix_add.cu

@@ -8,6 +8,19 @@ __global__ void vectorAdd(const float *A, const float *B, float *C, int numEleme
     }
 }
 
+// Cuda kernel for vector addition with memory coalescing
+__global__ void vectorAdd_coalesced(const float *A, const float *B, float *C, int numElements) {
+    int i = blockDim.x * blockIdx.x + threadIdx.x;
+    __shared__ float A_shared[16];
+    __shared__ float B_shared[16];
+    A_shared[threadIdx.x] = A[i];
+    B_shared[threadIdx.x] = B[i];
+    __syncthreads();
+    if (i < numElements) {
+        C[i] = A_shared[threadIdx.x] + B_shared[threadIdx.x];
+    }
+}
+
 int main(void) {
     int numElements = 50000;
     size_t size = numElements * sizeof(float);