Updated the naming of matrices to M x N and h_X

matAdd.cu

@@ -2,86 +2,86 @@
 #include <stdlib.h>
 
 // matrix addition kernel
-__global__ void matAdd(float *d_A, float *d_B, float *d_C, int N, int M) {
+__global__ void matAdd(float *d_A, float *d_B, float *d_C, int M, int N) {
 	int row = blockIdx.y * blockDim.y + threadIdx.y;
 	int col = blockIdx.x * blockDim.x + threadIdx.x;
 
 	// add matrix elements
-	if (row < N && col < M) {
-		d_C[row * M + col] = d_A[row * M + col] + d_B[row * M + col];
+	if (row < M && col < N) {
+		d_C[row * N + col] = d_A[row * N + col] + d_B[row * N + col];
 	}
 }
 
 int main() {
   	// var declaration
 	int N = 5;
 	int M = 5;
-	float *A, *B, *C;
+	float *h_A, *h_B, *h_C;
 	float *d_A, *d_B, *d_C;
 
 	// allocate host memory
-	A = (float *)malloc(N * M * sizeof(float));
-	B = (float *)malloc(N * M * sizeof(float));
-        C = (float *)malloc(N * M * sizeof(float));
+	h_A = (float *)malloc(M * N * sizeof(float));
+	h_B = (float *)malloc(M * N * sizeof(float));
+        h_C = (float *)malloc(M * N * sizeof(float));
 
 	// allocate device memory
-	cudaMalloc(&d_A, N * M * sizeof(float));
-	cudaMalloc(&d_B, N * M * sizeof(float));
-	cudaMalloc(&d_C, N * M * sizeof(float));
+	cudaMalloc(&d_A, M * N * sizeof(float));
+	cudaMalloc(&d_B, M * N * sizeof(float));
+	cudaMalloc(&d_C, M * N * sizeof(float));
 
 	// initialize data
-	for (int i = 0; i < N * M; ++i) {
-		A[i] = i - 3;
-		B[i] = i;
+	for (int i = 0; i < M * N; ++i) {
+		h_A[i] = i - 3;
+		h_B[i] = i;
 	}
 
 	// copy host data to device
-	cudaMemcpy(d_A, A, N * M * sizeof(float), cudaMemcpyHostToDevice);
-	cudaMemcpy(d_B, B, N * M * sizeof(float), cudaMemcpyHostToDevice);
-	cudaMemcpy(d_C, C, N * M * sizeof(float), cudaMemcpyHostToDevice);
+	cudaMemcpy(d_A, h_A, M * N * sizeof(float), cudaMemcpyHostToDevice);
+	cudaMemcpy(d_B, h_B, M * N * sizeof(float), cudaMemcpyHostToDevice);
+	cudaMemcpy(d_C, h_C, M * N * sizeof(float), cudaMemcpyHostToDevice);
 
 	// launch kernel instance
 	dim3 blockDim(16, 16);
-	dim3 gridDim((M + blockDim.x - 1)/blockDim.x, (N + blockDim.y - 1)/blockDim.y);
-	matAdd<<<gridDim, blockDim>>>(d_A, d_B, d_C, N, M);
+	dim3 gridDim((N + blockDim.x - 1)/blockDim.x, (M + blockDim.y - 1)/blockDim.y);
+	matAdd<<<gridDim, blockDim>>>(d_A, d_B, d_C, M, N);
 
 	// copy result back to host
-	cudaMemcpy(A, d_A, N * M * sizeof(float), cudaMemcpyDeviceToHost);
-	cudaMemcpy(B, d_B, N * M * sizeof(float), cudaMemcpyDeviceToHost);
-	cudaMemcpy(C, d_C, N * M * sizeof(float), cudaMemcpyDeviceToHost);
+	cudaMemcpy(h_A, d_A, M * N * sizeof(float), cudaMemcpyDeviceToHost);
+	cudaMemcpy(h_B, d_B, M * N * sizeof(float), cudaMemcpyDeviceToHost);
+	cudaMemcpy(h_C, d_C, M * N * sizeof(float), cudaMemcpyDeviceToHost);
 
   	// display results
 	printf("Matrix A: \n");
 	printf("----------\n");
-	for (int i = 0; i < N; ++i) {
-		for (int j = 0; j < M; ++j) {
-			printf("A: %f ", A[i * M + j]);
+	for (int i = 0; i < M; ++i) {
+		for (int j = 0; j < N; ++j) {
+			printf("A: %f ", h_A[i * N + j]);
 		}
 		printf("\n)
 	}
 
 	printf("----------\n");
 	printf("Matrix B: \n");
 	printf("----------\n");
-	for (int i = 0; i < N; ++i) {
-		for (int j = 0; j < M; ++j) {
-			printf("B: %f ", B[i * M + j]);
+	for (int i = 0; i < M; ++i) {
+		for (int j = 0; j < N; ++j) {
+			printf("B: %f ", h_B[i * N + j]);
 		}
 		printf("\n)
 	}
 
 	printf("----------\n");
 	printf("Matrix C: \n");
 	printf("----------\n");
-	for (int i = 0; i < N; ++i) {
-		for (int j = 0; j < M; ++j) {
-			printf("C: %f ", C[i * M + j]);
+	for (int i = 0; i < M; ++i) {
+		for (int j = 0; j < N; ++j) {
+			printf("C: %f ", h_C[i * N + j]);
 		}
 		printf("\n)
 	}
 
 	// clean up data
-	free(A); free(B); free(C);
+	free(h_A); free(h_B); free(h_C);
 	cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
 
 	return 0;