From a60649868d06915e0a7229c38cc93c7c3df8e12c Mon Sep 17 00:00:00 2001
From: Aditya Atluri <adityaavinash1@gmail.com>
Date: Wed, 16 Mar 2016 05:15:03 -0500
Subject: [PATCH 1/4] added cudaHostRegister test

---
 hipamd/tests/src/cudaRegister.cu | 80 ++++++++++++++++++++++++++++++++
 1 file changed, 80 insertions(+)
 create mode 100644 hipamd/tests/src/cudaRegister.cu

diff --git a/hipamd/tests/src/cudaRegister.cu b/hipamd/tests/src/cudaRegister.cu
new file mode 100644
index 0000000000..1fc456f5fa
--- /dev/null
+++ b/hipamd/tests/src/cudaRegister.cu
@@ -0,0 +1,80 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include<cuda.h>
+#include<cuda_runtime.h>
+#include<iostream>
+#include<unistd.h>
+#include<stdio.h>
+#include<malloc.h>
+
+#define LEN 1024*1024
+#define SIZE LEN * sizeof(float)
+#define ITER 1024*128
+
+#define check(msg, status){ \
+if(status != cudaSuccess) { \
+	printf("%s failed. \n", #msg); \
+} \
+}
+
+__global__ void Inc1(float *Ad){
+	int tx = threadIdx.x + blockIdx.x * blockDim.x;
+	if(tx < 1 ){
+		for(int i=0;i<ITER;i++){
+			Ad[tx] = Ad[tx] + 1.0f;
+		}
+	}
+}
+
+__global__ void Inc2(float *Ad){
+	int tx = threadIdx.x + blockIdx.x * blockDim.x;
+	if(tx < 1024){
+		for(int i=0;i<ITER;i++){
+			Ad[tx] = Ad[tx] + 1.0f;
+		}
+	}
+}
+
+int main(){
+	float *A, *Ad;
+	A = new float[LEN];
+	for(int i=0;i<LEN;i++){
+		A[i] = 0.0f;
+	}
+	cudaError_t status;
+	status = cudaHostRegister(A, SIZE, cudaHostRegisterMapped);
+	check("Registering A",status);
+	cudaHostGetDevicePointer(&Ad, A, 0);
+
+	dim3 dimGrid(LEN/512,1,1);
+	dim3 dimBlock(512,1,1);
+	Inc1<<<dimGrid, dimBlock>>>(Ad);
+	A[1] = -(ITER*1.0f);
+	cudaDeviceSynchronize();
+	std::cout<<A[1]<<std::endl;
+
+	for(int i=0;i<LEN;i++){
+		A[i] = 0.0f;
+	}
+	Inc2<<<dimGrid, dimBlock>>>(Ad);
+	A[1] = -(ITER*1.0f);
+	cudaDeviceSynchronize();
+	std::cout<<A[1]<<std::endl;
+}

From 8f527d2288aaf9fcf34fd3ab4525e6ab746001c5 Mon Sep 17 00:00:00 2001
From: Aditya Atluri <adityaavinash1@gmail.com>
Date: Wed, 16 Mar 2016 05:24:08 -0500
Subject: [PATCH 2/4] increased iteration size

---
 hipamd/tests/src/cudaRegister.cu | 11 ++++++-----
 1 file changed, 6 insertions(+), 5 deletions(-)

diff --git a/hipamd/tests/src/cudaRegister.cu b/hipamd/tests/src/cudaRegister.cu
index 1fc456f5fa..6ea038a71a 100644
--- a/hipamd/tests/src/cudaRegister.cu
+++ b/hipamd/tests/src/cudaRegister.cu
@@ -26,7 +26,7 @@ THE SOFTWARE.
 
 #define LEN 1024*1024
 #define SIZE LEN * sizeof(float)
