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

[ROCm/clr commit: 6baeb08e1c525e289b26c35d67835481fe947b80]
---
 projects/clr/hipamd/tests/src/cudaRegister.cu | 20 +++++++++++++------
 1 file changed, 14 insertions(+), 6 deletions(-)

diff --git a/projects/clr/hipamd/tests/src/cudaRegister.cu b/projects/clr/hipamd/tests/src/cudaRegister.cu
index 8d97c7ae59..b9953a5a3f 100644
--- a/projects/clr/hipamd/tests/src/cudaRegister.cu
+++ b/projects/clr/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();