Merge remote-tracking branch 'origin/master' into feature_use_module_based_dispatch_instead_of_pfe

[ROCm/hip commit: adaf6b8dff]

projects/hip/src/hip_event.cpp

@@ -45,10 +45,10 @@ ihipEvent_t::ihipEvent_t(unsigned flags)
 void ihipEvent_t::attachToCompletionFuture(const hc::completion_future *cf, 
                                            hipStream_t stream, ihipEventType_t eventType)
 {
-    _state  = hipEventStatusRecording;
     _marker = *cf;
     _type   = eventType;
     _stream = stream;
+    _state  = hipEventStatusRecording;
 }
 
 
@@ -157,13 +157,12 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream)
             event->_state = hipEventStatusComplete;
             return ihipLogStatus(hipSuccess);
         } else {
-            event->_state  = hipEventStatusRecording;
             // Clear timestamps
             event->_timestamp = 0;
 
             // Record the event in the stream:
             stream->locked_recordEvent(event);
-
+            event->_state  = hipEventStatusRecording;
             return ihipLogStatus(hipSuccess);
         }
     } else {
@@ -191,6 +190,10 @@ hipError_t hipEventSynchronize(hipEvent_t event)
 {
     HIP_INIT_SPECIAL_API(TRACE_SYNC, event);
 
+    if (!(event->_flags & hipEventReleaseToSystem)) {
+        tprintf(DB_WARN, "hipEventSynchronize on event without system-scope fence ; consider creating with hipEventReleaseToSystem\n");
+    }
+
     if (event) {
         if (event->_state == hipEventStatusUnitialized) {
             return ihipLogStatus(hipErrorInvalidResourceHandle);
@@ -268,6 +271,10 @@ hipError_t hipEventQuery(hipEvent_t event)
 {
     HIP_INIT_SPECIAL_API(TRACE_QUERY, event);
 
+    if (!(event->_flags & hipEventReleaseToSystem)) {
+        tprintf(DB_WARN, "hipEventQuery on event without system-scope fence ; consider creating with hipEventReleaseToSystem\n");
+    }
+
     if ((event->_state == hipEventStatusRecording) && !event->locked_isReady()) {
         return ihipLogStatus(hipErrorNotReady);
     } else {

projects/hip/src/hip_hcc_internal.h

@@ -218,7 +218,8 @@ extern const char *API_COLOR_END;
 #define DB_SYNC   1 /* 0x02 - trace synchronization pieces */
 #define DB_MEM    2 /* 0x04 - trace memory allocation / deallocation */
 #define DB_COPY   3 /* 0x08 - trace memory copy and peer commands. . */
-#define DB_MAX_FLAG 4
+#define DB_WARN   4 /* 0x10 - warn about sub-optimal or shady behavior */
+#define DB_MAX_FLAG 5
 // When adding a new debug flag, also add to the char name table below.
 //
 //
@@ -235,6 +236,7 @@ static const DbName dbName [] =
     {KYEL, "sync"},
     {KCYN, "mem"},
     {KMAG, "copy"},
+    {KRED, "warn"},
 };
 
 

projects/hip/src/hip_stream.cpp

@@ -103,7 +103,7 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
             event->locked_waitComplete((event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked : hc::hcWaitModeActive);
         } else {
             stream = ihipSyncAndResolveStream(stream);
-            // This will user create_blocking_marker to wait on the specified queue.
+            // This will use create_blocking_marker to wait on the specified queue.
             stream->locked_streamWaitEvent(event);
         }
 

projects/hip/tests/src/runtimeApi/device/hipDeviceComputeCapability.cpp

@@ -0,0 +1,46 @@
+/*
+Copyright (c) 2015-2017 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.
+*/
+
+/*
+ * Conformance test for checking functionality of
+ * hipError_t hipDeviceComputeCapability(int* major, int* minor, hipDevice_t device);
+ */
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+
+#include "test_common.h"
+
+int main()
+{
+    int numDevices = 0;
+    int major,minor;
+    hipDevice_t device;
+    HIPCHECK(hipGetDeviceCount(&numDevices));
+    for(int i=0;i<numDevices;i++){
+        HIPCHECK(hipDeviceGet(&device,i));
+        HIPCHECK(hipDeviceComputeCapability(&major, &minor, device));
+        HIPASSERT(major >= 0);
+        HIPASSERT(minor >= 0);
+    }
+    passed();
+}

projects/hip/tests/src/runtimeApi/device/hipDeviceGetName.cpp

@@ -0,0 +1,47 @@
+/*
+Copyright (c) 2015-2017 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.
+*/
+
+/*
+ * Conformance test for checking functionality of
+ * hipError_t hipDeviceGetName(char* name, int len, hipDevice_t device);
+ */
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+
+#include "test_common.h"
+
+#define len 256
+
+int main()
+{
+    int numDevices = 0;
+    char name[len];
+    hipDevice_t device;
+    HIPCHECK(hipGetDeviceCount(&numDevices));
+    for(int i=0;i<numDevices;i++){
+        HIPCHECK(hipDeviceGet(&device,i));
+        HIPCHECK(hipDeviceGetName(name,len,device));
+        HIPASSERT(name != "");
+    }
+    passed();
+}

projects/hip/tests/src/runtimeApi/device/hipDeviceSynchronize.cpp

@@ -0,0 +1,74 @@
+/*
+Copyright (c) 2015-2017 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.
+*/
+
+/*
+ * Test for checking the functionality of
+ * hipError_t hipDeviceSynchronize();
+*/
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+
+#include"test_common.h"
+
+#define _SIZE sizeof(int)*1024*1024
+#define NUM_STREAMS 2
+
+__global__ void Iter(hipLaunchParm lp, int *Ad, int num){
+    int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    // Kernel loop designed to execute very slowly... ... ...   so we can test timing-related behavior below
+    if(tx == 0){
+        for(int i = 0; i<num;i++){
+            Ad[tx] += 1;
+        }
+    }
+}
+
+int main(){
+    int *A[NUM_STREAMS];
+    int *Ad[NUM_STREAMS];
+    hipStream_t stream[NUM_STREAMS];
+    for(int i=0;i<NUM_STREAMS;i++){
+        HIPCHECK(hipHostMalloc((void**)&A[i], _SIZE, hipHostMallocDefault));
+        A[i][0] = 1;
+        HIPCHECK(hipMalloc((void**)&Ad[i], _SIZE));
+        HIPCHECK(hipStreamCreate(&stream[i]));
+    }
+    for(int i=0;i<NUM_STREAMS;i++){
+        HIPCHECK(hipMemcpyAsync(Ad[i], A[i], _SIZE, hipMemcpyHostToDevice, stream[i]));
+    }
+    for(int i=0;i<NUM_STREAMS;i++){
+        hipLaunchKernel(HIP_KERNEL_NAME(Iter), dim3(1), dim3(1), 0, stream[i], Ad[i], 1<<30);
+    }
+    for(int i=0;i<NUM_STREAMS;i++){
+        HIPCHECK(hipMemcpyAsync(A[i], Ad[i], _SIZE, hipMemcpyDeviceToHost, stream[i]));
+    }
+
+
+   // This first check but relies on the kernel running for so long that the D2H async memcopy has not started yet.
+   // This will be true in an optimal asynchronous implementation.
+   // Conservative implementations which synchronize the hipMemcpyAsync will fail, ie if HIP_LAUNCH_BLOCKING=true
+    HIPASSERT(1<<30 != A[NUM_STREAMS-1][0]-1);
+    HIPCHECK(hipDeviceSynchronize());
+    HIPASSERT(1<<30 == A[NUM_STREAMS-1][0]-1);
+    passed();
+}

projects/hip/tests/src/runtimeApi/device/hipDeviceTotalMem.cpp

@@ -0,0 +1,45 @@
+/*
+Copyright (c) 2015-2017 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.
+*/
+
+/*
+ * Conformance test for checking functionality of
+ * hipError_t hipDeviceGetName(char* name, int len, hipDevice_t device);
+ */
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+
+#include "test_common.h"
+
+int main()
+{
+    int numDevices = 0;
+    size_t totMem;
+    hipDevice_t device;
+    HIPCHECK(hipGetDeviceCount(&numDevices));
+    for(int i=0;i<numDevices;i++){
+        HIPCHECK(hipDeviceGet(&device,i));
+        HIPCHECK(hipDeviceTotalMem(&totMem,device));
+        HIPASSERT(totMem != 0);
+    }
+    passed();
+}

projects/hip/tests/src/runtimeApi/device/hipGetDevice.cpp

@@ -0,0 +1,44 @@
+/*
+Copyright (c) 2015-2017 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.
+*/
+
+/*
+ * Conformance test for checking functionality of
+ * hipError_t hipGetDevice(int *device);
+ */
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+
+#include "test_common.h"
+
+int main()
+{
+    int numDevices = 0;
+    int device;
+    HIPCHECK(hipGetDeviceCount(&numDevices));
+    for(int i=0;i<numDevices;i++){
+        HIPCHECK(hipSetDevice(i));
+        HIPCHECK(hipGetDevice(&device));
+        HIPASSERT(device == i);
+    }
+    passed();
+}

projects/hip/tests/src/runtimeApi/device/hipRuntimeGetVersion.cpp

@@ -0,0 +1,39 @@
+/*
+Copyright (c) 2015-2017 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.
+*/
+
+/*
+ * Conformance test for checking functionality of
+ * hipError_t hipRuntimeGetVersion(int* runtimeVersion);
+ * On HIP/HCC path this function returns HIP runtime patch version(a 5 digit code) however on HIP/NVCC path this function return CUDA runtime version.
+ */
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+
+#include "test_common.h"
+
+int main()
+{
+    int runtimeVersion;
+    HIPCHECK(hipRuntimeGetVersion(&runtimeVersion));
+    passed();
+}

projects/hip/tests/src/runtimeApi/device/hipSetCachceConfig.cpp

@@ -0,0 +1,36 @@
+/*
+Copyright (c) 2015-Present 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.
+*/
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+
+#include<hip/hip_runtime_api.h>
+#include<iostream>
+#include"test_common.h"
+
+int main(){
+  hipFuncCache_t cacheConfig;
+  void *func;
+  hipFuncSetCacheConfig(func, cacheConfig);
+  passed();
+}
+

projects/hip/tests/src/runtimeApi/device/hipSetDevice.cpp

@@ -0,0 +1,36 @@
+/*
+Copyright (c) 2015-2017 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.
+*/
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * RUN: %t EXCLUDE_HIP_PLATFORM
+ * HIT_END
+ */
+
+#include "test_common.h"
+
+int main(){
+    int numDevices = 0;
+    HIPCHECK(hipGetDeviceCount(&numDevices));
+    for(int i=0;i<numDevices;i++){
+        HIPCHECK(hipSetDevice(i));
+    }
+    HIPASSERT(hipErrorInvalidDevice == hipSetDevice(numDevices));
+    passed();
+}

projects/hip/tests/src/runtimeApi/device/hipSetDeviceFlags.cpp

@@ -0,0 +1,52 @@
+/*
+Copyright (c) 2015-2017 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.
+*/
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+
+#include "test_common.h"
+
+int main()
+{
+    unsigned flag = 0;
+    HIPCHECK(hipDeviceReset());
+
+    int deviceCount = 0;
+    HIPCHECK(hipGetDeviceCount(&deviceCount));
+
+    for(int j=0;j<deviceCount;j++){
+
+        HIPCHECK(hipSetDevice(j));
+
+        for(int i=0;i<4;i++){
+            flag = 1 << i;
+            printf ("Flag=%x\n", flag);
+            HIPCHECK(hipSetDeviceFlags(flag));
+            //HIPCHECK_API(hipSetDeviceFlags(flag), hipErrorInvalidValue);
+        }
+
+        flag = 0;
+
+    }
+
+    passed();
+}