-#define ITER 1024*128
+#define ITER 1024*1024
 
 #define check(msg, status){ \
 if(status != cudaSuccess) { \
@@ -66,15 +66,16 @@ int main(){
 	dim3 dimGrid(LEN/512,1,1);
 	dim3 dimBlock(512,1,1);
 	Inc1<<<dimGrid, dimBlock>>>(Ad);
-	A[1] = -(ITER*1.0f);
+	A[0] = -(ITER*1.0f);
+	std::cout<<A[0]<<std::endl;
 	cudaDeviceSynchronize();
-	std::cout<<A[1]<<std::endl;
+	std::cout<<A[0]<<std::endl;
 
 	for(int i=0;i<LEN;i++){
 		A[i] = 0.0f;
 	}
 	Inc2<<<dimGrid, dimBlock>>>(Ad);
-	A[1] = -(ITER*1.0f);
+	A[0] = -(ITER*1.0f);
 	cudaDeviceSynchronize();
-	std::cout<<A[1]<<std::endl;
+	std::cout<<A[0]<<std::endl;
 }

From 9b772c855384720d7d4dd65decccb7d07d8f176e Mon Sep 17 00:00:00 2001
From: Aditya Atluri <adityaavinash1@gmail.com>
Date: Wed, 16 Mar 2016 05:30:29 -0500
Subject: [PATCH 3/4] test/src [v3] clean up

---
 hipamd/tests/src/cudaRegister.cu | 9 +++++----
 1 file changed, 5 insertions(+), 4 deletions(-)

diff --git a/hipamd/tests/src/cudaRegister.cu b/hipamd/tests/src/cudaRegister.cu
index 6ea038a71a..8d97c7ae59 100644
--- a/hipamd/tests/src/cudaRegister.cu
+++ b/hipamd/tests/src/cudaRegister.cu
@@ -24,7 +24,7 @@ THE SOFTWARE.
 #include<stdio.h>
 #include<malloc.h>
 
-#define LEN 1024*1024
+#define LEN 1024
 #define SIZE LEN * sizeof(float)
 #define ITER 1024*1024
 
@@ -67,15 +67,16 @@ int main(){
 	dim3 dimBlock(512,1,1);
 	Inc1<<<dimGrid, dimBlock>>>(Ad);
 	A[0] = -(ITER*1.0f);
-	std::cout<<A[0]<<std::endl;
+	std::cout<<"Same cache line before completion: \t"<< A[0]<<std::endl;
 	cudaDeviceSynchronize();
-	std::cout<<A[0]<<std::endl;
+	std::cout<<"Same cache line after completion: \t"<< A[0]<<std::endl;
 
 	for(int i=0;i<LEN;i++){
 		A[i] = 0.0f;
 	}
 	Inc2<<<dimGrid, dimBlock>>>(Ad);
 	A[0] = -(ITER*1.0f);
+	std::cout<<"Diff cache line before completion: \t"<<A[0]<<std::endl;
 	cudaDeviceSynchronize();
-	std::cout<<A[0]<<std::endl;
+	std::cout<<"Diff cache line after completion: \t"<<A[0]<<std::endl;
 }

From 6baeb08e1c525e289b26c35d67835481fe947b80 Mon Sep 17 00:00:00 2001
From: Aditya Atluri <adityaavinash1@gmail.com>
Date: Wed, 16 Mar 2016 07:04:40 -0500
Subject: [PATCH 4/4] tests/src [v4] Added feature for partial writes on CPU

---
 hipamd/tests/src/cudaRegister.cu | 20 ++++++++++++++------
 1 file changed, 14 insertions(+), 6 deletions(-)

diff --git a/hipamd/tests/src/cudaRegister.cu b/hipamd/tests/src/cudaRegister.cu
index 8d97c7ae59..b9953a5a3f 100644
--- a/hipamd/tests/src/cudaRegister.cu
+++ b/hipamd/tests/src/cudaRegister.cu
@@ -34,26 +34,32 @@ if(status != cudaSuccess) { \
 } \
 }
 
-__global__ void Inc1(float *Ad){
+__global__ void Inc1(float *Ad, float *Bd){
 	int tx = threadIdx.x + blockIdx.x * blockDim.x;
 	if(tx < 1 ){
 		for(int i=0;i<ITER;i++){
 			Ad[tx] = Ad[tx] + 1.0f;
+			for(int j=0;j<256;j++){
+				Bd[tx] = Ad[tx];
+			}
 		}
 	}
 }
 
-__global__ void Inc2(float *Ad){
+__global__ void Inc2(float *Ad, float *Bd){
 	int tx = threadIdx.x + blockIdx.x * blockDim.x;
 	if(tx < 1024){
 		for(int i=0;i<ITER;i++){
 			Ad[tx] = Ad[tx] + 1.0f;
+			for(int j=0;j<256;j++){
+				Bd[tx] = Ad[tx];
+			}
 		}
 	}
 }
 
 int main(){
-	float *A, *Ad;
+	float *A, *Ad, *Bd;
 	A = new float[LEN];
 	for(int i=0;i<LEN;i++){
 		A[i] = 0.0f;
@@ -62,10 +68,11 @@ int main(){
 	status = cudaHostRegister(A, SIZE, cudaHostRegisterMapped);
 	check("Registering A",status);
 	cudaHostGetDevicePointer(&Ad, A, 0);
-
+	cudaMalloc((void**)&Bd, SIZE);
 	dim3 dimGrid(LEN/512,1,1);
 	dim3 dimBlock(512,1,1);
-	Inc1<<<dimGrid, dimBlock>>>(Ad);
+	Inc1<<<dimGrid, dimBlock>>>(Ad, Bd);
+	sleep(3);
 	A[0] = -(ITER*1.0f);
 	std::cout<<"Same cache line before completion: \t"<< A[0]<<std::endl;
 	cudaDeviceSynchronize();
@@ -74,7 +81,8 @@ int main(){
 	for(int i=0;i<LEN;i++){
 		A[i] = 0.0f;
 	}
-	Inc2<<<dimGrid, dimBlock>>>(Ad);
+	Inc2<<<dimGrid, dimBlock>>>(Ad, Bd);
+	sleep(3);
 	A[0] = -(ITER*1.0f);
 	std::cout<<"Diff cache line before completion: \t"<<A[0]<<std::endl;
 	cudaDeviceSynchronize();