changed __global__ attribute

1. Moved around tests and added them to HIT

Change-Id: I5d75280c42a5af852670ebabc7305ee56721ec7b


[ROCm/hip commit: 5e3d63c0a3]

projects/hip/include/hip/hcc_detail/hip_runtime.h

@@ -26,8 +26,8 @@ THE SOFTWARE.
  */
 
 //#pragma once
-#ifndef HIP_RUNTIME_H
-#define HIP_RUNTIME_H
+#ifndef HIP_HCC_DETAIL_RUNTIME_H
+#define HIP_HCC_DETAIL_RUNTIME_H
 
 //---
 // Top part of file can be compiled with any compiler
@@ -345,14 +345,31 @@ __device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask
 #define hipGridDim_y   (hc_get_num_groups(1))
 #define hipGridDim_z   (hc_get_num_groups(2))
 
-//extern "C" __device__ void* memcpy(void* dst, void* src, size_t size);
-//extern "C" __device__ void* memset(void* ptr, uint8_t val, size_t size);
-
+extern "C" __device__ void* __hip_hc_memcpy(void* dst, void* src, size_t size);
+extern "C" __device__ void* __hip_hc_memset(void* ptr, uint8_t val, size_t size);
 extern "C" __device__ void* __hip_hc_malloc(size_t);
 extern "C" __device__ void* __hip_hc_free(void *ptr);
 
-//extern "C" __device__ void* malloc(size_t size);
-//extern "C" __device__ void* free(void *ptr);
+static inline __device__ void* malloc(size_t size)
+{
+    return __hip_hc_malloc(size);
+}
+
+static inline __device__ void* free(void *ptr)
+{
+    return __hip_hc_free(ptr);
+}
+
+static inline __device__ void* memcpy(void* dst, void* src, size_t size)
+{
+  return __hip_hc_memcpy(dst, src, size);
+}
+
+static inline __device__ void* memset(void* ptr, uint8_t val, size_t size)
+{
+  return __hip_hc_memset(ptr, val, size);
+}
+
 
 
 #define __syncthreads() hc_barrier(CLK_LOCAL_MEM_FENCE)

projects/hip/include/hip/hcc_detail/host_defines.h

@@ -35,7 +35,7 @@ THE SOFTWARE.
 #define __host__     __attribute__((cpu))
 #define __device__   __attribute__((hc))
 
-#define __global__  __attribute__((hc_grid_launch))
+#define __global__  __attribute__((hc_grid_launch)) __attribute__((used))
 
 #define __noinline__      __attribute__((noinline))
 #define __forceinline__   __attribute__((always_inline))

projects/hip/src/device_util.cpp

@@ -99,7 +99,7 @@ __device__ void* __hip_hc_free(void *ptr)
 
 
 // loop unrolling
-__device__ void* memcpy(void* dst, void* src, size_t size)
+__device__ void* __hip_hc_memcpy(void* dst, void* src, size_t size)
 {
     uint8_t *dstPtr, *srcPtr;
     dstPtr = (uint8_t*)dst;
@@ -110,7 +110,7 @@ __device__ void* memcpy(void* dst, void* src, size_t size)
     return nullptr;
 }
 
-__device__ void* memset(void* ptr, uint8_t val, size_t size)
+__device__ void* __hip_hc_memset(void* ptr, uint8_t val, size_t size)
 {
     uint8_t *dstPtr;
     dstPtr = (uint8_t*)ptr;
@@ -120,16 +120,6 @@ __device__ void* memset(void* ptr, uint8_t val, size_t size)
     return nullptr;
 }
 
