Merge branch 'privatestaging' of https://github.com/AMDComputeLibraries/HIP-privatestaging into privatestaging

[ROCm/clr commit: bf83b949f6]

projects/clr/hipamd/src/hip_hcc.cpp

@@ -1151,8 +1151,6 @@ INLINE ihipDevice_t *ihipGetDevice(int deviceId)
 
 }
 
-
-
 //---
 // Get the stream to use for a command submission.
 //
@@ -1162,8 +1160,10 @@ inline hipStream_t ihipSyncAndResolveStream(hipStream_t stream)
 {
     if (stream == hipStreamNull ) {
         ihipDevice_t *device = ihipGetTlsDefaultDevice();
-        device->syncDefaultStream(false);
 
+#ifndef HIP_API_PER_THREAD_DEFAULT_STREAM
+        device->syncDefaultStream(false);
+#endif
         return device->_default_stream;
     } else {
         // Have to wait for legacy default stream to be empty:
@@ -2624,7 +2624,7 @@ hipError_t hipHostFree(void* ptr)
 // TODO - deprecated function.
 hipError_t hipFreeHost(void* ptr)
 {
-    hipHostFree(ptr);
+    return hipHostFree(ptr);
 }
 
 

projects/clr/hipamd/tests/src/CMakeLists.txt

@@ -109,6 +109,8 @@ macro (make_test_matches exe match_string)
     )
 endmacro()
 
+make_hip_executable (hipAPIStreamEnable hipAPIStreamEnable.cpp)
+make_hip_executable (hipAPIStreamDisable hipAPIStreamDisable.cpp)
 make_hip_executable (hip_ballot hip_ballot.cpp)
 make_hip_executable (hip_anyall hip_anyall.cpp) 
 make_hip_executable (hip_popc hip_popc.cpp) 
@@ -138,7 +140,6 @@ make_hip_executable (hipStreamL5 hipStreamL5.cpp)
 make_hip_executable (hipHostGetFlags hipHostGetFlags.cpp)
 make_hip_executable (hipHostRegister hipHostRegister.cpp)
 make_hip_executable (hipRandomMemcpyAsync hipRandomMemcpyAsync.cpp)
-
 make_test(hip_ballot " " )
 make_test(hip_anyall " " )
 make_test(hip_popc " " )
@@ -159,6 +160,7 @@ make_test(hipMemcpy_simple  " " )
 make_named_test(hipMemcpy "hipMemcpy-modes" --tests 0x1 )
 make_named_test(hipMemcpy "hipMemcpy-size" --tests 0x6 )
 make_named_test(hipMemcpy "hipMemcpy-multithreaded" --tests 0x8 )
+
 make_test(hipHostAlloc " ")
 make_test(hipMemcpyAsync  " " )
 make_test(hipHostGetFlags " ")
@@ -166,4 +168,6 @@ make_test(hipHcc  " " )
 make_test(hipHostRegister " ")
 make_test(hipStreamL5 " ")
 make_test(hipRandomMemcpyAsync " ")
+make_test(hipAPIStreamEnable " ")
+make_test(hipAPIStreamDisable " ")
 make_hipify_test(specialFunc.cu )

projects/clr/hipamd/tests/src/hipAPIStreamDisable.cpp

@@ -0,0 +1,61 @@
+/*
+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 WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include<iostream>
+#include"test_common.h"
+
+const int NN = 1 << 21;
+
+__global__ void kernel(hipLaunchParm lp, float *x, float *y, int n){
+	int tid = hipThreadIdx_x;
+	if(tid < 1){
+		for(int i=0;i<n;i++){
+			x[i] = sqrt(pow(3.14159,i));
+		}
+		y[tid] = y[tid] + 1.0f;
+	}
+}
+
+__global__ void nKernel(hipLaunchParm lp, float *y){
+	int tid = hipThreadIdx_x;
+	y[tid] = y[tid] + 1.0f;
+}
+
+int main(){
+	const int num_streams = 8;
+	hipStream_t streams[num_streams];
+	float *data[num_streams], *yd, *xd;
+	float y = 1.0f, x = 1.0f;
+	HIPCHECK(hipMalloc((void**)&yd, sizeof(float)));
+	HIPCHECK(hipMalloc((void**)&xd, sizeof(float)));
+	HIPCHECK(hipMemcpy(yd, &y, sizeof(float), hipMemcpyHostToDevice));
+	HIPCHECK(hipMemcpy(xd, &x, sizeof(float), hipMemcpyHostToDevice));
+	for(int i=0;i<num_streams;i++){
+		HIPCHECK(hipStreamCreate(&streams[i]));
+		HIPCHECK(hipMalloc(&data[i], NN * sizeof(float)));
+		hipLaunchKernel(HIP_KERNEL_NAME(kernel), dim3(1), dim3(1), 0, streams[i], data[i], xd, N);
+		hipLaunchKernel(HIP_KERNEL_NAME(nKernel), dim3(1), dim3(1), 0, 0, yd);
+	}
+
+	HIPCHECK(hipMemcpy(&x, xd, sizeof(float), hipMemcpyDeviceToHost));
+	HIPCHECK(hipMemcpy(&y, yd, sizeof(float), hipMemcpyDeviceToHost));
+	std::cout<<x<<" "<<y<<std::endl;
+	HIPASSERT(x == y);
+	passed();
+}

projects/clr/hipamd/tests/src/hipAPIStreamEnable.cpp

@@ -0,0 +1,63 @@
+/*
+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 WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+//#define HIP_API_PER_THREAD_DEFAULT_STREAM
+
+#include<iostream>
+#include"test_common.h"
+
+const int NN = 1 << 21;
+
+__global__ void kernel(hipLaunchParm lp, float *x, float *y, int n){
+	int tid = hipThreadIdx_x;
+	if(tid < 1){
+		for(int i=0;i<n;i++){
+			x[i] = sqrt(pow(3.14159,i));
+		}
+		y[tid] = y[tid] + 1.0f;
+	}
+}
+
+__global__ void nKernel(hipLaunchParm lp, float *y){
+	int tid = hipThreadIdx_x;
+	y[tid] = y[tid] + 1.0f;
+}
+
+int main(){
+	const int num_streams = 8;
+	hipStream_t streams[num_streams];
+	float *data[num_streams], *yd, *xd;
+	float y = 1.0f, x = 1.0f;
+	HIPCHECK(hipMalloc((void**)&yd, sizeof(float)));
+	HIPCHECK(hipMalloc((void**)&xd, sizeof(float)));
+	HIPCHECK(hipMemcpy(yd, &y, sizeof(float), hipMemcpyHostToDevice));
+	HIPCHECK(hipMemcpy(xd, &x, sizeof(float), hipMemcpyHostToDevice));
+	for(int i=0;i<num_streams;i++){
+		HIPCHECK(hipStreamCreate(&streams[i]));
+		HIPCHECK(hipMalloc(&data[i], NN * sizeof(float)));
+		hipLaunchKernel(HIP_KERNEL_NAME(kernel), dim3(1), dim3(1), 0, streams[i], data[i], xd, N);
+		hipLaunchKernel(HIP_KERNEL_NAME(nKernel), dim3(1), dim3(1), 0, 0, yd);
+	}
+
+	HIPCHECK(hipMemcpy(&x, xd, sizeof(float), hipMemcpyDeviceToHost));
+	HIPCHECK(hipMemcpy(&y, yd, sizeof(float), hipMemcpyDeviceToHost));
+	std::cout<<x<<" "<<y<<std::endl;
+	HIPASSERT(x<y);
+	passed();
+}

projects/clr/hipamd/tests/src/testAPIStream.sh

@@ -0,0 +1,8 @@
+#!/bin/bash
+
+rm $HIP_PATH/src/hip_hcc.o
+mkdir build
+$HIP_PATH/bin/hipcc hipAPIStreamDisable.cpp test_common.cpp -o ./build/hipAPIStreamDisable
+rm $HIP_PATH/src/hip_hcc.o
+$HIP_PATH/bin/hipcc hipAPIStreamEnable.cpp test_common.cpp -o ./build/hipAPIStreamEnable
+rm $HIP_PATH/src/hip_hcc.o