diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 822ee31f2b..aa6fe06337 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/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) diff --git a/include/hip/hcc_detail/host_defines.h b/include/hip/hcc_detail/host_defines.h index 906c39421e..e401cb24f3 100644 --- a/include/hip/hcc_detail/host_defines.h +++ b/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)) diff --git a/src/device_util.cpp b/src/device_util.cpp index d80d9e7ef5..fb207e3996 100644 --- a/src/device_util.cpp +++ b/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; diff --git a/tests/src/hipFuncDeviceSynchronize.cpp b/tests/src/Functional/device/hipFuncDeviceSynchronize.cpp similarity index 94% rename from tests/src/hipFuncDeviceSynchronize.cpp rename to tests/src/Functional/device/hipFuncDeviceSynchronize.cpp index 930bc37b8b..dac56bf709 100644 --- a/tests/src/hipFuncDeviceSynchronize.cpp +++ b/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(); } - diff --git a/tests/src/hipFuncGetDevice.cpp b/tests/src/Functional/device/hipFuncGetDevice.cpp similarity index 93% rename from tests/src/hipFuncGetDevice.cpp rename to tests/src/Functional/device/hipFuncGetDevice.cpp index f903149f91..a268ec0acd 100644 --- a/tests/src/hipFuncGetDevice.cpp +++ b/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 */ diff --git a/tests/src/hipFuncSetDevice.cpp b/tests/src/Functional/device/hipFuncSetDevice.cpp similarity index 93% rename from tests/src/hipFuncSetDevice.cpp rename to tests/src/Functional/device/hipFuncSetDevice.cpp index 030ceb7b0c..b1b7cac12d 100644 --- a/tests/src/hipFuncSetDevice.cpp +++ b/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 */ diff --git a/tests/src/hipFuncSetDeviceFlags.cpp b/tests/src/Functional/device/hipFuncSetDeviceFlags.cpp similarity index 67% rename from tests/src/hipFuncSetDeviceFlags.cpp rename to tests/src/Functional/device/hipFuncSetDeviceFlags.cpp index 9bda91f113..eacd4d2b63 100644 --- a/tests/src/hipFuncSetDeviceFlags.cpp +++ b/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 */ diff --git a/tests/src/Functional/Negative/Device/hipDeviceGetAttribute.cpp b/tests/src/Negative/Device/hipDeviceGetAttribute.cpp similarity index 100% rename from tests/src/Functional/Negative/Device/hipDeviceGetAttribute.cpp rename to tests/src/Negative/Device/hipDeviceGetAttribute.cpp diff --git a/tests/src/Functional/Negative/Device/hipDeviceUtil.h b/tests/src/Negative/Device/hipDeviceUtil.h similarity index 100% rename from tests/src/Functional/Negative/Device/hipDeviceUtil.h rename to tests/src/Negative/Device/hipDeviceUtil.h diff --git a/tests/src/Functional/Negative/Device/hipGetDevice.cpp b/tests/src/Negative/Device/hipGetDevice.cpp similarity index 100% rename from tests/src/Functional/Negative/Device/hipGetDevice.cpp rename to tests/src/Negative/Device/hipGetDevice.cpp diff --git a/tests/src/Functional/Negative/Device/hipGetDeviceCount.cpp b/tests/src/Negative/Device/hipGetDeviceCount.cpp similarity index 100% rename from tests/src/Functional/Negative/Device/hipGetDeviceCount.cpp rename to tests/src/Negative/Device/hipGetDeviceCount.cpp diff --git a/tests/src/Functional/Negative/Device/hipGetDeviceProperties.cpp b/tests/src/Negative/Device/hipGetDeviceProperties.cpp similarity index 100% rename from tests/src/Functional/Negative/Device/hipGetDeviceProperties.cpp rename to tests/src/Negative/Device/hipGetDeviceProperties.cpp diff --git a/tests/src/Functional/Negative/Device/hipSetDevice.cpp b/tests/src/Negative/Device/hipSetDevice.cpp similarity index 100% rename from tests/src/Functional/Negative/Device/hipSetDevice.cpp rename to tests/src/Negative/Device/hipSetDevice.cpp diff --git a/tests/src/hipPerfMemcpy.cpp b/tests/src/Performance/memory/hipPerfMemcpy.cpp similarity index 100% rename from tests/src/hipPerfMemcpy.cpp rename to tests/src/Performance/memory/hipPerfMemcpy.cpp diff --git a/tests/src/buildHIPC.sh b/tests/src/buildHIPC.sh deleted file mode 100755 index b2c9ce2a07..0000000000 --- a/tests/src/buildHIPC.sh +++ /dev/null @@ -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 diff --git a/tests/src/hipDrvGetPCIBusId.cpp b/tests/src/context/hipDrvGetPCIBusId.cpp similarity index 100% rename from tests/src/hipDrvGetPCIBusId.cpp rename to tests/src/context/hipDrvGetPCIBusId.cpp diff --git a/tests/src/hipDrvMemcpy.cpp b/tests/src/context/hipDrvMemcpy.cpp similarity index 100% rename from tests/src/hipDrvMemcpy.cpp rename to tests/src/context/hipDrvMemcpy.cpp diff --git a/tests/src/hipComplex.cpp b/tests/src/deviceLib/hipComplex.cpp similarity index 100% rename from tests/src/hipComplex.cpp rename to tests/src/deviceLib/hipComplex.cpp diff --git a/tests/src/hipDeviceMemcpy.cpp b/tests/src/deviceLib/hipDeviceMemcpy.cpp similarity index 100% rename from tests/src/hipDeviceMemcpy.cpp rename to tests/src/deviceLib/hipDeviceMemcpy.cpp diff --git a/tests/src/deviceLib/hipTestHalf.cpp b/tests/src/deviceLib/hipTestHalf.cpp index 12b7e2e270..7455037923 100644 --- a/tests/src/deviceLib/hipTestHalf.cpp +++ b/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 +#include #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; } diff --git a/tests/src/hipTestHost.cpp b/tests/src/deviceLib/hipTestHost.cpp similarity index 100% rename from tests/src/hipTestHost.cpp rename to tests/src/deviceLib/hipTestHost.cpp diff --git a/tests/src/hipC.c b/tests/src/hipC.c index 50177ac6c2..cdecc2cf9d 100644 --- a/tests/src/hipC.c +++ b/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 #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(); } diff --git a/tests/src/hipC.cpp b/tests/src/hipC.cpp index 3380e8abd9..8abb877808 100644 --- a/tests/src/hipC.cpp +++ b/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 +#include "test_common.h" +#include #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(); } diff --git a/tests/src/hipCKernel.c b/tests/src/hipCKernel.c index 19a034d843..7a72cf84ca 100644 --- a/tests/src/hipCKernel.c +++ b/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(){ diff --git a/tests/src/hipChooseDevice.cpp b/tests/src/hipChooseDevice.cpp deleted file mode 100644 index 4f289b9eb8..0000000000 --- a/tests/src/hipChooseDevice.cpp +++ /dev/null @@ -1,17 +0,0 @@ -#include -#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 ); -} diff --git a/tests/src/hipPeerToPeer_simple.cpp b/tests/src/hipPeerToPeer_simple.cpp index 1dfbdafdfc..1ea594f4bb 100644 --- a/tests/src/hipPeerToPeer_simple.cpp +++ b/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__); } diff --git a/tests/src/hipTestHalf.cpp b/tests/src/hipTestHalf.cpp deleted file mode 100644 index 7455037923..0000000000 --- a/tests/src/hipTestHalf.cpp +++ /dev/null @@ -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 -#include -#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; -} diff --git a/tests/src/hipDynamicShared.cpp b/tests/src/kernel/hipDynamicShared.cpp similarity index 97% rename from tests/src/hipDynamicShared.cpp rename to tests/src/kernel/hipDynamicShared.cpp index d5aed7b24f..522572a191 100644 --- a/tests/src/hipDynamicShared.cpp +++ b/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(tid); + sdata[tid] = static_cast(tid); } // prefix sum inside dynamic shared memory @@ -146,4 +146,3 @@ int main(int argc, char *argv[]) { passed(); } - diff --git a/tests/src/hipKernel.cpp b/tests/src/kernel/hipEmptyKernel.cpp similarity index 89% rename from tests/src/hipKernel.cpp rename to tests/src/kernel/hipEmptyKernel.cpp index bb6135ef5d..37245053be 100644 --- a/tests/src/hipKernel.cpp +++ b/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(); } diff --git a/tests/src/kernel/hipGridLaunch.cpp b/tests/src/kernel/hipGridLaunch.cpp index 99c6a29557..6cd724a070 100644 --- a/tests/src/kernel/hipGridLaunch.cpp +++ b/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(); } - diff --git a/tests/src/kernel/hipLanguageExtensions.cpp b/tests/src/kernel/hipLanguageExtensions.cpp index cc170f9c8a..e519a1c2a8 100644 --- a/tests/src/kernel/hipLanguageExtensions.cpp +++ b/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; } diff --git a/tests/src/hipLaunchParm.cpp b/tests/src/kernel/hipLaunchParm.cpp similarity index 96% rename from tests/src/hipLaunchParm.cpp rename to tests/src/kernel/hipLaunchParm.cpp index 5ffb557662..4dfd8a42ed 100644 --- a/tests/src/hipLaunchParm.cpp +++ b/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 */ diff --git a/tests/src/kernel/hipTestConstant.cpp b/tests/src/kernel/hipTestConstant.cpp index f86e8ace4f..2d637e3f1a 100644 --- a/tests/src/kernel/hipTestConstant.cpp +++ b/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 #include #include diff --git a/tests/src/kernel/hipTestMallocKernel.cpp b/tests/src/kernel/hipTestMallocKernel.cpp index 826f6164c3..bd0c17d898 100644 --- a/tests/src/kernel/hipTestMallocKernel.cpp +++ b/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 #include #include diff --git a/tests/src/kernel/hipTestMemKernel.cpp b/tests/src/kernel/hipTestMemKernel.cpp index bf97fc1dc8..9298e20e1b 100644 --- a/tests/src/kernel/hipTestMemKernel.cpp +++ b/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 #include #include +#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(); } diff --git a/tests/src/kernel/launch_bounds.cpp b/tests/src/kernel/launch_bounds.cpp index 2e10a9204a..3b1476fb11 100644 --- a/tests/src/kernel/launch_bounds.cpp +++ b/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) diff --git a/tests/src/runtimeApi/device/hipChooseDevice.cpp b/tests/src/runtimeApi/device/hipChooseDevice.cpp new file mode 100644 index 0000000000..4ddffe8aad --- /dev/null +++ b/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 +#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(); +} diff --git a/tests/src/hipGetDeviceAttribute.cpp b/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp similarity index 99% rename from tests/src/hipGetDeviceAttribute.cpp rename to tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp index a67296476f..8e55e2f699 100644 --- a/tests/src/hipGetDeviceAttribute.cpp +++ b/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(); }; - diff --git a/tests/src/hipEventRecord.cpp b/tests/src/runtimeApi/event/hipEventRecord.cpp similarity index 96% rename from tests/src/hipEventRecord.cpp rename to tests/src/runtimeApi/event/hipEventRecord.cpp index f8ef36b0bb..5606b4ab9b 100644 --- a/tests/src/hipEventRecord.cpp +++ b/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"); diff --git a/tests/src/hipArray.cpp b/tests/src/runtimeApi/memory/hipArray.cpp similarity index 87% rename from tests/src/hipArray.cpp rename to tests/src/runtimeApi/memory/hipArray.cpp index 9f9875a8d2..8f831bf5e0 100644 --- a/tests/src/hipArray.cpp +++ b/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 diff --git a/tests/src/hipHostGetFlags.cpp b/tests/src/runtimeApi/memory/hipHostGetFlags.cpp similarity index 98% rename from tests/src/hipHostGetFlags.cpp rename to tests/src/runtimeApi/memory/hipHostGetFlags.cpp index a58ba7aa43..a989b879ac 100644 --- a/tests/src/hipHostGetFlags.cpp +++ b/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 */ diff --git a/tests/src/hipHostMalloc.cpp b/tests/src/runtimeApi/memory/hipHostMalloc.cpp similarity index 98% rename from tests/src/hipHostMalloc.cpp rename to tests/src/runtimeApi/memory/hipHostMalloc.cpp index d1950f825f..d6b3b05a1d 100644 --- a/tests/src/hipHostMalloc.cpp +++ b/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 */ diff --git a/tests/src/hipHostRegister.cpp b/tests/src/runtimeApi/memory/hipHostRegister.cpp similarity index 98% rename from tests/src/hipHostRegister.cpp rename to tests/src/runtimeApi/memory/hipHostRegister.cpp index e9044d5a86..37ee9b1b78 100644 --- a/tests/src/hipHostRegister.cpp +++ b/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 */ diff --git a/tests/src/runtimeApi/memory/hipMemcpyAsync2.cpp b/tests/src/runtimeApi/memory/hipMemcpyAsync2.cpp index 4d7af21c72..7f19db559d 100644 --- a/tests/src/runtimeApi/memory/hipMemcpyAsync2.cpp +++ b/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 diff --git a/tests/src/hipRandomMemcpyAsync.cpp b/tests/src/runtimeApi/memory/hipRandomMemcpyAsync.cpp similarity index 98% rename from tests/src/hipRandomMemcpyAsync.cpp rename to tests/src/runtimeApi/memory/hipRandomMemcpyAsync.cpp index 8a5067ac8d..9a5636c24b 100644 --- a/tests/src/hipRandomMemcpyAsync.cpp +++ b/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 */ diff --git a/tests/src/hipTestMemcpyPin.cpp b/tests/src/runtimeApi/memory/hipTestMemcpyPin.cpp similarity index 97% rename from tests/src/hipTestMemcpyPin.cpp rename to tests/src/runtimeApi/memory/hipTestMemcpyPin.cpp index 3ac4e09765..c6634ebbbb 100644 --- a/tests/src/hipTestMemcpyPin.cpp +++ b/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 */ diff --git a/tests/src/hipModule.cpp b/tests/src/runtimeApi/module/hipModule.cpp similarity index 100% rename from tests/src/hipModule.cpp rename to tests/src/runtimeApi/module/hipModule.cpp diff --git a/tests/src/hipModuleUnload.cpp b/tests/src/runtimeApi/module/hipModuleUnload.cpp similarity index 100% rename from tests/src/hipModuleUnload.cpp rename to tests/src/runtimeApi/module/hipModuleUnload.cpp diff --git a/tests/src/sampleModule.cpp b/tests/src/sampleModule.cpp deleted file mode 100644 index 7ee9147e45..0000000000 --- a/tests/src/sampleModule.cpp +++ /dev/null @@ -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 -#include -#include - -#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;iargBuffer(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; - .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; -} - -