enable building cooperative directed tests on NV platform

Change-Id: I9a658b28010176d4f9401010b9c4b4be7ae86b92

hipamd/tests/src/runtimeApi/cooperativeGrps/api_failure_tests.cpp

@@ -42,7 +42,7 @@ runtime and that nothing breaks.*/
 
 
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD: %t %s ../../test_common.cpp
  * TEST: %t
  * HIT_END
  */

hipamd/tests/src/runtimeApi/cooperativeGrps/cooperative_streams.cpp

@@ -50,7 +50,7 @@ regular kernels can properly run in parallel with cooperative kernels.
 */
 
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11
  * TEST: %t
  * HIT_END
  */

hipamd/tests/src/runtimeApi/cooperativeGrps/grid_group_data_sharing.cpp

@@ -42,7 +42,7 @@ writing from thread X and then hitting a barrier allows thread Y to see the
 values.*/
 
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD: %t %s ../../test_common.cpp
  * TEST: %t
  * HIT_END
  */

hipamd/tests/src/runtimeApi/cooperativeGrps/hipLaunchCoopMultiKernel.cpp

@@ -20,7 +20,7 @@ THE SOFTWARE.
 // Simple test for hipLaunchCooperativeKernelMultiDevice API.
 
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 -rdc=true -gencode arch=compute_60,code=sm_60
  * TEST: %t
  * HIT_END
  */

hipamd/tests/src/runtimeApi/cooperativeGrps/hipLaunchCooperativeKernel.cpp

@@ -22,7 +22,7 @@ THE SOFTWARE.
 // Simple test for hipLaunchCooperativeKernel API.
 
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11
  * TEST: %t
  * HIT_END
  */
@@ -128,7 +128,7 @@ int main() {
     params[3] = (void*)&dC;
 
     std::cout << "Testing with grid size = " << dimGrid.x << " and block size = " << dimBlock.x << "\n";
-    HIPCHECK(hipLaunchCooperativeKernel(test_gws, dimGrid, dimBlock, params, dimBlock.x * sizeof(long), stream));
+    HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(test_gws), dimGrid, dimBlock, params, dimBlock.x * sizeof(long), stream));
 
     HIPCHECK(hipMemcpy(init, dC, sizeof(long), hipMemcpyDeviceToHost));
 

hipamd/tests/src/runtimeApi/cooperativeGrps/multi_gpu_api_failure_tests.cpp

@@ -37,7 +37,7 @@ should fail or succeed and compares this to what actually happens.
 */
 
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD: %t %s ../../test_common.cpp
  * TEST: %t
  * HIT_END
  */

hipamd/tests/src/runtimeApi/cooperativeGrps/multi_gpu_streams.cpp

@@ -80,7 +80,7 @@ overlap-allowing flags work as expected.
 */
 
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 -rdc=true -gencode arch=compute_60,code=sm_60
  * TEST: %t
  * HIT_END
  */

hipamd/tests/src/runtimeApi/cooperativeGrps/multi_grid_group_all_gpus.cpp

@@ -40,7 +40,7 @@ will sleep on the barrier waiting for the last warp to finally catch up.
 */
 
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -rdc=true -gencode arch=compute_60,code=sm_60
  * TEST: %t
  * HIT_END
  */

hipamd/tests/src/runtimeApi/cooperativeGrps/simple_grid_group_barrier.cpp

@@ -40,7 +40,7 @@ will sleep on the barrier waiting for the last warp to finally catch up.
 */
 
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD: %t %s ../../test_common.cpp
  * TEST: %t
  * HIT_END
  */

hipamd/tests/src/runtimeApi/cooperativeGrps/simple_multi_grid_group_barrier.cpp

@@ -40,7 +40,7 @@ will sleep on the barrier waiting for the last warp to finally catch up.
 */
 
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -rdc=true -gencode arch=compute_60,code=sm_60
  * TEST: %t
  * HIT_END
  */