From 02d0e936012ef91fc0ed2173b753254f48f87b24 Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Tue, 24 Jul 2018 18:12:32 -0400 Subject: [PATCH 01/25] Support malloc/free for hip-clang --- docs/markdown/hip_programming_guide.md | 16 +- include/hip/hcc_detail/device_functions.h | 1 + include/hip/hcc_detail/hip_memory.h | 102 ++++++++++++ include/hip/hcc_detail/hip_runtime.h | 13 +- src/device_util.cpp | 64 -------- src/device_util.h | 11 -- src/hip_device.cpp | 2 +- tests/src/deviceLib/hipDeviceMalloc.cpp | 190 ++++++++++++++++++++++ 8 files changed, 317 insertions(+), 82 deletions(-) create mode 100644 include/hip/hcc_detail/hip_memory.h create mode 100644 tests/src/deviceLib/hipDeviceMalloc.cpp diff --git a/docs/markdown/hip_programming_guide.md b/docs/markdown/hip_programming_guide.md index 9313eb22e1..52d250cab5 100644 --- a/docs/markdown/hip_programming_guide.md +++ b/docs/markdown/hip_programming_guide.md @@ -91,8 +91,22 @@ Setting HCC_UNPINNED_COPY_MODE = 3, forces all unpinned transfer to use direct m Following environment variables can be used to control the transfer thresholds: -- HCC_H2D_STAGING_THRESHOLD - Threshold in KB for H2D copy. For sizes smaller than threshold direct copy logic would be used else staging buffers logic. By default it is set to 64. +- HCC_H2D_STAGING_THRESHOLD - Threshold in KB for H2D copy. For sizes smaller than threshold direct copy logic would be used else staging buffers logic. By default it is set to 64. - HCC_H2D_PININPLACE_THRESHOLD - Threshold in KB for H2D copy. For sizes smaller than threshold staging buffers logic would be used else PinInPlace logic. By default it is set to 4096. - HCC_D2H_PININPLACE_THRESHOLD - Threshold in KB for D2H copy. For sizes smaller than threshold staging buffer logic would be used else PinInPlace logic. By default it is set to 1024. + +## Device-Side Malloc + +hip-hcc and hip-clang supports device-side malloc and free. Users can allocate +memory dynamically in a kernel. The allocated memory are in global address +space, however, different threads get different memory allocations for the same +call of malloc. The allocated memory can be accessed or freed by other threads +or other kernels. It persists in the life time of the HIP program until it is +freed. + +The memory are allocated in pages. Users can define macro +`__HIP_SIZE_OF_PAGE` for controlling the page size in bytes and macro +`__HIP_NUM_PAGES` for controlling the total number of pages that can be +allocated. \ No newline at end of file diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index 6455fa6cd1..34a9a194e2 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/include/hip/hcc_detail/device_functions.h @@ -1029,4 +1029,5 @@ static inline __device__ void* memset(void* ptr, int val, size_t size) { unsigned char val8 = static_cast(val); return __hip_hc_memset(ptr, val8, size); } + #endif diff --git a/include/hip/hcc_detail/hip_memory.h b/include/hip/hcc_detail/hip_memory.h new file mode 100644 index 0000000000..9167baba38 --- /dev/null +++ b/include/hip/hcc_detail/hip_memory.h @@ -0,0 +1,102 @@ +/* +Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_MEMORY_H +#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_MEMORY_H + +// Implementation of malloc and free device functions. +// HIP heap is implemented as a global array with fixed size. Users may define +// __HIP_SIZE_OF_PAGE and __HIP_NUM_PAGES to have a larger heap. + +// Size of page in bytes. +#ifndef __HIP_SIZE_OF_PAGE +#define __HIP_SIZE_OF_PAGE 64 +#endif + +// Total number of pages +#ifndef __HIP_NUM_PAGES +#define __HIP_NUM_PAGES (16 * 64 * 64) +#endif + +#define __HIP_SIZE_OF_HEAP (__HIP_NUM_PAGES * __HIP_SIZE_OF_PAGE) + +__attribute__((weak)) __device__ char __hip_device_heap[__HIP_SIZE_OF_HEAP]; +__attribute__((weak)) __device__ + uint32_t __hip_device_page_flag[__HIP_NUM_PAGES]; + +extern "C" inline __device__ void* __hip_malloc(size_t size) { + char* heap = (char*)__hip_device_heap; + if (size > __HIP_SIZE_OF_HEAP) { + return (void*)nullptr; + } + uint32_t totalThreads = + hipBlockDim_x * hipGridDim_x * hipBlockDim_y + * hipGridDim_y * hipBlockDim_z * hipGridDim_z; + uint32_t currentWorkItem = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x; + + uint32_t numHeapsPerWorkItem = __HIP_NUM_PAGES / totalThreads; + uint32_t heapSizePerWorkItem = __HIP_SIZE_OF_HEAP / totalThreads; + + uint32_t stride = size / __HIP_SIZE_OF_PAGE; + uint32_t start = numHeapsPerWorkItem * currentWorkItem; + + uint32_t k = 0; + + while (__hip_device_page_flag[k] > 0) { + k++; + } + + for (uint32_t i = 0; i < stride - 1; i++) { + __hip_device_page_flag[i + start + k] = 1; + } + + __hip_device_page_flag[start + stride - 1 + k] = 2; + + void* ptr = (void*)(heap + + heapSizePerWorkItem * currentWorkItem + k * __HIP_SIZE_OF_PAGE); + + return ptr; +} + +extern "C" inline __device__ void* __hip_free(void* ptr) { + if (ptr == nullptr) { + return nullptr; + } + + uint32_t offsetByte = (uint64_t)ptr - (uint64_t)__hip_device_heap; + uint32_t offsetPage = offsetByte / __HIP_SIZE_OF_PAGE; + + while (__hip_device_page_flag[offsetPage] != 0) { + if (__hip_device_page_flag[offsetPage] == 2) { + __hip_device_page_flag[offsetPage] = 0; + offsetPage++; + break; + } else { + __hip_device_page_flag[offsetPage] = 0; + offsetPage++; + } + } + + return nullptr; +} + +#endif // HIP_INCLUDE_HIP_HCC_DETAIL_HIP_MEMORY_H diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index c2ae6e8e4f..8734feaf5d 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -260,11 +260,11 @@ static constexpr Coordinates threadIdx; #endif // defined __HCC__ #if __HCC_OR_HIP_CLANG__ -extern "C" __device__ void* __hip_hc_malloc(size_t); -extern "C" __device__ void* __hip_hc_free(void* ptr); +extern "C" __device__ void* __hip_malloc(size_t); +extern "C" __device__ void* __hip_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* malloc(size_t size) { return __hip_malloc(size); } +static inline __device__ void* free(void* ptr) { return __hip_free(ptr); } #ifdef __HCC_ACCELERATOR__ @@ -438,6 +438,8 @@ extern const __device__ __attribute__((weak)) __hip_builtin_gridDim_t gridDim; #define hipGridDim_y gridDim.y #define hipGridDim_z gridDim.z +#include + // Support std::complex. #pragma push_macro("__CUDA__") #define __CUDA__ @@ -448,8 +450,9 @@ extern const __device__ __attribute__((weak)) __hip_builtin_gridDim_t gridDim; #undef __CUDA__ #pragma pop_macro("__CUDA__") -#include #endif +#include + #endif // HIP_HCC_DETAIL_RUNTIME_H diff --git a/src/device_util.cpp b/src/device_util.cpp index 87fbe0fcbc..7fa77dc5fe 100644 --- a/src/device_util.cpp +++ b/src/device_util.cpp @@ -28,70 +28,6 @@ THE SOFTWARE. #include "hip/hip_runtime.h" #include -//================================================================================================= -/* - Implementation of malloc and free device functions. - - This is the best place to put them because the device - global variables need to be initialized at the start. -*/ -__device__ char gpuHeap[SIZE_OF_HEAP]; -__device__ uint32_t gpuFlags[NUM_PAGES]; - -__device__ void* __hip_hc_malloc(size_t size) { - char* heap = (char*)gpuHeap; - if (size > SIZE_OF_HEAP) { - return (void*)nullptr; - } - uint32_t totalThreads = - blockDim.x * gridDim.x * blockDim.y * gridDim.y * blockDim.z * gridDim.z; - uint32_t currentWorkItem = threadIdx.x + blockDim.x * blockIdx.x; - - uint32_t numHeapsPerWorkItem = NUM_PAGES / totalThreads; - uint32_t heapSizePerWorkItem = SIZE_OF_HEAP / totalThreads; - - uint32_t stride = size / SIZE_OF_PAGE; - uint32_t start = numHeapsPerWorkItem * currentWorkItem; - - uint32_t k = 0; - - while (gpuFlags[k] > 0) { - k++; - } - - for (uint32_t i = 0; i < stride - 1; i++) { - gpuFlags[i + start + k] = 1; - } - - gpuFlags[start + stride - 1 + k] = 2; - - void* ptr = (void*)(heap + heapSizePerWorkItem * currentWorkItem + k * SIZE_OF_PAGE); - - return ptr; -} - -__device__ void* __hip_hc_free(void* ptr) { - if (ptr == nullptr) { - return nullptr; - } - - uint32_t offsetByte = (uint64_t)ptr - (uint64_t)gpuHeap; - uint32_t offsetPage = offsetByte / SIZE_OF_PAGE; - - while (gpuFlags[offsetPage] != 0) { - if (gpuFlags[offsetPage] == 2) { - gpuFlags[offsetPage] = 0; - offsetPage++; - break; - } else { - gpuFlags[offsetPage] = 0; - offsetPage++; - } - } - - return nullptr; -} - // abort __device__ void abort() { return hc::abort(); } diff --git a/src/device_util.h b/src/device_util.h index 8fa96da9d9..84dbbf71ed 100644 --- a/src/device_util.h +++ b/src/device_util.h @@ -29,14 +29,6 @@ THE SOFTWARE. Heap size computation for malloc and free device functions. */ -#define NUM_PAGES_PER_THREAD 16 -#define SIZE_OF_PAGE 64 -#define NUM_THREADS_PER_CU 64 -#define NUM_CUS_PER_GPU 64 // Specific for r9 Nano -#define NUM_PAGES NUM_PAGES_PER_THREAD* NUM_THREADS_PER_CU* NUM_CUS_PER_GPU -#define SIZE_MALLOC NUM_PAGES* SIZE_OF_PAGE -#define SIZE_OF_HEAP SIZE_MALLOC - #define HIP_SQRT_2 1.41421356237 #define HIP_SQRT_PI 1.77245385091 @@ -62,9 +54,6 @@ THE SOFTWARE. #define HIP_PI 3.14159265358979323846 -__device__ void* __hip_hc_malloc(size_t size); -__device__ void* __hip_hc_free(void* ptr); - __device__ float __hip_erfinvf(float x); __device__ double __hip_erfinv(double x); diff --git a/src/hip_device.cpp b/src/hip_device.cpp index 72150c3f54..2aae7cf2a8 100644 --- a/src/hip_device.cpp +++ b/src/hip_device.cpp @@ -99,7 +99,7 @@ hipError_t hipDeviceGetLimit(size_t* pValue, hipLimit_t limit) { return ihipLogStatus(hipErrorInvalidValue); } if (limit == hipLimitMallocHeapSize) { - *pValue = (size_t)SIZE_OF_HEAP; + *pValue = (size_t)__HIP_SIZE_OF_HEAP; return ihipLogStatus(hipSuccess); } else { return ihipLogStatus(hipErrorUnsupportedLimit); diff --git a/tests/src/deviceLib/hipDeviceMalloc.cpp b/tests/src/deviceLib/hipDeviceMalloc.cpp new file mode 100644 index 0000000000..8eb8cdcc3c --- /dev/null +++ b/tests/src/deviceLib/hipDeviceMalloc.cpp @@ -0,0 +1,190 @@ +/* +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. +*/ +/* HIT_START + * BUILD: %t %s NVCC_OPTIONS -std=c++11 + * RUN: %t EXCLUDE_HIP_PLATFORM nvcc + * HIT_END + */ +#include "test_common.h" +#include +#include + +// Tolerance for error +const double tolerance = 1e-6; +const bool verbose = false; + +#define LEN 64 + +#define ALL_FUN \ + OP(add) \ + OP(sub) \ + OP(mul) \ + OP(div) + +#define OP(x) CK_##x, +enum CalcKind { + ALL_FUN +}; +#undef OP + +#define OP(x) case CK_##x: return #x; +std::string getName(enum CalcKind CK) { + switch(CK){ + ALL_FUN + } +} +#undef OP + +// Calculates function. +// If the function has one argument, B is ignored. +// If the function returns real number, converts it to a complex number. +#define ONE_ARG(func) \ + case CK_##func: \ + return std::complex(std::func(A)); + +template +__device__ __host__ std::complex calc(std::complex A, + std::complex B, + enum CalcKind CK) { + switch(CK) { + case CK_add: + return A + B; + case CK_sub: + return A - B; + case CK_mul: + return A * B; + case CK_div: + return A / B; + + } +} + +// Allocate memory in kernel and save the address to pA and pB. +// Copy value from A, B to allocated memory. +template +__global__ void kernel_alloc(std::complex* A, + std::complex* B, + std::complex** pA, + std::complex** pB) { + typedef std::complex CFloatT; + int tx = threadIdx.x + blockIdx.x * blockDim.x; + if (tx == 0) { + *pA = (CFloatT*)malloc(sizeof(CFloatT)*LEN); + *pB = (CFloatT*)malloc(sizeof(CFloatT)*LEN); + for (int i = 0; i < LEN; i++) { + (*pA)[i] = A[i]; + (*pB)[i] = B[i]; + } + } +} + +// Do calculation using values saved in allocated memmory. pA, pB are buffers +// containing the address of the device-side allocated array. +template +__global__ void kernel_free(std::complex** pA, + std::complex** pB, std::complex* C, + enum CalcKind CK) { + typedef std::complex CFloatT; + int tx = threadIdx.x + blockIdx.x * blockDim.x; + C[tx] = calc((*pA)[tx], (*pB)[tx], CK); + if (tx == 0) { + free(*pA); + free(*pB); + } +} + +template +void test() { + typedef std::complex ComplexT; + ComplexT *A, *Ad, *B, *Bd, *C, *Cd, *D; + A = new ComplexT[LEN]; + B = new ComplexT[LEN]; + C = new ComplexT[LEN]; + D = new ComplexT[LEN]; + hipMalloc((void**)&Ad, sizeof(ComplexT)*LEN); + hipMalloc((void**)&Bd, sizeof(ComplexT)*LEN); + hipMalloc((void**)&Cd, sizeof(ComplexT)*LEN); + + for (uint32_t i = 0; i < LEN; i++) { + A[i] = ComplexT((i + 1) * 1.0f, (i + 2) * 1.0f); + B[i] = A[i]; + C[i] = A[i]; + } + hipMemcpy(Ad, A, sizeof(ComplexT)*LEN, hipMemcpyHostToDevice); + hipMemcpy(Bd, B, sizeof(ComplexT)*LEN, hipMemcpyHostToDevice); + + // Run kernel for a calculation kind and verify by comparing with host + // calculation result. Returns false if fails. + auto test_fun = [&](enum CalcKind CK) { + // kernel_alloc allocates memory on device side and initialize it. + // kernel_free uses allocated memory from kernel_alloc and does the + // calculation then free the memory. + // pA and pB are buffers to pass the device-side allocated memory address + // from kernel_alloc to kernel_free. + ComplexT **pA, **pB; + hipMalloc((ComplexT***)&pA, sizeof(ComplexT*)); + hipMalloc((ComplexT***)&pB, sizeof(ComplexT*)); + hipLaunchKernelGGL(kernel_alloc, dim3(1), dim3(LEN), 0, 0, + Ad, Bd, pA, pB); + hipDeviceSynchronize(); + hipLaunchKernelGGL(kernel_free, dim3(1), dim3(LEN), 0, 0, + pA, pB, Cd, CK); + hipMemcpy(C, Cd, sizeof(ComplexT)*LEN, hipMemcpyDeviceToHost); + hipFree(pA); + hipFree(pB); + for (int i = 0; i < LEN; i++) { + ComplexT Expected = calc(A[i], B[i], CK); + FloatT error = std::abs(C[i] - Expected); + if (std::abs(Expected) > tolerance) + error /= std::abs(Expected); + bool pass = error < tolerance; + if (verbose || !pass) { + std::cout << "Function: " << getName(CK) + << " Operands: " << A[i] << " " << B[i] + << " Result: " << C[i] + << " Expected: " << Expected + << " Error: " << error + << " Pass: " << pass + << std::endl; + } + if (!pass) + return false; + } + return true; + }; + +#define OP(x) assert(test_fun(CK_##x)); + ALL_FUN +#undef OP + + hipFree(Ad); + hipFree(Bd); + hipFree(Cd); + delete[] A; + delete[] B; + delete[] C; + delete[] D; +} + +int main() { + test(); + test(); + passed(); + return 0; +} From 03320890de9c1499737a66662af0df1bb4cf01ba Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Wed, 25 Jul 2018 16:56:39 -0400 Subject: [PATCH 02/25] Fix thread index calculation in __hip_malloc --- include/hip/hcc_detail/hip_memory.h | 5 ++++- tests/src/deviceLib/hipDeviceMalloc.cpp | 25 ++++++++++++++++++++----- 2 files changed, 24 insertions(+), 6 deletions(-) diff --git a/include/hip/hcc_detail/hip_memory.h b/include/hip/hcc_detail/hip_memory.h index 9167baba38..2394a05d0f 100644 --- a/include/hip/hcc_detail/hip_memory.h +++ b/include/hip/hcc_detail/hip_memory.h @@ -51,7 +51,10 @@ extern "C" inline __device__ void* __hip_malloc(size_t size) { uint32_t totalThreads = hipBlockDim_x * hipGridDim_x * hipBlockDim_y * hipGridDim_y * hipBlockDim_z * hipGridDim_z; - uint32_t currentWorkItem = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x; + uint32_t currentWorkItem = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x + + (hipThreadIdx_y + hipBlockDim_y * hipBlockIdx_y) * hipBlockDim_x + + (hipThreadIdx_z + hipBlockDim_z * hipBlockIdx_z) * hipBlockDim_x + * hipBlockDim_y; uint32_t numHeapsPerWorkItem = __HIP_NUM_PAGES / totalThreads; uint32_t heapSizePerWorkItem = __HIP_SIZE_OF_HEAP / totalThreads; diff --git a/tests/src/deviceLib/hipDeviceMalloc.cpp b/tests/src/deviceLib/hipDeviceMalloc.cpp index 8eb8cdcc3c..a9d62db025 100644 --- a/tests/src/deviceLib/hipDeviceMalloc.cpp +++ b/tests/src/deviceLib/hipDeviceMalloc.cpp @@ -29,7 +29,14 @@ THE SOFTWARE. const double tolerance = 1e-6; const bool verbose = false; -#define LEN 64 +#define BLKDIM_X 64 +#define BLKDIM_Y 1 +#define BLKDIM_Z 1 +#define NUM_BLK_X 1 +#define NUM_BLK_Y 1 +#define NUM_BLK_Z 1 + +#define LEN (BLKDIM_X * BLKDIM_Y * BLKDIM_Z * NUM_BLK_X * NUM_BLK_Y * NUM_BLK_Z) #define ALL_FUN \ OP(add) \ @@ -83,7 +90,10 @@ __global__ void kernel_alloc(std::complex* A, std::complex** pA, std::complex** pB) { typedef std::complex CFloatT; - int tx = threadIdx.x + blockIdx.x * blockDim.x; + int tx = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x + + (hipThreadIdx_y + hipBlockDim_y * hipBlockIdx_y) * hipBlockDim_x + + (hipThreadIdx_z + hipBlockDim_z * hipBlockIdx_z) * hipBlockDim_x + * hipBlockDim_y; if (tx == 0) { *pA = (CFloatT*)malloc(sizeof(CFloatT)*LEN); *pB = (CFloatT*)malloc(sizeof(CFloatT)*LEN); @@ -101,7 +111,10 @@ __global__ void kernel_free(std::complex** pA, std::complex** pB, std::complex* C, enum CalcKind CK) { typedef std::complex CFloatT; - int tx = threadIdx.x + blockIdx.x * blockDim.x; + int tx = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x + + (hipThreadIdx_y + hipBlockDim_y * hipBlockIdx_y) * hipBlockDim_x + + (hipThreadIdx_z + hipBlockDim_z * hipBlockIdx_z) * hipBlockDim_x + * hipBlockDim_y; C[tx] = calc((*pA)[tx], (*pB)[tx], CK); if (tx == 0) { free(*pA); @@ -140,10 +153,12 @@ void test() { ComplexT **pA, **pB; hipMalloc((ComplexT***)&pA, sizeof(ComplexT*)); hipMalloc((ComplexT***)&pB, sizeof(ComplexT*)); - hipLaunchKernelGGL(kernel_alloc, dim3(1), dim3(LEN), 0, 0, + dim3 blkDim(BLKDIM_X, BLKDIM_Y, BLKDIM_Z); + dim3 numBlk(NUM_BLK_X, NUM_BLK_Y, NUM_BLK_Z); + hipLaunchKernelGGL(kernel_alloc, numBlk, blkDim, 0, 0, Ad, Bd, pA, pB); hipDeviceSynchronize(); - hipLaunchKernelGGL(kernel_free, dim3(1), dim3(LEN), 0, 0, + hipLaunchKernelGGL(kernel_free, numBlk, blkDim, 0, 0, pA, pB, Cd, CK); hipMemcpy(C, Cd, sizeof(ComplexT)*LEN, hipMemcpyDeviceToHost); hipFree(pA); From f06894e2f0abb89c852ff6c3bdf27b9758cc4922 Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Fri, 27 Jul 2018 17:07:00 -0400 Subject: [PATCH 03/25] Do not use std::complex in test hipDeviceMalloc --- tests/src/deviceLib/hipDeviceMalloc.cpp | 105 ++++++++++-------------- 1 file changed, 44 insertions(+), 61 deletions(-) diff --git a/tests/src/deviceLib/hipDeviceMalloc.cpp b/tests/src/deviceLib/hipDeviceMalloc.cpp index a9d62db025..4ec10077c5 100644 --- a/tests/src/deviceLib/hipDeviceMalloc.cpp +++ b/tests/src/deviceLib/hipDeviceMalloc.cpp @@ -60,43 +60,35 @@ std::string getName(enum CalcKind CK) { // Calculates function. // If the function has one argument, B is ignored. -// If the function returns real number, converts it to a complex number. -#define ONE_ARG(func) \ - case CK_##func: \ - return std::complex(std::func(A)); +#define ONE_ARG(func) \ + case CK_##func: \ + return std::func(A); -template -__device__ __host__ std::complex calc(std::complex A, - std::complex B, - enum CalcKind CK) { - switch(CK) { - case CK_add: - return A + B; - case CK_sub: - return A - B; - case CK_mul: - return A * B; - case CK_div: - return A / B; - - } +template +__device__ __host__ FloatT calc(FloatT A, FloatT B, enum CalcKind CK) { + switch (CK) { + case CK_add: + return A + B; + case CK_sub: + return A - B; + case CK_mul: + return A * B; + case CK_div: + return A / B; + } } // Allocate memory in kernel and save the address to pA and pB. // Copy value from A, B to allocated memory. -template -__global__ void kernel_alloc(std::complex* A, - std::complex* B, - std::complex** pA, - std::complex** pB) { - typedef std::complex CFloatT; +template +__global__ void kernel_alloc(FloatT* A, FloatT* B, FloatT** pA, FloatT** pB) { int tx = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x + (hipThreadIdx_y + hipBlockDim_y * hipBlockIdx_y) * hipBlockDim_x + (hipThreadIdx_z + hipBlockDim_z * hipBlockIdx_z) * hipBlockDim_x * hipBlockDim_y; if (tx == 0) { - *pA = (CFloatT*)malloc(sizeof(CFloatT)*LEN); - *pB = (CFloatT*)malloc(sizeof(CFloatT)*LEN); + *pA = (FloatT*)malloc(sizeof(FloatT) * LEN); + *pB = (FloatT*)malloc(sizeof(FloatT) * LEN); for (int i = 0; i < LEN; i++) { (*pA)[i] = A[i]; (*pB)[i] = B[i]; @@ -106,11 +98,8 @@ __global__ void kernel_alloc(std::complex* A, // Do calculation using values saved in allocated memmory. pA, pB are buffers // containing the address of the device-side allocated array. -template -__global__ void kernel_free(std::complex** pA, - std::complex** pB, std::complex* C, - enum CalcKind CK) { - typedef std::complex CFloatT; +template +__global__ void kernel_free(FloatT** pA, FloatT** pB, FloatT* C, enum CalcKind CK) { int tx = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x + (hipThreadIdx_y + hipBlockDim_y * hipBlockIdx_y) * hipBlockDim_x + (hipThreadIdx_z + hipBlockDim_z * hipBlockIdx_z) * hipBlockDim_x @@ -124,23 +113,22 @@ __global__ void kernel_free(std::complex** pA, template void test() { - typedef std::complex ComplexT; - ComplexT *A, *Ad, *B, *Bd, *C, *Cd, *D; - A = new ComplexT[LEN]; - B = new ComplexT[LEN]; - C = new ComplexT[LEN]; - D = new ComplexT[LEN]; - hipMalloc((void**)&Ad, sizeof(ComplexT)*LEN); - hipMalloc((void**)&Bd, sizeof(ComplexT)*LEN); - hipMalloc((void**)&Cd, sizeof(ComplexT)*LEN); + FloatT *A, *Ad, *B, *Bd, *C, *Cd, *D; + A = new FloatT[LEN]; + B = new FloatT[LEN]; + C = new FloatT[LEN]; + D = new FloatT[LEN]; + hipMalloc((void**)&Ad, sizeof(FloatT) * LEN); + hipMalloc((void**)&Bd, sizeof(FloatT) * LEN); + hipMalloc((void**)&Cd, sizeof(FloatT) * LEN); for (uint32_t i = 0; i < LEN; i++) { - A[i] = ComplexT((i + 1) * 1.0f, (i + 2) * 1.0f); + A[i] = (i + 1) * 1.0f; B[i] = A[i]; C[i] = A[i]; } - hipMemcpy(Ad, A, sizeof(ComplexT)*LEN, hipMemcpyHostToDevice); - hipMemcpy(Bd, B, sizeof(ComplexT)*LEN, hipMemcpyHostToDevice); + hipMemcpy(Ad, A, sizeof(FloatT) * LEN, hipMemcpyHostToDevice); + hipMemcpy(Bd, B, sizeof(FloatT) * LEN, hipMemcpyHostToDevice); // Run kernel for a calculation kind and verify by comparing with host // calculation result. Returns false if fails. @@ -150,9 +138,9 @@ void test() { // calculation then free the memory. // pA and pB are buffers to pass the device-side allocated memory address // from kernel_alloc to kernel_free. - ComplexT **pA, **pB; - hipMalloc((ComplexT***)&pA, sizeof(ComplexT*)); - hipMalloc((ComplexT***)&pB, sizeof(ComplexT*)); + FloatT **pA, **pB; + hipMalloc((FloatT***)&pA, sizeof(FloatT*)); + hipMalloc((FloatT***)&pB, sizeof(FloatT*)); dim3 blkDim(BLKDIM_X, BLKDIM_Y, BLKDIM_Z); dim3 numBlk(NUM_BLK_X, NUM_BLK_Y, NUM_BLK_Z); hipLaunchKernelGGL(kernel_alloc, numBlk, blkDim, 0, 0, @@ -160,23 +148,18 @@ void test() { hipDeviceSynchronize(); hipLaunchKernelGGL(kernel_free, numBlk, blkDim, 0, 0, pA, pB, Cd, CK); - hipMemcpy(C, Cd, sizeof(ComplexT)*LEN, hipMemcpyDeviceToHost); + hipMemcpy(C, Cd, sizeof(FloatT) * LEN, hipMemcpyDeviceToHost); hipFree(pA); hipFree(pB); for (int i = 0; i < LEN; i++) { - ComplexT Expected = calc(A[i], B[i], CK); - FloatT error = std::abs(C[i] - Expected); - if (std::abs(Expected) > tolerance) - error /= std::abs(Expected); - bool pass = error < tolerance; - if (verbose || !pass) { - std::cout << "Function: " << getName(CK) - << " Operands: " << A[i] << " " << B[i] - << " Result: " << C[i] - << " Expected: " << Expected - << " Error: " << error - << " Pass: " << pass - << std::endl; + FloatT Expected = calc(A[i], B[i], CK); + FloatT error = std::abs(C[i] - Expected); + if (std::abs(Expected) > tolerance) error /= std::abs(Expected); + bool pass = error < tolerance; + if (verbose || !pass) { + std::cout << "Function: " << getName(CK) << " Operands: " << A[i] << " " << B[i] + << " Result: " << C[i] << " Expected: " << Expected << " Error: " << error + << " Pass: " << pass << std::endl; } if (!pass) return false; From a6c7aeed7237b71122c85c18663f8d740902e9ea Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Sat, 28 Jul 2018 09:02:38 -0400 Subject: [PATCH 04/25] Add HIP directed test hipTestGlobalVariable.cpp --- tests/src/kernel/hipTestGlobalVariable.cpp | 97 ++++++++++++++++++++++ 1 file changed, 97 insertions(+) create mode 100644 tests/src/kernel/hipTestGlobalVariable.cpp diff --git a/tests/src/kernel/hipTestGlobalVariable.cpp b/tests/src/kernel/hipTestGlobalVariable.cpp new file mode 100644 index 0000000000..2209e2c254 --- /dev/null +++ b/tests/src/kernel/hipTestGlobalVariable.cpp @@ -0,0 +1,97 @@ +/* +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 + * RUN: %t + * HIT_END + */ + +#include +#include +#include +#include "test_common.h" + +#define HIP_ASSERT(status) assert(status == hipSuccess) + +#define LEN 512 +#define SIZE 2048 + +struct TestConstantGlobalVar { + static __constant__ int ConstantGlobalVar; + + static __global__ void kernel(int* Ad) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + Ad[tid] = ConstantGlobalVar; + } + + void run() { + int *A, *Ad; + A = new int[LEN]; + for (unsigned i = 0; i < LEN; i++) { + A[i] = 0; + } + + HIP_ASSERT(hipMalloc((void**)&Ad, SIZE)); + hipLaunchKernelGGL(kernel, dim3(1, 1, 1), dim3(LEN, 1, 1), 0, 0, Ad); + HIP_ASSERT(hipMemcpy(A, Ad, SIZE, hipMemcpyDeviceToHost)); + + for (unsigned i = 0; i < LEN; i++) { + assert(123 == A[i]); + } + } +}; +__constant__ int TestConstantGlobalVar::ConstantGlobalVar = 123; + +struct TestGlobalArray { + static __device__ int GlobalArray[LEN]; + + static __global__ void kernelWrite() { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + GlobalArray[tid] = tid; + } + static __global__ void kernelRead(int* Ad) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + Ad[tid] = GlobalArray[tid]; + } + + void run() { + int *A, *Ad; + A = new int[LEN]; + for (unsigned i = 0; i < LEN; i++) { + A[i] = 0; + } + + HIP_ASSERT(hipMalloc((void**)&Ad, SIZE)); + hipLaunchKernelGGL(kernelWrite, dim3(1, 1, 1), dim3(LEN, 1, 1), 0, 0); + hipLaunchKernelGGL(kernelRead, dim3(1, 1, 1), dim3(LEN, 1, 1), 0, 0, Ad); + HIP_ASSERT(hipMemcpy(A, Ad, SIZE, hipMemcpyDeviceToHost)); + + for (unsigned i = 0; i < LEN; i++) { + assert(i == A[i]); + } + } +}; +__device__ int TestGlobalArray::GlobalArray[LEN]; + +int main() { + TestConstantGlobalVar().run(); + TestGlobalArray().run(); + passed(); +} From dbcede2ab85e0311677c20a464e7b85d18f4247b Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Wed, 15 Aug 2018 21:53:20 +0000 Subject: [PATCH 05/25] Remove few hcc specific cmake for hip-clang --- CMakeLists.txt | 6 ++++-- hip-config.cmake.in | 4 +++- packaging/hip-targets-release.cmake | 14 ++++++++++++++ packaging/hip-targets.cmake | 6 ++++++ 4 files changed, 27 insertions(+), 3 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 818c75783e..edb8c2d238 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -217,9 +217,11 @@ if(HIP_PLATFORM STREQUAL "hcc") set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -L/opt/rocm/profiler/CXLActivityLogger/bin/x86_64 -lCXLActivityLogger") endif() add_library(hip_hcc SHARED ${SOURCE_FILES_RUNTIME}) - target_link_libraries(hip_hcc PRIVATE hc_am) add_library(hip_hcc_static STATIC ${SOURCE_FILES_RUNTIME}) - target_link_libraries(hip_hcc_static PRIVATE hc_am) + if(HIP_COMPILER STREQUAL "hcc") + target_link_libraries(hip_hcc PRIVATE hc_am) + target_link_libraries(hip_hcc_static PRIVATE hc_am) + endif() add_library(hip_device STATIC ${SOURCE_FILES_DEVICE}) string(REPLACE " " ";" HCC_CXX_FLAGS_LIST ${HCC_CXX_FLAGS}) diff --git a/hip-config.cmake.in b/hip-config.cmake.in index efcdf708bb..d5dc6803fc 100644 --- a/hip-config.cmake.in +++ b/hip-config.cmake.in @@ -48,7 +48,9 @@ set_and_check( hip_BIN_INSTALL_DIR "@PACKAGE_BIN_INSTALL_DIR@" ) set_and_check(hip_HIPCC_EXECUTABLE "${hip_BIN_INSTALL_DIR}/hipcc") set_and_check(hip_HIPCONFIG_EXECUTABLE "${hip_BIN_INSTALL_DIR}/hipconfig") -find_dependency(hcc) +if(HIP_COMPILER STREQUAL "hcc") + find_dependency(hcc) +endif() include( "${CMAKE_CURRENT_LIST_DIR}/hip-targets.cmake" ) set( hip_LIBRARIES hip::host hip::device) diff --git a/packaging/hip-targets-release.cmake b/packaging/hip-targets-release.cmake index ba0a5005f5..0ae7405cde 100644 --- a/packaging/hip-targets-release.cmake +++ b/packaging/hip-targets-release.cmake @@ -7,22 +7,36 @@ set(CMAKE_IMPORT_FILE_VERSION 1) # Import target "hip::hip_hcc_static" for configuration "Release" set_property(TARGET hip::hip_hcc_static APPEND PROPERTY IMPORTED_CONFIGURATIONS RELEASE) +if(HIP_COMPILER STREQUAL "clang") +set_target_properties(hip::hip_hcc_static PROPERTIES + IMPORTED_LINK_INTERFACE_LANGUAGES_RELEASE "CXX" + IMPORTED_LOCATION_RELEASE "/opt/rocm/hip/lib/libhip_hcc_static.a" + ) +else() set_target_properties(hip::hip_hcc_static PROPERTIES IMPORTED_LINK_INTERFACE_LANGUAGES_RELEASE "CXX" IMPORTED_LINK_INTERFACE_LIBRARIES_RELEASE "hc_am" IMPORTED_LOCATION_RELEASE "/opt/rocm/hip/lib/libhip_hcc_static.a" ) +endif() list(APPEND _IMPORT_CHECK_TARGETS hip::hip_hcc_static ) list(APPEND _IMPORT_CHECK_FILES_FOR_hip::hip_hcc_static "/opt/rocm/hip/lib/libhip_hcc_static.a" ) # Import target "hip::hip_hcc" for configuration "Release" set_property(TARGET hip::hip_hcc APPEND PROPERTY IMPORTED_CONFIGURATIONS RELEASE) +if(HIP_COMPILER STREQUAL "clang") +set_target_properties(hip::hip_hcc PROPERTIES + IMPORTED_LOCATION_RELEASE "/opt/rocm/hip/lib/libhip_hcc.so" + IMPORTED_SONAME_RELEASE "libhip_hcc.so" + ) +else() set_target_properties(hip::hip_hcc PROPERTIES IMPORTED_LINK_INTERFACE_LIBRARIES_RELEASE "hcc::hccrt;hcc::hc_am" IMPORTED_LOCATION_RELEASE "/opt/rocm/hip/lib/libhip_hcc.so" IMPORTED_SONAME_RELEASE "libhip_hcc.so" ) +endif() list(APPEND _IMPORT_CHECK_TARGETS hip::hip_hcc ) list(APPEND _IMPORT_CHECK_FILES_FOR_hip::hip_hcc "/opt/rocm/hip/lib/libhip_hcc.so" ) diff --git a/packaging/hip-targets.cmake b/packaging/hip-targets.cmake index 5aff2ca1ed..ec2fa716a6 100644 --- a/packaging/hip-targets.cmake +++ b/packaging/hip-targets.cmake @@ -75,9 +75,15 @@ set_target_properties(hip::host PROPERTIES # Create imported target hip::device add_library(hip::device INTERFACE IMPORTED) +if(HIP_COMPILER STREQUAL "clang") +set_target_properties(hip::device PROPERTIES + INTERFACE_LINK_LIBRARIES "hip::host;hip::hip_device" +) +else() set_target_properties(hip::device PROPERTIES INTERFACE_LINK_LIBRARIES "hip::host;hip::hip_device;hcc::hccrt;hcc::hc_am" ) +endif() if(CMAKE_VERSION VERSION_LESS 3.0.0) message(FATAL_ERROR "This file relies on consumers using CMake 3.0.0 or greater.") From 1299b65e15935b48f3a46eb429ce2d902bf8c976 Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Fri, 17 Aug 2018 11:33:11 -0400 Subject: [PATCH 06/25] Add HIP_DB=fatbin for debugging fat binary issues --- src/hip_clang.cpp | 2 ++ src/hip_hcc.cpp | 2 +- src/hip_hcc_internal.h | 4 +++- 3 files changed, 6 insertions(+), 2 deletions(-) diff --git a/src/hip_clang.cpp b/src/hip_clang.cpp index f7573e2819..15a96d298a 100644 --- a/src/hip_clang.cpp +++ b/src/hip_clang.cpp @@ -58,6 +58,7 @@ __hipRegisterFatBinary(const void* data) { HIP_INIT(); + tprintf(DB_FB, "Enter __hipRegisterFatBinary(%p)\n", data); const __CudaFatBinaryWrapper* fbwrapper = reinterpret_cast(data); if (fbwrapper->magic != __hipFatMAGIC2 || fbwrapper->version != 1) { return nullptr; @@ -113,6 +114,7 @@ __hipRegisterFatBinary(const void* data) } } + tprintf(DB_FB, "__hipRegisterFatBinary succeeds and returns %p\n", modules); return modules; } diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index aecf75b717..eff93da847 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -1228,7 +1228,7 @@ void HipReadEnv() { READ_ENV_C(release, HIP_DB, 0, "Print debug info. Bitmask (HIP_DB=0xff) or flags separated by '+' " - "(HIP_DB=api+sync+mem+copy)", + "(HIP_DB=api+sync+mem+copy+fatbin)", HIP_DB_callback); if ((HIP_DB & (1 << DB_API)) && (HIP_TRACE_API == 0)) { // Set HIP_TRACE_API default before we read it, so it is printed correctly. diff --git a/src/hip_hcc_internal.h b/src/hip_hcc_internal.h index be257aff4f..d64a4a4cbe 100644 --- a/src/hip_hcc_internal.h +++ b/src/hip_hcc_internal.h @@ -223,7 +223,8 @@ extern const char* API_COLOR_END; #define DB_MEM 2 /* 0x04 - trace memory allocation / deallocation */ #define DB_COPY 3 /* 0x08 - trace memory copy and peer commands. . */ #define DB_WARN 4 /* 0x10 - warn about sub-optimal or shady behavior */ -#define DB_MAX_FLAG 5 +#define DB_FB 5 /* 0x20 - trace loading fat binary */ +#define DB_MAX_FLAG 6 // When adding a new debug flag, also add to the char name table below. // // @@ -237,6 +238,7 @@ struct DbName { static const DbName dbName[] = { {KGRN, "api"}, // not used, {KYEL, "sync"}, {KCYN, "mem"}, {KMAG, "copy"}, {KRED, "warn"}, + {KBLU, "fatbin"}, }; From 8e0e373f69b33dd1d9f4b0ef8e4bf974d3e08188 Mon Sep 17 00:00:00 2001 From: Saleel Kudchadker Date: Fri, 17 Aug 2018 12:14:42 -0700 Subject: [PATCH 07/25] Check for hipEnvVar at the same level if directed_tests location fails --- tests/src/hipEnvVarDriver.cpp | 11 +++++++++-- 1 file changed, 9 insertions(+), 2 deletions(-) diff --git a/tests/src/hipEnvVarDriver.cpp b/tests/src/hipEnvVarDriver.cpp index 599b138565..e52aa46063 100644 --- a/tests/src/hipEnvVarDriver.cpp +++ b/tests/src/hipEnvVarDriver.cpp @@ -38,7 +38,10 @@ int getDeviceNumber() { string str; std::this_thread::sleep_for(std::chrono::milliseconds(10)); if (!(in = popen("./directed_tests/hipEnvVar -c", "r"))) { - return 1; + // Check at same level + if (!(in = popen("./hipEnvVar -c", "r"))) { + return 1; + } } while (fgets(buff, 512, in) != NULL) { cout << buff; @@ -54,7 +57,11 @@ void getDevicePCIBusNumRemote(int deviceID, char* pciBusID) { str += std::to_string(deviceID); std::this_thread::sleep_for(std::chrono::milliseconds(10)); if (!(in = popen(str.c_str(), "r"))) { - exit(1); + // Check at same level + if (!(in = popen("./hipEnvVar -d ", "r"))) { + exit(1); + } + } while (fgets(pciBusID, 100, in) != NULL) { cout << pciBusID; From 1daee67eb682c6962013109f679a65d20a5615b3 Mon Sep 17 00:00:00 2001 From: Saleel Kudchadker Date: Fri, 31 Aug 2018 12:54:23 -0700 Subject: [PATCH 08/25] Fix record_event and hipStreamSync2 tests. The test should expect null stream to complete if synchrionize is called as per the spec --- tests/src/runtimeApi/event/record_event.cpp | 4 ++-- tests/src/runtimeApi/stream/hipStreamSync2.cpp | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/tests/src/runtimeApi/event/record_event.cpp b/tests/src/runtimeApi/event/record_event.cpp index afd6bef2ef..3eb54e1735 100644 --- a/tests/src/runtimeApi/event/record_event.cpp +++ b/tests/src/runtimeApi/event/record_event.cpp @@ -170,8 +170,8 @@ void runTests(int64_t numElements) { // for (int waitStart=0; waitStart<2; waitStart++) { for (int waitStart = 1; waitStart >= 0; waitStart--) { unsigned W = waitStart ? 0x1000 : 0; - test(W | 0x01, C_d, C_h, numElements, 0, waitStart, syncNone); - test(W | 0x02, C_d, C_h, numElements, stream, waitStart, syncNone); + test(W | 0x01, C_d, C_h, numElements, 0, 0, syncNone); + test(W | 0x02, C_d, C_h, numElements, stream, 0, syncNone); test(W | 0x04, C_d, C_h, numElements, 0, waitStart, syncStream); test(W | 0x08, C_d, C_h, numElements, stream, waitStart, syncStream); test(W | 0x10, C_d, C_h, numElements, 0, waitStart, syncStopEvent); diff --git a/tests/src/runtimeApi/stream/hipStreamSync2.cpp b/tests/src/runtimeApi/stream/hipStreamSync2.cpp index cf25d0bd2b..652c799792 100644 --- a/tests/src/runtimeApi/stream/hipStreamSync2.cpp +++ b/tests/src/runtimeApi/stream/hipStreamSync2.cpp @@ -178,7 +178,7 @@ void runTests(int64_t numElements) { { test(0x01, C_d, C_h, numElements, syncNone, true /*expectMismatch*/); test(0x02, C_d, C_h, numElements, syncNullStream, false /*expectMismatch*/); - test(0x04, C_d, C_h, numElements, syncOtherStream, true /*expectMismatch*/); + test(0x04, C_d, C_h, numElements, syncOtherStream, false /*expectMismatch*/); test(0x08, C_d, C_h, numElements, syncDevice, false /*expectMismatch*/); // Sending a marker to to null stream may synchronize the otherStream From 338eaefa8418c194ef99a041b3e2a11321dc7a4b Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Tue, 28 Aug 2018 19:13:27 +0000 Subject: [PATCH 09/25] Support placement new in hip-clang --- include/hip/hcc_detail/hip_runtime.h | 1 + tests/src/deviceLib/hipTestNew.cpp | 71 ++++++++++++++++++++++++++++ 2 files changed, 72 insertions(+) create mode 100644 tests/src/deviceLib/hipTestNew.cpp diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 41710e4165..7bd89e2815 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -471,6 +471,7 @@ hc_get_workitem_absolute_id(int dim) #include <__clang_cuda_complex_builtins.h> #include #include +#include #undef __CUDA__ #pragma pop_macro("__CUDA__") diff --git a/tests/src/deviceLib/hipTestNew.cpp b/tests/src/deviceLib/hipTestNew.cpp new file mode 100644 index 0000000000..60774ff21d --- /dev/null +++ b/tests/src/deviceLib/hipTestNew.cpp @@ -0,0 +1,71 @@ +/* +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 + * RUN: %t + * HIT_END + */ + +#include +#include +#include +#include "test_common.h" + +#define HIP_ASSERT(status) assert(status == hipSuccess) + +#define LEN 512 +#define SIZE 2048 + +struct TestPlacementNew { + class A { + public: + __device__ A() { + a = threadIdx.x + blockIdx.x * blockDim.x; + } + private: + int a; + }; + + static __global__ void kernel(int* Ad) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + new(Ad+tid) A(); + } + + void run() { + int *A, *Ad; + A = new int[LEN]; + for (unsigned i = 0; i < LEN; i++) { + A[i] = 0; + } + + HIP_ASSERT(hipMalloc((void**)&Ad, SIZE)); + hipLaunchKernelGGL(kernel, dim3(1, 1, 1), dim3(LEN, 1, 1), 0, 0, Ad); + HIP_ASSERT(hipMemcpy(A, Ad, SIZE, hipMemcpyDeviceToHost)); + + for (unsigned i = 0; i < LEN; i++) { + assert(i == A[i]); + } + } +}; + +int main() { + TestPlacementNew().run(); + passed(); +} From 1bb28375bcf3313b403f4c42c4f88a89f6387255 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Mon, 27 Aug 2018 14:51:25 -0400 Subject: [PATCH 10/25] Document kernel launching for hip-clang --- docs/markdown/hip_porting_driver_api.md | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/docs/markdown/hip_porting_driver_api.md b/docs/markdown/hip_porting_driver_api.md index 2e61f0eb32..7d02666af3 100644 --- a/docs/markdown/hip_porting_driver_api.md +++ b/docs/markdown/hip_porting_driver_api.md @@ -102,6 +102,15 @@ hip-clang links device code from different translation units together. For each hip-clang generates initializatiion and termination functions for each translation unit for host code compilation. The initialization functions call `__hipRegisterFatBinary` to register the fatbinary embeded in the ELF file. They also call `__hipRegisterFunction` and `__hipRegisterVar` to register kernel functions and device side global variables. The termination functions call `__hipUnregisterFatBinary`. hip-clang emits a global variable `__hip_gpubin_handle` of void** type with linkonce linkage and inital value 0 for each host translation unit. Each initialization function checks `__hip_gpubin_handle` and register the fatbinary only if `__hip_gpubin_handle` is 0 and saves the return value of `__hip_gpubin_handle` to `__hip_gpubin_handle`. This is to guarantee that the fatbinary is only registered once. Similar check is done in the termination functions. +#### Kernel Launching +hip-clang supports kernel launching by CUDA `<<<>>>` syntax, hipLaunchKernel, and hipLaunchKernelGGL. The latter two are macros which expand to CUDA `<<<>>>` syntax. + +In host code, hip-clang emits a stub function with the same name and arguments as the kernel. In the body of this function, hipSetupArgument is called for each kernel argument, then hipLaunchByPtr is called with a function pointer to the stub function. + +When the executable or shared library is loaded by the dynamic linker, the initilization functions are called. In the initialization functions, when `__hipRegisterFatBinary` is called, the code objects containing all kernels are loaded; when `__hipRegisterFunction` is called, the stub functions are associated with the corresponding kernels in code objects. + +In the host code, for the `<<<>>>` statement, hip-clang first emits call of hipConfigureCall to set up the threads and grids, then emits call of the stub function with the given arguments. In the stub function, when the runtime host API function hipLaunchByPtr is called, the real kernel associated with the stub function is launched. + ### NVCC Implementation Notes #### Interoperation between HIP and CUDA Driver From 1587b18a918910d516bee2dac68b625aad3059cb Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Wed, 22 Aug 2018 15:42:54 -0400 Subject: [PATCH 11/25] Let hipcc handle obj files in linker response file for hip-clang If obj files in linker response file contains device code, pass them to hip-clang, otherwise keep them in the linker response file. --- bin/hipcc | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/bin/hipcc b/bin/hipcc index 92ec2578cb..4610a6be89 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -453,6 +453,15 @@ foreach $arg (@ARGV) system("cd $tmpdir; ar c $libBaseName $realObjs"); print $out "$tmpdir/$libBaseName\n"; } + } elsif ($line =~ m/\.o$/) { + my $fileType = `file $line`; + my $isObj = ($fileType =~ m/ELF/ or $fileType =~ m/COFF/); + if ($isObj) { + print $out "$line\n"; + } else { + push (@inputs, $line); + $new_arg = "$new_arg $line"; + } } else { print $out "$line\n"; } From 4f17b4877e7b8c1a5d2978ef6fdb175882ef22ba Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Fri, 17 Aug 2018 11:34:07 -0400 Subject: [PATCH 12/25] Let hipcc link -lpthread -lm by default --- bin/hipcc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/bin/hipcc b/bin/hipcc index 92ec2578cb..3ff7cd6ead 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -725,7 +725,7 @@ if ($needLDFLAGS and not $compileOnly) { } $CMD .= " $toolArgs"; if ($needLDFLAGS and not $compileOnly and $HIP_PLATFORM eq "clang") { - $CMD .= " -lgcc_s -lgcc"; + $CMD .= " -lgcc_s -lgcc -lpthread -lm"; } if ($verbose & 0x1) { From a8bc26344f778e7149f641c2ec104c446962efab Mon Sep 17 00:00:00 2001 From: fpadmin Date: Tue, 4 Sep 2018 11:58:19 +0530 Subject: [PATCH 13/25] [HIP-Clang]Add friend class/function test --- tests/src/compiler/hipClassKernel.cpp | 89 ++++++++++++++++++++++++++- tests/src/compiler/hipClassKernel.h | 56 ++++++++++++++++- 2 files changed, 142 insertions(+), 3 deletions(-) diff --git a/tests/src/compiler/hipClassKernel.cpp b/tests/src/compiler/hipClassKernel.cpp index b3502ca9ae..edc3808d83 100644 --- a/tests/src/compiler/hipClassKernel.cpp +++ b/tests/src/compiler/hipClassKernel.cpp @@ -26,6 +26,66 @@ THE SOFTWARE. */ #include "hipClassKernel.h" +#ifdef ENABLE_OVERLOAD_OVERRIDE_TESTS +__global__ void +ovrdClassKernel(bool* result_ecd){ + int tid = threadIdx.x + blockIdx.x * blockDim.x; + testOvrD tobj1; + result_ecd[tid] = (tobj1.ovrdFunc1() == 30); +} + +void HipClassTests::TestForOverride(void){ + bool *result_ecd, *result_ech; + result_ech = HipClassTests::AllocateHostMemory(); + result_ecd = HipClassTests::AllocateDeviceMemory(); + + hipLaunchKernelGGL(ovrdClassKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0, + result_ecd); + + HipClassTests::VerifyResult(result_ech,result_ecd); + HipClassTests::FreeMem(result_ech,result_ecd); +} + + +__global__ void +ovldClassKernel(bool* result_ecd){ + int tid = threadIdx.x + blockIdx.x * blockDim.x; + testFuncOvld tfo1; + result_ecd[tid] = (tfo1.func1(10) == 20) + && (tfo1.func1(10,10) == 30); +} + +void HipClassTests::TestForOverload(void){ + bool *result_ecd, *result_ech; + result_ech = HipClassTests::AllocateHostMemory(); + result_ecd = HipClassTests::AllocateDeviceMemory(); + + hipLaunchKernelGGL(ovldClassKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0, + result_ecd); + + HipClassTests::VerifyResult(result_ech,result_ecd); + HipClassTests::FreeMem(result_ech,result_ecd); +} +#endif + +#ifdef ENABLE_FRIEND_TEST +// check for friend +__global__ void +friendClassKernel(bool* result_ecd){ + int tid = threadIdx.x + blockIdx.x * blockDim.x; + testFrndB tfb1; + result_ecd[tid] = (tfb1.showA() == 10); +} +#endif + // check sizeof empty class is 1 __global__ void emptyClassKernel(bool* result_ecd) { @@ -209,6 +269,20 @@ void HipClassTests::TestForConsrtDesrt(){ } #endif +#ifdef ENABLE_FRIEND_TEST +void HipClassTests::TestForFriend(void){ + bool *result_ecd, *result_ech; + result_ech = HipClassTests::AllocateHostMemory(); + result_ecd = HipClassTests::AllocateDeviceMemory(); + hipLaunchKernelGGL(friendClassKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0, + result_ecd); +} +#endif + bool* HipClassTests::AllocateHostMemory(void){ bool *result_ech; HIPCHECK(hipHostMalloc(&result_ech, @@ -253,6 +327,19 @@ int main(){ test_passed(TestForClassSize); classTests.TestForPassByValue(); test_passed(TestForPassByValue); + +#ifdef ENABLE_OVERLOAD_OVERRIDE_TESTS + classTests.TestForOverload(); + test_passed(TestForOverload); + classTests.TestForOverride(); + test_passed(TestForOverride); +#endif + +#ifdef ENABLE_FRIEND_TEST + classTests.TestForFriend(); + test_passed(TestForFriend); +#endif + // classTests.TestForMallocPassByValue(); // test_passed(TestForMallocPassByValue); #this test is crashing @@ -261,8 +348,6 @@ int main(){ test_passed(TestForVirtualClassSize); #endif - - #ifdef ENABLE_DESTRUCTOR_TEST classTests.TestForConsrtDesrt(); test_passed(TestForConsrtDesrt); diff --git a/tests/src/compiler/hipClassKernel.h b/tests/src/compiler/hipClassKernel.h index 8538276662..09d5f22bea 100644 --- a/tests/src/compiler/hipClassKernel.h +++ b/tests/src/compiler/hipClassKernel.h @@ -31,10 +31,62 @@ static const int BLOCKS = 512; static const int THREADS_PER_BLOCK = 1; static const int ENABLE_DESTRUCTOR_TEST = 0; static const int ENABLE_VIRTUAL_TESTS = 0; +static const int ENABLE_FRIEND_TEST = 0; +static const int ENABLE_OVERLAD_OVERRIDE_TESTS = 0; size_t NBOOL = BLOCKS * sizeof(bool); #define test_passed(test_name) printf("%s %s PASSED!%s\n", KGRN, #test_name, KNRM); +#ifdef ENABLE_OVERLOAD_OVERRIDE_TESTS +class testFuncOvld{ + public: + int __host__ __device__ func1(int a){ + return a + 10; + } + + int __host__ __device__ func1(int a , int b){ + return a + b + 10; + } + +}; + + +class testOvrB{ + public: + int __host__ __device__ ovrdFunc1(){ + return 10; + } + + +}; + + +class testOvrD: public testOvrB{ + public: + int __host__ __device__ ovrdFunc1(){ + int x = testOvrB::ovrdFunc1(); + return x + 20; + } + +}; +#endif + +#ifdef ENABLE_FRIEND_TEST +class testFrndA{ + private: + int fa = 10; + public: + friend class testFrndB; +}; + +class testFrndB{ + public: + __host__ __device__ int showA(){ + testFrndA x; + return x.fa; + } +}; +#endif class testClassEmpty {}; @@ -177,7 +229,9 @@ class HipClassTests{ void TestForPassByValue(void); void TestForMallocPassByValue(void); void TestForConsrtDesrt(void); - + void TestForOverload(void); + void TestForOverride(void); + bool* AllocateHostMemory(void); bool* AllocateDeviceMemory(void); void VerifyResult(bool* result_ech, bool* result_ecd); From 2984c020b8b9020a61326580a452c45cd6545175 Mon Sep 17 00:00:00 2001 From: lcskrishna Date: Thu, 6 Sep 2018 15:02:52 -0700 Subject: [PATCH 14/25] added host for float2half and hlaf2float --- include/hip/hcc_detail/hip_fp16.h | 26 +++++++++++++------------- 1 file changed, 13 insertions(+), 13 deletions(-) diff --git a/include/hip/hcc_detail/hip_fp16.h b/include/hip/hcc_detail/hip_fp16.h index 777113a9de..68f0e35f5f 100644 --- a/include/hip/hcc_detail/hip_fp16.h +++ b/include/hip/hcc_detail/hip_fp16.h @@ -635,37 +635,37 @@ THE SOFTWARE. // TODO: rounding behaviour is not correct. // float -> half | half2 inline - __device__ + __device__ __host__ __half __float2half(float x) { return __half_raw{static_cast<_Float16>(x)}; } inline - __device__ + __device__ __host__ __half __float2half_rn(float x) { return __half_raw{static_cast<_Float16>(x)}; } inline - __device__ + __device__ __host__ __half __float2half_rz(float x) { return __half_raw{static_cast<_Float16>(x)}; } inline - __device__ + __device__ __host__ __half __float2half_rd(float x) { return __half_raw{static_cast<_Float16>(x)}; } inline - __device__ + __device__ __host__ __half __float2half_ru(float x) { return __half_raw{static_cast<_Float16>(x)}; } inline - __device__ + __device__ __host__ __half2 __float2half2_rn(float x) { return __half2_raw{ @@ -673,14 +673,14 @@ THE SOFTWARE. static_cast<_Float16>(x), static_cast<_Float16>(x)}}; } inline - __device__ + __device__ __host__ __half2 __floats2half2_rn(float x, float y) { return __half2_raw{_Float16_2{ static_cast<_Float16>(x), static_cast<_Float16>(y)}}; } inline - __device__ + __device__ __host__ __half2 __float22half2_rn(float2 x) { return __floats2half2_rn(x.x, x.y); @@ -688,25 +688,25 @@ THE SOFTWARE. // half | half2 -> float inline - __device__ + __device__ __host__ float __half2float(__half x) { return static_cast<__half_raw>(x).data; } inline - __device__ + __device__ __host__ float __low2float(__half2 x) { return static_cast<__half2_raw>(x).data.x; } inline - __device__ + __device__ __host__ float __high2float(__half2 x) { return static_cast<__half2_raw>(x).data.y; } inline - __device__ + __device__ __host__ float2 __half22float2(__half2 x) { return make_float2( @@ -1633,4 +1633,4 @@ THE SOFTWARE. #endif // defined(__cplusplus) #elif defined(__GNUC__) #include "hip_fp16_gcc.h" -#endif // !defined(__clang__) && defined(__GNUC__) \ No newline at end of file +#endif // !defined(__clang__) && defined(__GNUC__) From 3e1833ca52b72dd9b0abef92b28e2b18391be8ae Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Fri, 7 Sep 2018 16:15:08 -0400 Subject: [PATCH 15/25] Add empty printf for hip-clang --- include/hip/hcc_detail/hip_runtime.h | 7 ++----- 1 file changed, 2 insertions(+), 5 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 41710e4165..aecceea163 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -264,19 +264,16 @@ extern "C" __device__ void* __hip_hc_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); } -#ifdef __HCC_ACCELERATOR__ - -#ifdef HC_FEATURE_PRINTF +#if defined(__HCC_ACCELERATOR__) && defined(HC_FEATURE_PRINTF) template static inline __device__ void printf(const char* format, All... all) { hc::printf(format, all...); } -#else +#elif defined(__HCC_ACCELERATOR__) || __HIP__ template static inline __device__ void printf(const char* format, All... all) {} #endif -#endif #endif //__HCC_OR_HIP_CLANG__ #ifdef __HCC__ From 9e9a93e10a5ffde6a027ef5ee9ba79cdc34e1a9b Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Fri, 7 Sep 2018 16:20:00 -0400 Subject: [PATCH 16/25] Use template for hipLaunchKernelGGL for hip-clang --- include/hip/hcc_detail/hip_runtime.h | 19 +++++++++++-------- 1 file changed, 11 insertions(+), 8 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 41710e4165..e0ad6befdd 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -346,15 +346,18 @@ extern void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, gri typedef int hipLaunchParm; -#define hipLaunchKernel(kernelName, numblocks, numthreads, memperblock, streamId, ...) \ - do { \ - kernelName<<>>(0, ##__VA_ARGS__); \ - } while (0) +template +inline void hipLaunchKernelGGL(F kernelName, const dim3& numblocks, const dim3& numthreads, + unsigned memperblock, hipStream_t streamId, Args... args) { + kernelName<<>>(args...); +} -#define hipLaunchKernelGGL(kernelName, numblocks, numthreads, memperblock, streamId, ...) \ - do { \ - kernelName<<>>(__VA_ARGS__); \ - } while (0) +template +inline void hipLaunchKernel(F kernel, const dim3& numBlocks, const dim3& dimBlocks, + std::uint32_t groupMemBytes, hipStream_t stream, Args... args) { + hipLaunchKernelGGL(kernel, numBlocks, dimBlocks, groupMemBytes, stream, hipLaunchParm{}, + std::move(args)...); +} #include From 0121ec13aa4fe825d77c648e814996f0b2be3605 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Tue, 11 Sep 2018 17:55:28 +0000 Subject: [PATCH 17/25] Use templates for min to prevent ambiguity --- include/hip/hcc_detail/hip_runtime.h | 10 --------- include/hip/hcc_detail/math_functions.h | 28 +++++++++++++++++++++++++ 2 files changed, 28 insertions(+), 10 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 41710e4165..f5294dfc56 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -199,16 +199,6 @@ __device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask #endif //__HIP_ARCH_GFX803__ == 1 -__device__ inline static int min(int arg1, int arg2) { - return (arg1 < arg2) ? arg1 : arg2; -} -__device__ inline static int max(int arg1, int arg2) { - return (arg1 > arg2) ? arg1 : arg2; -} - -__host__ inline static int min(int arg1, int arg2) { return std::min(arg1, arg2); } -__host__ inline static int max(int arg1, int arg2) { return std::max(arg1, arg2); } - #endif // __HCC_OR_HIP_CLANG__ #if defined __HCC__ diff --git a/include/hip/hcc_detail/math_functions.h b/include/hip/hcc_detail/math_functions.h index 10967605d4..5457a99cfd 100644 --- a/include/hip/hcc_detail/math_functions.h +++ b/include/hip/hcc_detail/math_functions.h @@ -1304,6 +1304,32 @@ float func(float x, int y) \ } __DEF_FLOAT_FUN2I(scalbn) +#if __HCC__ +template +__DEVICE__ __host__ inline static const T min(T arg1, T arg2) { + return std::min(arg1, arg2); +} + +template +__DEVICE__ __host__ inline static const T max(T arg1, T arg2) { + return std::max(arg1, arg2); +} +#else +__device__ inline static int min(int arg1, int arg2) { + return (arg1 < arg2) ? arg1 : arg2; +} +__device__ inline static int max(int arg1, int arg2) { + return (arg1 > arg2) ? arg1 : arg2; +} + +__host__ inline static int min(int arg1, int arg2) { + return std::min(arg1, arg2); +} + +__host__ inline static int max(int arg1, int arg2) { + return std::max(arg1, arg2); +} + __DEVICE__ inline float max(float x, float y) { @@ -1331,6 +1357,8 @@ double min(double x, double y) { __HIP_OVERLOAD2(double, max) __HIP_OVERLOAD2(double, min) +#endif + #pragma pop_macro("__DEF_FLOAT_FUN") #pragma pop_macro("__DEF_FLOAT_FUN2") #pragma pop_macro("__DEF_FLOAT_FUN2I") From ffd89dde9cac20a83bc82c9a07f698f57dcbfcba Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Tue, 11 Sep 2018 18:48:31 +0000 Subject: [PATCH 18/25] Avoid host min func conflict with gcc min --- include/hip/hcc_detail/math_functions.h | 25 +++++++++++++------------ 1 file changed, 13 insertions(+), 12 deletions(-) diff --git a/include/hip/hcc_detail/math_functions.h b/include/hip/hcc_detail/math_functions.h index 5457a99cfd..b44d3f00bd 100644 --- a/include/hip/hcc_detail/math_functions.h +++ b/include/hip/hcc_detail/math_functions.h @@ -1306,30 +1306,22 @@ __DEF_FLOAT_FUN2I(scalbn) #if __HCC__ template -__DEVICE__ __host__ inline static const T min(T arg1, T arg2) { +__DEVICE__ inline static T min(T arg1, T arg2) { return std::min(arg1, arg2); } template -__DEVICE__ __host__ inline static const T max(T arg1, T arg2) { +__DEVICE__ inline static T max(T arg1, T arg2) { return std::max(arg1, arg2); } #else -__device__ inline static int min(int arg1, int arg2) { +__DEVICE__ inline static int min(int arg1, int arg2) { return (arg1 < arg2) ? arg1 : arg2; } -__device__ inline static int max(int arg1, int arg2) { +__DEVICE__ inline static int max(int arg1, int arg2) { return (arg1 > arg2) ? arg1 : arg2; } -__host__ inline static int min(int arg1, int arg2) { - return std::min(arg1, arg2); -} - -__host__ inline static int max(int arg1, int arg2) { - return std::max(arg1, arg2); -} - __DEVICE__ inline float max(float x, float y) { @@ -1359,6 +1351,15 @@ __HIP_OVERLOAD2(double, min) #endif +__host__ inline static int min(int arg1, int arg2) { + return std::min(arg1, arg2); +} + +__host__ inline static int max(int arg1, int arg2) { + return std::max(arg1, arg2); +} + + #pragma pop_macro("__DEF_FLOAT_FUN") #pragma pop_macro("__DEF_FLOAT_FUN2") #pragma pop_macro("__DEF_FLOAT_FUN2I") From cb6cf6584e9ae3c79886b4e8adbe1ab89d627752 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 12 Sep 2018 10:32:15 +0530 Subject: [PATCH 19/25] [ci] Disable hipVectorTypes* tests in CI Disable directed_tests/deviceLib/hipVectorTypes.tst & directed_tests/deviceLib/hipVectorTypesDevice.tst in CI due to HCC regressions. Once HCC fixes are in, the tests can be re-enabled in CI. --- Jenkinsfile | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/Jenkinsfile b/Jenkinsfile index 15ca1c9935..2432cea38e 100644 --- a/Jenkinsfile +++ b/Jenkinsfile @@ -177,7 +177,7 @@ def docker_build_inside_image( def build_image, String inside_args, String platf cd ${build_dir_rel} make install -j\$(nproc) make build_tests -i -j\$(nproc) - ctest + ctest -E hipVectorTypes """ // If unit tests output a junit or xunit file in the future, jenkins can parse that file // to display test results on the dashboard From d577f27d1a311d57de4326723acfdcf57e94f75c Mon Sep 17 00:00:00 2001 From: carlushuang Date: Wed, 12 Sep 2018 13:24:56 +0800 Subject: [PATCH 20/25] fix __longlong_as_double() problem, return the double value previous version return a long long valus *as* double, hence we may get the wrong result. this also affect atomicAdd(double * ...), which use long long pointer to mimic double pointer. Signed-off-by: carlushuang --- include/hip/hcc_detail/device_functions.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index adf631d782..508a301344 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/include/hip/hcc_detail/device_functions.h @@ -593,7 +593,7 @@ __device__ static inline double __longlong_as_double(long long int x) { double tmp; __builtin_memcpy(&tmp, &x, sizeof(tmp)); - return x; + return tmp; } __device__ static inline double __uint2double_rn(int x) { return (double)x; } From cdfea3ef7bb840218d6134b011d78e51069c833f Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Wed, 12 Sep 2018 10:30:48 +0100 Subject: [PATCH 21/25] Remove potential for mismatch between runtime passed actuals and defined formals. --- .../hip/hcc_detail/functional_grid_launch.hpp | 64 +++++++++++++------ 1 file changed, 44 insertions(+), 20 deletions(-) diff --git a/include/hip/hcc_detail/functional_grid_launch.hpp b/include/hip/hcc_detail/functional_grid_launch.hpp index de943f310d..66e5873f3a 100644 --- a/include/hip/hcc_detail/functional_grid_launch.hpp +++ b/include/hip/hcc_detail/functional_grid_launch.hpp @@ -51,30 +51,53 @@ inline T round_up_to_next_multiple_nonnegative(T x, T y) { return tmp - tmp % y; } -inline std::vector make_kernarg() { return {}; } - -inline std::vector make_kernarg(std::vector kernarg) { return kernarg; } - -template -inline std::vector make_kernarg(std::vector kernarg, T x) { - kernarg.resize(round_up_to_next_multiple_nonnegative(kernarg.size(), alignof(T)) + sizeof(T)); - - new (kernarg.data() + kernarg.size() - sizeof(T)) T{std::move(x)}; - +template < + std::size_t n, + typename... Ts, + typename std::enable_if::type* = nullptr> +inline std::vector make_kernarg( + std::vector kernarg, const std::tuple&) { return kernarg; } -template -inline std::vector make_kernarg(std::vector kernarg, T x, Ts... xs) { - return make_kernarg(make_kernarg(std::move(kernarg), std::move(x)), std::move(xs)...); +template < + std::size_t n, + typename... Ts, + typename std::enable_if::type* = nullptr> +inline std::vector make_kernarg( + std::vector kernarg, const std::tuple& formals) { + using T = typename std::tuple_element>::type; + + static_assert( + !std::is_reference{}, + "A __global__ function cannot have a reference as one of its " + "arguments."); + #if defined(HIP_STRICT) + static_assert( + std::is_trivially_copyable{}, + "Only TriviallyCopyable types can be arguments to a __global__ " + "function"); + #endif + + kernarg.resize(round_up_to_next_multiple_nonnegative( + kernarg.size(), alignof(T)) + sizeof(T)); + + new (kernarg.data() + kernarg.size() - sizeof(T)) T{std::get(formals)}; + + return make_kernarg(std::move(kernarg), formals); } -template -inline std::vector make_kernarg(Ts... xs) { - std::vector kernarg; - kernarg.reserve(sizeof(std::tuple)); +template +inline std::vector make_kernarg( + void (*)(Formals...), std::tuple actuals) { + static_assert(sizeof...(Formals) == sizeof...(Actuals), + "The count of formal arguments must match the count of actuals."); - return make_kernarg(std::move(kernarg), std::move(xs)...); + std::tuple to_formals{std::move(actuals)}; + std::vector kernarg; + kernarg.reserve(sizeof(to_formals)); + + return make_kernarg<0>(std::move(kernarg), to_formals); } void hipLaunchKernelGGLImpl(std::uintptr_t function_address, const dim3& numBlocks, @@ -85,7 +108,8 @@ void hipLaunchKernelGGLImpl(std::uintptr_t function_address, const dim3& numBloc template inline void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, std::uint32_t sharedMemBytes, hipStream_t stream, Args... args) { - auto kernarg = hip_impl::make_kernarg(std::move(args)...); + auto kernarg = hip_impl::make_kernarg( + kernel, std::tuple{std::move(args)...}); std::size_t kernarg_size = kernarg.size(); void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, kernarg.data(), HIP_LAUNCH_PARAM_BUFFER_SIZE, @@ -100,4 +124,4 @@ inline void hipLaunchKernel(F kernel, const dim3& numBlocks, const dim3& dimBloc std::uint32_t groupMemBytes, hipStream_t stream, Args... args) { hipLaunchKernelGGL(kernel, numBlocks, dimBlocks, groupMemBytes, stream, hipLaunchParm{}, std::move(args)...); -} +} \ No newline at end of file From fb706902cc688e4777df2d5b4f6e1a4c390ba482 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 12 Sep 2018 16:41:24 +0530 Subject: [PATCH 22/25] Remove adipose extn from launchKernelHcc sample --- samples/0_Intro/module_api/launchKernelHcc.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/samples/0_Intro/module_api/launchKernelHcc.cpp b/samples/0_Intro/module_api/launchKernelHcc.cpp index b5f40d68d9..f6bb9e5361 100644 --- a/samples/0_Intro/module_api/launchKernelHcc.cpp +++ b/samples/0_Intro/module_api/launchKernelHcc.cpp @@ -34,7 +34,7 @@ THE SOFTWARE. #define LEN 64 #define SIZE LEN << 2 -#define fileName "vcpy_kernel.code.adipose" +#define fileName "vcpy_kernel.code" #define kernel_name "hello_world" #define HIP_CHECK(status) \ From 894cbdd749786f01b950e422c93502310105693b Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Wed, 12 Sep 2018 14:54:31 +0000 Subject: [PATCH 23/25] Avoid AMP-retrict call to CPU-restrict --- include/hip/hcc_detail/math_functions.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/include/hip/hcc_detail/math_functions.h b/include/hip/hcc_detail/math_functions.h index b44d3f00bd..059aabcc53 100644 --- a/include/hip/hcc_detail/math_functions.h +++ b/include/hip/hcc_detail/math_functions.h @@ -1307,12 +1307,12 @@ __DEF_FLOAT_FUN2I(scalbn) #if __HCC__ template __DEVICE__ inline static T min(T arg1, T arg2) { - return std::min(arg1, arg2); + return (arg1 < arg2) ? arg1 : arg2; } template __DEVICE__ inline static T max(T arg1, T arg2) { - return std::max(arg1, arg2); + return (arg1 > arg2) ? arg1 : arg2; } #else __DEVICE__ inline static int min(int arg1, int arg2) { From 6b811ca6d158afacd94877aa54898aa35a982c21 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Thu, 13 Sep 2018 23:16:20 +0000 Subject: [PATCH 24/25] Fix Tensorflow ambiguous min issue --- include/hip/hcc_detail/math_functions.h | 42 +++++++++++++++++++++++++ 1 file changed, 42 insertions(+) diff --git a/include/hip/hcc_detail/math_functions.h b/include/hip/hcc_detail/math_functions.h index 059aabcc53..6c8510fcbb 100644 --- a/include/hip/hcc_detail/math_functions.h +++ b/include/hip/hcc_detail/math_functions.h @@ -1310,10 +1310,52 @@ __DEVICE__ inline static T min(T arg1, T arg2) { return (arg1 < arg2) ? arg1 : arg2; } +__DEVICE__ inline static uint32_t min(uint32_t arg1, int32_t arg2) { + return min(arg1, (uint32_t) arg2); +} +/*__DEVICE__ inline static uint32_t min(int32_t arg1, uint32_t arg2) { + return min((uint32_t) arg1, arg2); +} + +__DEVICE__ inline static uint64_t min(uint64_t arg1, int64_t arg2) { + return min(arg1, (uint64_t) arg2); +} +__DEVICE__ inline static uint64_t min(int64_t arg1, uint64_t arg2) { + return min((uint64_t) arg1, arg2); +} + +__DEVICE__ inline static unsigned long long min(unsigned long long arg1, long long arg2) { + return min(arg1, (unsigned long long) arg2); +} +__DEVICE__ inline static unsigned long long min(long long arg1, unsigned long long arg2) { + return min((unsigned long long) arg1, arg2); +}*/ + template __DEVICE__ inline static T max(T arg1, T arg2) { return (arg1 > arg2) ? arg1 : arg2; } + +__DEVICE__ inline static uint32_t max(uint32_t arg1, int32_t arg2) { + return max(arg1, (uint32_t) arg2); +} +__DEVICE__ inline static uint32_t max(int32_t arg1, uint32_t arg2) { + return max((uint32_t) arg1, arg2); +} + +/*__DEVICE__ inline static uint64_t max(uint64_t arg1, int64_t arg2) { + return max(arg1, (uint64_t) arg2); +} +__DEVICE__ inline static uint64_t max(int64_t arg1, uint64_t arg2) { + return max((uint64_t) arg1, arg2); +} + +__DEVICE__ inline static unsigned long long max(unsigned long long arg1, long long arg2) { + return max(arg1, (unsigned long long) arg2); +} +__DEVICE__ inline static unsigned long long max(long long arg1, unsigned long long arg2) { + return max((unsigned long long) arg1, arg2); +}*/ #else __DEVICE__ inline static int min(int arg1, int arg2) { return (arg1 < arg2) ? arg1 : arg2; From 9b2107749cd31c34434be03388a136fd17aee69e Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Sat, 15 Sep 2018 13:23:38 +0530 Subject: [PATCH 25/25] Revert changes to runtime/stream/hipStreamSync2 --- tests/src/runtimeApi/stream/hipStreamSync2.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/src/runtimeApi/stream/hipStreamSync2.cpp b/tests/src/runtimeApi/stream/hipStreamSync2.cpp index 652c799792..cf25d0bd2b 100644 --- a/tests/src/runtimeApi/stream/hipStreamSync2.cpp +++ b/tests/src/runtimeApi/stream/hipStreamSync2.cpp @@ -178,7 +178,7 @@ void runTests(int64_t numElements) { { test(0x01, C_d, C_h, numElements, syncNone, true /*expectMismatch*/); test(0x02, C_d, C_h, numElements, syncNullStream, false /*expectMismatch*/); - test(0x04, C_d, C_h, numElements, syncOtherStream, false /*expectMismatch*/); + test(0x04, C_d, C_h, numElements, syncOtherStream, true /*expectMismatch*/); test(0x08, C_d, C_h, numElements, syncDevice, false /*expectMismatch*/); // Sending a marker to to null stream may synchronize the otherStream