-__device__ void* malloc(size_t size)
-{
-    return __hip_hc_malloc(size);
-}
-
-__device__ void* free(void *ptr)
-{
-    return __hip_hc_free(ptr);
-}
-
 __device__ float __hip_erfinvf(float x){
     float ret;
     int  sign;

projects/hip/tests/src/Functional/device/hipFuncDeviceSynchronize.cpp

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+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
@@ -23,7 +23,7 @@ THE SOFTWARE.
 */
 
 /* HIT_START
- * BUILD: %t %s test_common.cpp
+ * BUILD: %t %s ../../test_common.cpp
  * RUN: %t
  * HIT_END
  */
@@ -63,13 +63,12 @@ int main(){
         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.  
+
+   // 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/Functional/device/hipFuncGetDevice.cpp

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+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
@@ -23,7 +23,7 @@ THE SOFTWARE.
  */
 
 /* HIT_START
- * BUILD: %t %s test_common.cpp
+ * BUILD: %t %s ../../test_common.cpp
  * RUN: %t
  * HIT_END
  */

projects/hip/tests/src/Functional/device/hipFuncSetDevice.cpp

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+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
@@ -18,7 +18,7 @@ THE SOFTWARE.
 */
 
 /* HIT_START
- * BUILD: %t %s test_common.cpp
+ * BUILD: %t %s ../../test_common.cpp
  * RUN: %t EXCLUDE_HIP_PLATFORM
  * HIT_END
  */

projects/hip/tests/src/Functional/device/hipFuncSetDeviceFlags.cpp

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+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
@@ -8,17 +8,17 @@ 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 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
+ * BUILD: %t %s ../../test_common.cpp
  * RUN: %t
  * HIT_END
  */

projects/hip/tests/src/buildHIPC.sh

@@ -1,3 +0,0 @@
-#!/bin/bash
-
-$HCC_HOME/bin/hcc -I$HCC_HOME/include -I$HSA_PATH/include -I$HIP_PATH/include -std=c11 -c hipC.c

projects/hip/tests/src/deviceLib/hipTestHalf.cpp

@@ -1,16 +1,13 @@
 /*
-Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
-
+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
@@ -20,113 +17,40 @@ 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"
-#include "hip/hip_runtime.h"
+#include <stdio.h>
+#include <hip/hip_fp16.h>
 #include "hip/hip_runtime_api.h"
-#include "hip/hip_fp16.h"
 
-#define hInf     0x7C00
-#define hInfPK   0x7C007C00
-#define h65504   0xF7FF
-#define h65504PK 0xF7FFF7FF
-#define h27      0x4EC0
-#define h27PK    0x4EC04EC0
-#define h7       0x4700
-#define h7PK     0x47004700
-#define h3       0x4200
-#define h3PK     0x42004200
-#define h1       0x3C00
-#define h1PK     0x3C003C00
-#define hPoint5     0x3800
-#define hPoint5PK   0x38003800
-#define hZero    0x0000
-#define hNeg1    0xBC00
-#define hNeg1PK 0xBC00BC00
+#define DSIZE 4
+#define SCF 0.5f
+#define nTPB 256
+__global__ void half_scale_kernel(hipLaunchParm lp, float *din, float *dout, int dsize){
 
-__global__ void CheckHalf(hipLaunchParm lp, __half* In1, __half* In2, __half* In3, __half* Out){
-  Out[0] = __hadd(In1[0], In2[0]);
-  Out[1] = __hadd_sat(In1[1], In2[1]);
-  Out[2] = __hfma(In1[2], In2[2],In3[2]);
-  Out[3] = __hfma_sat(In1[3], In2[3], In3[3]);
-  Out[4] = __hmul(In1[4], In2[4]);
-  Out[5] = __hmul_sat(In1[5], In2[5]);
-  Out[6] = __hneg(In1[6]);
-  Out[7] = __hsub(In1[7], In2[7]);
-  Out[8] = __hsub_sat(In1[8], In2[8]);
-  Out[9] = hdiv(In1[9], In2[9]);
-  Out[10] = hceil(In1[10]);
-  Out[11] = hcos(In1[11]);
-  Out[12] = hexp(In1[12]);
-  Out[13] = hexp10(In1[13]);
-  Out[14] = hexp2(In1[14]);
-  Out[15] = hfloor(In1[15]);
-  Out[16] = hlog(In1[16]);
-  Out[17] = hlog10(In1[17]);
-  Out[18] = hlog2(In1[18]);
-//  Out[19] = hrcp(In1[19]);
-  Out[20] = hrint(In1[20]);
-  Out[21] = hrsqrt(In1[21]);
-  Out[22] = hsin(In1[22]);
-  Out[23] = hsqrt(In1[23]);
-  Out[24] = htrunc(In1[24]);
-}
+  int idx = hipThreadIdx_x+ hipBlockDim_x*hipBlockIdx_x;
+  if (idx < dsize){
+    __half scf = __float2half(SCF);
+    __half kin = __float2half(din[idx]);
+    __half kout;
 
-__global__ void CheckHalf2(hipLaunchParm lp, __half2* In1, __half2* In2, __half2* In3, __half2* Out){
-  Out[0] = __hadd2(In1[0], In2[0]);
-  Out[1] = __hadd2_sat(In1[1], In2[1]);
-  Out[2] = __hfma2(In1[2], In2[2],In3[2]);
-  Out[3] = __hfma2_sat(In1[3], In2[3], In3[3]);
-  Out[4] = __hmul2(In1[4], In2[4]);
-  Out[5] = __hmul2_sat(In1[5], In2[5]);
-  Out[6] = __hneg2(In1[6]);
-  Out[7] = __hsub2(In1[7], In2[7]);
-  Out[8] = __hsub2_sat(In1[8], In2[8]);
-  Out[9] = h2div(In1[9], In2[9]);
-  Out[10] = h2ceil(In1[10]);
-  Out[11] = h2cos(In1[11]);
-  Out[12] = h2exp(In1[12]);
-  Out[13] = h2exp10(In1[13]);
-  Out[14] = h2exp2(In1[14]);
-  Out[15] = h2floor(In1[15]);
-  Out[16] = h2log(In1[16]);
-  Out[17] = h2log10(In1[17]);
-  Out[18] = h2log2(In1[18]);
-  Out[19] = h2rcp(In1[19]);
-//  Out[20] = h2rint(In1[20]);
-  Out[21] = h2rsqrt(In1[21]);
-  Out[22] = h2sin(In1[22]);
-  Out[23] = h2sqrt(In1[23]);
-  Out[24] = h2trunc(In1[24]);
-}
+    kout = __hmul(kin, scf);
 
-__global__ void CheckCmpHalf(hipLaunchParm lp, __half* In1, __half* In2, bool* Out) {
-  Out[0] = __heq(In1[0], In2[0]);
-  Out[1] = __hge(In1[1], In2[1]);
-  Out[2] = __hgt(In1[2], In2[2]);
-  Out[3] = __hisinf(In1[3]);
-  Out[4] = __hisnan(In1[4]);
-  Out[5] = __hle(In1[5], In2[5]);
-  Out[6] = __hlt(In1[6], In2[6]);
-  Out[7] = __hne(In1[7], In2[7]);
-}
-
-__global__ void CheckCmpHalf2(hipLaunchParm lp, __half2* In1, __half2* In2, __half2* Out) {
-  Out[0] = __heq2(In1[0], In2[0]);
-  Out[1] = __hge2(In1[1], In2[1]);
-  Out[2] = __hgt2(In1[2], In2[2]);
-  Out[4] = __hisnan2(In1[4]);
-  Out[5] = __hle2(In1[5], In2[5]);
-  Out[6] = __hlt2(In1[6], In2[6]);
-  Out[7] = __hne2(In1[7], In2[7]);
+//    kout = cvt_float_to_half(cvt_half_to_float(kin)*cvt_half_to_float(scf));
 
+    dout[idx] = __half2float(kout);
+    }
 }
 
 int main(){
-  passed();
+
+  float *hin, *hout, *din, *dout;
+  hin  = (float *)malloc(DSIZE*sizeof(float));
+  hout = (float *)malloc(DSIZE*sizeof(float));
+  for (int i = 0; i < DSIZE; i++) hin[i] = i;
+  hipMalloc(&din,  DSIZE*sizeof(float));
+  hipMalloc(&dout, DSIZE*sizeof(float));
+  hipMemcpy(din, hin, DSIZE*sizeof(float), hipMemcpyHostToDevice);
+  hipLaunchKernel(half_scale_kernel, dim3((DSIZE+nTPB-1)/nTPB),dim3(nTPB), 0, 0, din, dout, DSIZE);
+  hipMemcpy(hout, dout, DSIZE*sizeof(float), hipMemcpyDeviceToHost);
+  for (int i = 0; i < DSIZE; i++) printf("%f\n", hout[i]);
+  return 0;
 }

projects/hip/tests/src/hipC.c

@@ -1,4 +1,33 @@
+/*
+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
+ * RUN: %t
+ * HIT_END
+ */
+
 #include "hip/hip_runtime.h"
+#include "test_common.h"
 #include<stdio.h>
 
 #define ITER 1<<20
@@ -22,4 +51,5 @@ int main(){
     dimBlock.x = 1, dimBlock.y = 1, dimGrid.z = 1;
     hipLaunchKernel(HIP_KERNEL_NAME(Iter), dimGrid, dimBlock, 0, 0, Ad);
     hipMemcpy(&A, Ad, SIZE, hipMemcpyDeviceToHost);
+    passed();
 }

projects/hip/tests/src/hipC.cpp

@@ -1,5 +1,29 @@
+/*
+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.
+*/
+
+
 #include "hip/hip_runtime.h"
-#include<stdio.h>
+#include "test_common.h"
+#include <stdio.h>
 
 #define ITER 1<<20
 #define SIZE 1024*1024*sizeof(int)
@@ -19,4 +43,5 @@ int main(){
     hipMemcpy(Ad, &A, SIZE, hipMemcpyHostToDevice);
     hipLaunchKernel(HIP_KERNEL_NAME(Iter), dim3(1), dim3(1), 0, 0, Ad);
     hipMemcpy(&A, Ad, SIZE, hipMemcpyDeviceToHost);
+    passed();
 }

projects/hip/tests/src/hipCKernel.c

@@ -3,6 +3,7 @@
 
 __global__ void Kernel(hipLaunchParm lp, float *Ad){
     int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    Ad[tx] += Ad[tx-1];
 }
 
 int main(){

projects/hip/tests/src/hipChooseDevice.cpp

@@ -1,17 +0,0 @@
-#include <stdio.h>
-#include "hip/hip_runtime.h"
-int main( void ) {
-    hipDeviceProp_t  prop;
-    int dev;
-
-    hipGetDevice( &dev ) ;
-    printf( "ID of current HIP device:  %d\n", dev );
-
-    memset( &prop, 0, sizeof( hipDeviceProp_t ) );
-    prop.major = 1;
-    prop.minor = 3;
-    hipChooseDevice( &dev, &prop );
-    printf( "ID of hip device closest to revision 1.3:  %d\n", dev );
-
-    hipSetDevice( dev );
-}

projects/hip/tests/src/hipPeerToPeer_simple.cpp

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+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
@@ -50,7 +50,7 @@ void help(char *argv[])
 };
 
 
-static hipError_t myHipMemcpy(void *dest, const void *src, size_t sizeBytes, hipMemcpyKind kind,  hipStream_t stream, bool async) 
+static hipError_t myHipMemcpy(void *dest, const void *src, size_t sizeBytes, hipMemcpyKind kind,  hipStream_t stream, bool async)
 {
     if (async) {
 		hipError_t e = hipMemcpyAsync(dest, src, sizeBytes, kind, stream);
@@ -78,7 +78,7 @@ void parseMyArguments(int argc, char *argv[])
             p_mirrorPeers = true;
         } else if (!strcmp(arg, "--peerDevice")) {
             if (++i >= argc || !HipTest::parseInt(argv[i], &p_peerDevice)) {
-               failed("Bad peerDevice argument"); 
+               failed("Bad peerDevice argument");
             }
         } else {
             failed("Bad argument '%s'", arg);
@@ -101,7 +101,7 @@ void syncBothDevices()
 
 
 // Sets globals g_currentDevice, g_peerDevice
-void setupPeerTests() 
+void setupPeerTests()
 {
     int deviceCnt;
 
@@ -159,17 +159,17 @@ void enablePeerFirst(bool useAsyncCopy)
     // allocate and initialize memory on device0
     HIPCHECK (hipSetDevice(g_currentDevice));
     HIPCHECK (hipMalloc(&A_d0, Nbytes) );
-    HIPCHECK (hipMemset(A_d0, memsetval, Nbytes) ); 
+    HIPCHECK (hipMemset(A_d0, memsetval, Nbytes) );
 
     // allocate and initialize memory on peer device
     HIPCHECK (hipSetDevice(g_peerDevice));
     HIPCHECK (hipMalloc(&A_d1, Nbytes) );
-    HIPCHECK (hipMemset(A_d1, 0x13, Nbytes) ); 
+    HIPCHECK (hipMemset(A_d1, 0x13, Nbytes) );
 
 
 
     // Device0 push to device1, using P2P:
-    // NOTE : if p_mirrorPeers=0 and p_memcpyWithPeer=1, then peer device does not have mapping for A_d1 and we need to use a 
+    // NOTE : if p_mirrorPeers=0 and p_memcpyWithPeer=1, then peer device does not have mapping for A_d1 and we need to use a
     //        a host staging copy for the P2P access.
     HIPCHECK (hipSetDevice(p_memcpyWithPeer ? g_peerDevice : g_currentDevice));
     HIPCHECK (myHipMemcpy(A_d1, A_d0, Nbytes, hipMemcpyDefault, 0/*stream*/, useAsyncCopy)); // This is P2P copy.
@@ -177,7 +177,7 @@ void enablePeerFirst(bool useAsyncCopy)
     // Copy data back to host:
     // Have to wait for previous operation to finish, since we are switching to another one:
     HIPCHECK(hipDeviceSynchronize());
-    
+
     HIPCHECK (hipSetDevice(g_peerDevice));
     HIPCHECK (myHipMemcpy(A_h, A_d1, Nbytes, hipMemcpyDeviceToHost, 0/*stream*/, useAsyncCopy));
     HIPCHECK(hipDeviceSynchronize());
@@ -215,12 +215,12 @@ void allocMemoryFirst(bool useAsyncCopy)
     // allocate and initialize memory on device0
     HIPCHECK (hipSetDevice(g_currentDevice));
     HIPCHECK (hipMalloc(&A_d0, Nbytes) );
-    HIPCHECK ( hipMemset(A_d0, memsetval, Nbytes) ); 
+    HIPCHECK ( hipMemset(A_d0, memsetval, Nbytes) );
 
     // allocate and initialize memory on peer device
     HIPCHECK (hipSetDevice(g_peerDevice));
     HIPCHECK (hipMalloc(&A_d1, Nbytes) );
-    HIPCHECK ( hipMemset(A_d1, 0x13, Nbytes) ); 
+    HIPCHECK ( hipMemset(A_d1, 0x13, Nbytes) );
 
 
     //---
@@ -268,7 +268,7 @@ void allocMemoryFirst(bool useAsyncCopy)
 // Test which tests peer H2D copy - ie: copy-engine=1, dst=1, src=0 (Host)
 // A_d0 is pinned host on dev0 (this)
 // A_d1 is device memory on dev1 (peer)
-// 
+//
 void testPeerHostToDevice(bool useAsyncCopy)
 {
     printf ("\n==testing: %s useAsyncCopy=%d\n", __func__, useAsyncCopy);
@@ -299,12 +299,12 @@ void testPeerHostToDevice(bool useAsyncCopy)
     // allocate and initialize memory on device0
     HIPCHECK (hipSetDevice(g_currentDevice));
     HIPCHECK (hipHostMalloc(&A_host_d0, Nbytes) );
-    HIPCHECK (hipMemset(A_host_d0, memsetval, Nbytes) ); 
+    HIPCHECK (hipMemset(A_host_d0, memsetval, Nbytes) );
 
     // allocate and initialize memory on peer device
     HIPCHECK (hipSetDevice(g_peerDevice));
     HIPCHECK (hipMalloc(&A_d1, Nbytes) );
-    HIPCHECK (hipMemset(A_d1, 0x13, Nbytes) ); 
+    HIPCHECK (hipMemset(A_d1, 0x13, Nbytes) );
 
     bool firstAsyncCopy = useAsyncCopy; /*TODO - should be useAsyncCopy*/
 
@@ -313,17 +313,17 @@ void testPeerHostToDevice(bool useAsyncCopy)
 
 
     // Device0 push to device1, using P2P:
-    // NOTE : if p_mirrorPeers=0 and p_memcpyWithPeer=1, then peer device does not have mapping for A_d1 and we need to use a 
+    // NOTE : if p_mirrorPeers=0 and p_memcpyWithPeer=1, then peer device does not have mapping for A_d1 and we need to use a
     //        a host staging copy for the P2P access.
     if (p_memcpyWithPeer) {
         // p_memcpyWithPeer=1 case is HostToDevice.
-        // if p_mirrorPeers = 1, this is accelerated copy over PCIe.  
+        // if p_mirrorPeers = 1, this is accelerated copy over PCIe.
         // if p_mirrorPeers = 0, this should fall back to host (because peer can't see A_host_d0)
         HIPCHECK (hipSetDevice(g_peerDevice));
         HIPCHECK (myHipMemcpy(A_d1, A_host_d0, Nbytes, hipMemcpyHostToDevice, 0/*stream*/, firstAsyncCopy)); // This is P2P copy.
     } else {
         // p_memcpyWithPeer=0 case is HostToDevice.
-        // if p_mirrorPeers = 1, this is accelerated copy over PCIe.  
+        // if p_mirrorPeers = 1, this is accelerated copy over PCIe.
         // if p_mirrorPeers = 0, this should fall back to host (because device0 can't see A_d1)
         HIPCHECK (hipSetDevice(g_currentDevice));
         HIPCHECK (myHipMemcpy(A_d1, A_host_d0, Nbytes, hipMemcpyHostToDevice, 0/*stream*/, firstAsyncCopy)); // This is P2P copy.
@@ -367,7 +367,7 @@ void simpleNegative()
     HIPASSERT( e == hipSuccess);  // no error returned, it doesn't hurt to ask.
     HIPASSERT (canAccessPeer == 0); // but self is not a peer.
 
-    e = hipSuccess; 
+    e = hipSuccess;
     //---
     // Enable same device twice in a row:
     HIPCHECK(hipSetDevice(g_currentDevice));
@@ -381,7 +381,7 @@ void simpleNegative()
     e =(hipDeviceDisablePeerAccess(g_peerDevice));
     HIPASSERT (e == hipErrorPeerAccessNotEnabled);
 
-    
+
     // More tests here:
     printf ("==done: %s\n\n", __func__);
 }

projects/hip/tests/src/hipTestHalf.cpp

@@ -1,56 +0,0 @@
-/*
-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 <stdio.h>
-#include <hip/hip_fp16.h>
-#include "hip/hip_runtime_api.h"
-
-#define DSIZE 4
-#define SCF 0.5f
-#define nTPB 256
-__global__ void half_scale_kernel(hipLaunchParm lp, float *din, float *dout, int dsize){
-
-  int idx = hipThreadIdx_x+ hipBlockDim_x*hipBlockIdx_x;
-  if (idx < dsize){
-    __half scf = __float2half(SCF);
-    __half kin = __float2half(din[idx]);
-    __half kout;
-
-    kout = __hmul(kin, scf);
-
-//    kout = cvt_float_to_half(cvt_half_to_float(kin)*cvt_half_to_float(scf));
-
-    dout[idx] = __half2float(kout);
-    }
-}
-
-int main(){
-
-  float *hin, *hout, *din, *dout;
-  hin  = (float *)malloc(DSIZE*sizeof(float));
-  hout = (float *)malloc(DSIZE*sizeof(float));
-  for (int i = 0; i < DSIZE; i++) hin[i] = i;
-  hipMalloc(&din,  DSIZE*sizeof(float));
-  hipMalloc(&dout, DSIZE*sizeof(float));
-  hipMemcpy(din, hin, DSIZE*sizeof(float), hipMemcpyHostToDevice);
-  hipLaunchKernel(half_scale_kernel, dim3((DSIZE+nTPB-1)/nTPB),dim3(nTPB), 0, 0, din, dout, DSIZE);
-  hipMemcpy(hout, dout, DSIZE*sizeof(float), hipMemcpyDeviceToHost);
-  for (int i = 0; i < DSIZE; i++) printf("%f\n", hout[i]);
-  return 0;
-}

projects/hip/tests/src/kernel/hipDynamicShared.cpp

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+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
@@ -21,7 +21,7 @@ THE SOFTWARE.
 */
 
 /* HIT_START
- * BUILD: %t %s test_common.cpp
+ * BUILD: %t %s ../test_common.cpp
  * RUN: %t EXCLUDE_HIP_PLATFORM nvcc
  * HIT_END
  */
@@ -45,7 +45,7 @@ __global__ void testExternSharedKernel(hipLaunchParm lp, const T* A_d, const T*
 
     // initialize dynamic shared memory
     if (tid < groupElements) {
-        sdata[tid] = static_cast<T>(tid); 
+        sdata[tid] = static_cast<T>(tid);
     }
 
     // prefix sum inside dynamic shared memory
@@ -146,4 +146,3 @@ int main(int argc, char *argv[]) {
 
     passed();
 }
-

projects/hip/tests/src/kernel/hipEmptyKernel.cpp

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+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
@@ -17,6 +17,12 @@ 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"
 
 __global__ void Empty(hipLaunchParm lp, int param){}
@@ -24,4 +30,5 @@ __global__ void Empty(hipLaunchParm lp, int param){}
 int main(){
 hipLaunchKernel(HIP_KERNEL_NAME(Empty), dim3(1), dim3(1), 0, 0, 0);
 hipDeviceSynchronize();
+passed();
 }

projects/hip/tests/src/kernel/hipGridLaunch.cpp

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+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
@@ -33,7 +33,7 @@ THE SOFTWARE.
 
 
 // __device__ maps to __attribute__((hc))
-__device__ int foo(int i) 
+__device__ int foo(int i)
 {
     return i+1;
 }
@@ -96,4 +96,3 @@ int main(int argc, char *argv[])
 
     passed();
 }
-

projects/hip/tests/src/kernel/hipLanguageExtensions.cpp

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+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
@@ -42,7 +42,7 @@ __device__ int deviceVar;
 // TODO-HCC __constant__ not working yet.
 __constant__ int constantVar1;
 
-__constant__ __device__ int constantVar2; 
+__constant__ __device__ int constantVar2;
 
 // Test HOST space:
 __host__ void foo() {
@@ -53,7 +53,7 @@ __device__ __noinline__    int sum1_noinline(int a)    { return a+1;};
 __device__ __forceinline__ int sum1_forceinline(int a) { return a+1;};
 
 
-__device__ __host__ float PlusOne(float x) 
+__device__ __host__ float PlusOne(float x)
 {
     return x + 1.0;
 }

projects/hip/tests/src/kernel/hipLaunchParm.cpp

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+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
@@ -18,7 +18,7 @@ THE SOFTWARE.
 */
 
 /* HIT_START
- * BUILD: %t %s test_common.cpp
+ * BUILD: %t %s ../test_common.cpp
  * RUN: %t
  * HIT_END
  */

projects/hip/tests/src/kernel/hipTestConstant.cpp

@@ -17,6 +17,12 @@ 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.h>
 #include<hip/hip_runtime_api.h>
 #include<iostream>

projects/hip/tests/src/kernel/hipTestMallocKernel.cpp

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+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
@@ -17,6 +17,12 @@ 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.h>
 #include<hip/hip_runtime_api.h>
 #include<iostream>

projects/hip/tests/src/kernel/hipTestMemKernel.cpp

@@ -1,6 +1,32 @@
+/*
+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<hip/hip_runtime_api.h>
 #include<hip/hip_runtime.h>
 #include<iostream>
+#include"test_common.h"
 
 #define LEN8  8  * 4
 #define LEN9  9  * 4
@@ -184,4 +210,6 @@ int main(){
     delete A;
     delete B;
     delete C;
+
+    passed();
 }

projects/hip/tests/src/kernel/launch_bounds.cpp

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+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
@@ -20,12 +20,6 @@ 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
- */
-
 // Test launch bounds and initialization conditions.
 
 #include "hip/hip_runtime.h"
@@ -34,7 +28,7 @@ THE SOFTWARE.
 int p_blockSize = 256;
 
 
-__global__ 
+__global__
 void
 __launch_bounds__(256, 2)
 myKern(hipLaunchParm lp, int *C, const int *A, int N, int xfactor)

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

@@ -0,0 +1,49 @@
+/*
+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 <stdio.h>
+#include "hip/hip_runtime.h"
+#include "test_common.h"
+
+int main( void ) {
+    hipDeviceProp_t  prop;
+    int dev;
+
+    hipGetDevice( &dev ) ;
+    printf( "ID of current HIP device:  %d\n", dev );
+
+    memset( &prop, 0, sizeof( hipDeviceProp_t ) );
+    prop.major = 1;
+    prop.minor = 3;
+    hipChooseDevice( &dev, &prop );
+    printf( "ID of hip device closest to revision 1.3:  %d\n", dev );
+
+    hipSetDevice( dev );
+
+    passed();
+}

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

@@ -22,7 +22,7 @@ THE SOFTWARE.
 // Test the device info API extensions for HIP:
 
 /* HIT_START
- * BUILD: %t %s test_common.cpp
+ * BUILD: %t %s ../../test_common.cpp
  * RUN: %t EXCLUDE_HIP_PLATFORM nvcc
  * HIT_END
  */
@@ -89,4 +89,3 @@ int main(int argc, char *argv[])
     passed();
 
 };
-

projects/hip/tests/src/runtimeApi/event/hipEventRecord.cpp

@@ -20,11 +20,11 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
 // Test hipEventRecord serialization behavior.
-// Through manual inspection of the reported timestamps, can determine if recording a NULL event forces synchronization : 
-// set 
+// Through manual inspection of the reported timestamps, can determine if recording a NULL event forces synchronization :
+// set
 
 /* HIT_START
- * BUILD: %t %s test_common.cpp
+ * BUILD: %t %s ../../test_common.cpp
  * RUN: %t --iterations 10
  * HIT_END
  */
@@ -39,7 +39,7 @@ int main(int argc, char *argv[])
     unsigned blocks = (N+threadsPerBlock-1)/threadsPerBlock;
     if (blocks > 1024)
         blocks = 1024;
-    if (blocks ==0 ) 
+    if (blocks ==0 )
         blocks = 1;
 
 	printf ("N=%zu (A+B+C= %6.1f MB total) blocks=%u threadsPerBlock=%u iterations=%d\n", N, ((double)3*N*sizeof(float))/1024/1024, blocks, threadsPerBlock, iterations);
@@ -81,7 +81,7 @@ int main(int argc, char *argv[])
 		float eventMs = 1.0f;
 		HIPCHECK (hipEventElapsedTime(&eventMs, start, stop));
 		float hostMs = HipTest::elapsed_time(hostStart, hostStop);
-		
+
 		printf ("host_time (gettimeofday)          =%6.3fms\n", hostMs);
 		printf ("kernel_time (hipEventElapsedTime) =%6.3fms\n", eventMs);
 		printf ("\n");

projects/hip/tests/src/runtimeApi/memory/hipArray.cpp

@@ -1,5 +1,27 @@
+/*
+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 EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
  * RUN: %t EXCLUDE_HIP_PLATFORM
  * HIT_END
  */
@@ -170,7 +192,7 @@ void memcpyArraytest_size(size_t maxElem=0, size_t offset=0)
       maxElem = free/sizeof(T)/5;
   }
 
-  printf ("  device#%d: hipMemGetInfo: free=%zu (%4.2fMB) total=%zu (%4.2fMB)    maxSize=%6.1fMB offset=%lu\n", 
+  printf ("  device#%d: hipMemGetInfo: free=%zu (%4.2fMB) total=%zu (%4.2fMB)    maxSize=%6.1fMB offset=%lu\n",
           deviceId, free, (float)(free/1024.0/1024.0), total, (float)(total/1024.0/1024.0), maxElem*sizeof(T)/1024.0/1024.0, offset);
 
   // Test 1D

projects/hip/tests/src/runtimeApi/memory/hipHostGetFlags.cpp

@@ -21,7 +21,7 @@ THE SOFTWARE.
 */
 
 /* HIT_START
- * BUILD: %t %s test_common.cpp
+ * BUILD: %t %s ../../test_common.cpp
  * RUN: %t
  * HIT_END
  */

projects/hip/tests/src/runtimeApi/memory/hipHostMalloc.cpp

@@ -21,7 +21,7 @@
    */
 
 /* HIT_START
- * BUILD: %t %s test_common.cpp
+ * BUILD: %t %s ../../test_common.cpp
  * RUN: %t
  * HIT_END
  */

projects/hip/tests/src/runtimeApi/memory/hipHostRegister.cpp

@@ -18,7 +18,7 @@ THE SOFTWARE.
 */
 
 /* HIT_START
- * BUILD: %t %s test_common.cpp
+ * BUILD: %t %s ../../test_common.cpp
  * RUN: %t
  * HIT_END
  */

projects/hip/tests/src/runtimeApi/memory/hipMemcpyAsync2.cpp

@@ -1,3 +1,27 @@
+/*
+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.
+*/
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * HIT_END
+ */
+
 #include"test_common.h"
 
 #define SIZE 1024*1024

projects/hip/tests/src/runtimeApi/memory/hipRandomMemcpyAsync.cpp

@@ -18,7 +18,7 @@ THE SOFTWARE.
 */
 
 /* HIT_START
- * BUILD: %t %s test_common.cpp
+ * BUILD: %t %s ../../test_common.cpp
  * RUN: %t
  * HIT_END
  */

projects/hip/tests/src/runtimeApi/memory/hipTestMemcpyPin.cpp

@@ -18,7 +18,7 @@ THE SOFTWARE.
 */
 
 /* HIT_START
- * BUILD: %t %s test_common.cpp
+ * BUILD: %t %s ../../test_common.cpp
  * RUN: %t
  * HIT_END
  */

projects/hip/tests/src/sampleModule.cpp

@@ -1,94 +0,0 @@
-/*
-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 "hip/hip_runtime.h"
-#include "hip/hip_runtime_api.h"
-#include<iostream>
-#include<fstream>
-#include<vector>
-
-#define LEN 64
-#define SIZE LEN<<2
-
-#ifdef __HIP_PLATFORM_HCC__ 
-#define fileName "vcpy_isa.co"
-#endif
-
-#ifdef __HIP_PLATFORM_NVCC__
-#define fileName "vcpy_isa.ptx"
-#endif
-
-#define kernel_name "hello_world"
-
-int main(){
-    float *A, *B;
-    hipDeviceptr_t Ad, Bd;
-    A = new float[LEN];
-    B = new float[LEN];
-
-    for(uint32_t i=0;i<LEN;i++){
-        A[i] = i*1.0f;
-        B[i] = 0.0f;
-        std::cout<<A[i] << " "<<B[i]<<std::endl;
-    }
-
-
-#ifdef __HIP_PLATFORM_NVCC__
-  	hipInit(0);
-	  hipDevice_t device;
-	  hipCtx_t context;
-	  hipDeviceGet(&device, 0);
-	  hipCtxCreate(&context, 0, device);
-#endif
-
-    hipMalloc((void**)&Ad, SIZE);
-    hipMalloc((void**)&Bd, SIZE);
-
-    hipMemcpyHtoD(Ad, A, SIZE);
-    hipMemcpyHtoD(Bd, B, SIZE);
-    hipModule_t Module;
-    hipFunction_t Function;
-    hipModuleLoad(&Module, fileName);
-    hipModuleGetFunction(&Function, Module, kernel_name);
-
-    std::vector<void*>argBuffer(2);
-    memcpy(&argBuffer[0], &Ad, sizeof(void*));
-    memcpy(&argBuffer[1], &Bd, sizeof(void*));
-
-    size_t size = argBuffer.size()*sizeof(void*);
-
-    void *config[] = {
-      HIP_LAUNCH_PARAM_BUFFER_POINTER, &argBuffer[0],
-      HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
-      HIP_LAUNCH_PARAM_END
-    };
-
-    hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config); 
-
-    hipMemcpyDtoH(B, Bd, SIZE);
-    for(uint32_t i=0;i<LEN;i++){
-        std::cout<<A[i]<<" - "<<B[i]<<std::endl;
-    }
-
-#ifdef __HIP_PLATFORM_NVCC__
-	  hipCtxDetach(context);
-#endif 
-
-    return 0;
-}

projects/hip/tests/src/vcpy_isa.cu

@@ -1,6 +0,0 @@
-
-extern "C" __global__ void hello_world(float *a, float *b)
-{
-	int tx = threadIdx.x;
-	b[tx] = a[tx];
-}

projects/hip/tests/src/vcpy_isa.ptx

@@ -1,38 +0,0 @@
-//
-// Generated by NVIDIA NVVM Compiler
-//
-// Compiler Build ID: CL-19856038
-// Cuda compilation tools, release 7.5, V7.5.17
-// Based on LLVM 3.4svn
-//
-
-.version 4.3
-.target sm_20
-.address_size 64
-
-	// .globl	hello_world
-
-.visible .entry hello_world(
-	.param .u64 hello_world_param_0,
-	.param .u64 hello_world_param_1
-)
-{
-	.reg .f32 	%f<2>;
-	.reg .b32 	%r<2>;
-	.reg .b64 	%rd<8>;
-
-
-	ld.param.u64 	%rd1, [hello_world_param_0];
-	ld.param.u64 	%rd2, [hello_world_param_1];
-	cvta.to.global.u64 	%rd3, %rd2;
-	cvta.to.global.u64 	%rd4, %rd1;
-	mov.u32 	%r1, %tid.x;
-	mul.wide.s32 	%rd5, %r1, 4;
-	add.s64 	%rd6, %rd4, %rd5;
-	ld.global.f32 	%f1, [%rd6];
-	add.s64 	%rd7, %rd3, %rd5;
-	st.global.f32 	[%rd7], %f1;
-	ret;
-}
-
-