From 089ab3b9474a7643f4bd85493335def68f8db4d8 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Sat, 2 Jun 2018 12:27:17 +0100 Subject: [PATCH 01/18] Switch the atomic implementation to use Clang builtins. --- include/hip/hcc_detail/device_functions.h | 3 +- include/hip/hcc_detail/hip_atomic.h | 265 ++++++++++++++++ include/hip/hcc_detail/hip_runtime.h | 77 +---- include/hip/hip_runtime.h | 2 +- src/device_util.cpp | 123 -------- tests/src/deviceLib/hipSimpleAtomicsTest.cpp | 316 +++++++++++-------- 6 files changed, 447 insertions(+), 339 deletions(-) create mode 100644 include/hip/hcc_detail/hip_atomic.h diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index 28d874b27a..20a365ebbe 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/include/hip/hcc_detail/device_functions.h @@ -23,7 +23,8 @@ THE SOFTWARE. #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_DEVICE_FUNCTIONS_H #define HIP_INCLUDE_HIP_HCC_DETAIL_DEVICE_FUNCTIONS_H -#include +#include "host_defines.h" + #include diff --git a/include/hip/hcc_detail/hip_atomic.h b/include/hip/hcc_detail/hip_atomic.h new file mode 100644 index 0000000000..4af1794ba0 --- /dev/null +++ b/include/hip/hcc_detail/hip_atomic.h @@ -0,0 +1,265 @@ +#pragma once + +#include "device_functions.h" + +__device__ +inline +int atomicCAS(int* address, int compare, int val) +{ + __atomic_compare_exchange_n( + address, &compare, val, false, __ATOMIC_RELAXED, __ATOMIC_RELAXED); + + return compare; +} +__device__ +inline +unsigned int atomicCAS( + unsigned int* address, unsigned int compare, unsigned int val) +{ + __atomic_compare_exchange_n( + address, &compare, val, false, __ATOMIC_RELAXED, __ATOMIC_RELAXED); + + return compare; +} +__device__ +inline +unsigned long long atomicCAS( + unsigned long long* address, + unsigned long long compare, + unsigned long long val) +{ + __atomic_compare_exchange_n( + address, &compare, val, false, __ATOMIC_RELAXED, __ATOMIC_RELAXED); + + return compare; +} + +__device__ +inline +int atomicAdd(int* address, int val) +{ + return __atomic_fetch_add(address, val, __ATOMIC_RELAXED); +} +__device__ +inline +unsigned int atomicAdd(unsigned int* address, unsigned int val) +{ + return __atomic_fetch_add(address, val, __ATOMIC_RELAXED); +} +__device__ +inline +unsigned long long atomicAdd( + unsigned long long* address, unsigned long long val) +{ + return __atomic_fetch_add(address, val, __ATOMIC_RELAXED); +} +__device__ +inline +float atomicAdd(float* address, float val) +{ + unsigned int* uaddr{reinterpret_cast(uaddr)}; + unsigned int old{__atomic_load_n(uaddr, __ATOMIC_RELAXED)}; + unsigned int r; + + do { + r = old; + old = atomicCAS(uaddr, r, __float_as_uint(val + __uint_as_float(r))); + } while (r != old); + + return __uint_as_float(r); +} +__device__ +inline +double atomicAdd(double* address, double val) +{ + unsigned long long* uaddr{reinterpret_cast(uaddr)}; + unsigned long long old{__atomic_load_n(uaddr, __ATOMIC_RELAXED)}; + unsigned long long r; + + do { + r = old; + old = atomicCAS( + uaddr, r, __double_as_longlong(val + __longlong_as_double(r))); + } while (r != old); + + return __longlong_as_double(r); +} + +__device__ +inline +int atomicSub(int* address, int val) +{ + return __atomic_fetch_sub(address, val, __ATOMIC_RELAXED); +} +__device__ +inline +unsigned int atomicSub(unsigned int* address, unsigned int val) +{ + return __atomic_fetch_sub(address, val, __ATOMIC_RELAXED); +} + +__device__ +inline +int atomicExch(int* address, int val) +{ + return __atomic_exchange_n(address, val, __ATOMIC_RELAXED); +} +__device__ +inline +unsigned int atomicExch(unsigned int* address, unsigned int val) +{ + return __atomic_exchange_n(address, val, __ATOMIC_RELAXED); +} +__device__ +inline +unsigned int atomicExch(unsigned long long* address, unsigned long long val) +{ + return __atomic_exchange_n(address, val, __ATOMIC_RELAXED); +} +__device__ +inline +float atomicExch(float* address, float val) +{ + return __uint_as_float(__atomic_exchange_n( + reinterpret_cast(address), + __float_as_uint(val), + __ATOMIC_RELAXED)); +} + +__device__ +inline +int atomicMin(int* address, int val) +{ + return __sync_fetch_and_min(address, val); +} +__device__ +inline +unsigned int atomicMin(unsigned int* address, unsigned int val) +{ + return __sync_fetch_and_umin(address, val); +} +__device__ +inline +unsigned long long atomicMin( + unsigned long long* address, unsigned long long val) +{ + unsigned long long tmp{__atomic_load_n(address, __ATOMIC_RELAXED)}; + while (val < tmp) { tmp = atomicCAS(address, tmp, val); } + + return tmp; +} + +__device__ +inline +int atomicMax(int* address, int val) +{ + return __sync_fetch_and_max(address, val); +} +__device__ +inline +unsigned int atomicMax(unsigned int* address, unsigned int val) +{ + return __sync_fetch_and_umax(address, val); +} +__device__ +inline +unsigned long long atomicMax( + unsigned long long* address, unsigned long long val) +{ + unsigned long long tmp{__atomic_load_n(address, __ATOMIC_RELAXED)}; + while (tmp < val) { tmp = atomicCAS(address, tmp, val); } + + return tmp; +} + +__device__ +inline +unsigned int atomicInc(unsigned int* address, unsigned int val) +{ + extern unsigned int __builtin_amdgcn_atomic_inc( + unsigned int*, + unsigned int, + unsigned int, + unsigned int, + bool) __asm("llvm.amdgcn.atomic.inc.i32.p0i32"); + + return __builtin_amdgcn_atomic_inc( + address, val, __ATOMIC_RELAXED, 1 /* Device scope */, false); +} + +__device__ +inline +unsigned int atomicDec(unsigned int* address, unsigned int val) +{ + extern unsigned int __builtin_amdgcn_atomic_dec( + unsigned int*, + unsigned int, + unsigned int, + unsigned int, + bool) __asm("llvm.amdgcn.atomic.dec.i32.p0i32"); + + return __builtin_amdgcn_atomic_dec( + address, val, __ATOMIC_RELAXED, 1 /* Device scope */, false); +} + +__device__ +inline +int atomicAnd(int* address, int val) +{ + return __atomic_fetch_and(address, val, __ATOMIC_RELAXED); +} +__device__ +inline +unsigned int atomicAnd(unsigned int* address, unsigned int val) +{ + return __atomic_fetch_and(address, val, __ATOMIC_RELAXED); +} +__device__ +inline +unsigned long long atomicAnd( + unsigned long long* address, unsigned long long val) +{ + return __atomic_fetch_and(address, val, __ATOMIC_RELAXED); +} + +__device__ +inline +int atomicOr(int* address, int val) +{ + return __atomic_fetch_or(address, val, __ATOMIC_RELAXED); +} +__device__ +inline +unsigned int atomicOr(unsigned int* address, unsigned int val) +{ + return __atomic_fetch_or(address, val, __ATOMIC_RELAXED); +} +__device__ +inline +unsigned long long atomicOr( + unsigned long long* address, unsigned long long val) +{ + return __atomic_fetch_or(address, val, __ATOMIC_RELAXED); +} + +__device__ +inline +int atomicXor(int* address, int val) +{ + return __atomic_fetch_xor(address, val, __ATOMIC_RELAXED); +} +__device__ +inline +unsigned int atomicXor(unsigned int* address, unsigned int val) +{ + return __atomic_fetch_xor(address, val, __ATOMIC_RELAXED); +} +__device__ +inline +unsigned long long atomicXor( + unsigned long long* address, unsigned long long val) +{ + return __atomic_fetch_xor(address, val, __ATOMIC_RELAXED); +} + +// TODO: add scoped atomics i.e. atomic{*}_system && atomic{*}_block. \ No newline at end of file diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 92f06e9174..c77b4b85bd 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -88,6 +88,7 @@ extern int HIP_TRACE_API; #ifdef __cplusplus #include #endif +#include #include #include #include @@ -175,82 +176,6 @@ __device__ clock_t clock(); // abort __device__ void abort(); -// atomicAdd() -__device__ int atomicAdd(int* address, int val); -__device__ unsigned int atomicAdd(unsigned int* address, unsigned int val); - -__device__ unsigned long long int atomicAdd(unsigned long long int* address, - unsigned long long int val); - -__device__ float atomicAdd(float* address, float val); - - -// atomicSub() -__device__ int atomicSub(int* address, int val); - -__device__ unsigned int atomicSub(unsigned int* address, unsigned int val); - - -// atomicExch() -__device__ int atomicExch(int* address, int val); - -__device__ unsigned int atomicExch(unsigned int* address, unsigned int val); - -__device__ unsigned long long int atomicExch(unsigned long long int* address, - unsigned long long int val); - -__device__ float atomicExch(float* address, float val); - - -// atomicMin() -__device__ int atomicMin(int* address, int val); -__device__ unsigned int atomicMin(unsigned int* address, unsigned int val); -__device__ unsigned long long int atomicMin(unsigned long long int* address, - unsigned long long int val); - - -// atomicMax() -__device__ int atomicMax(int* address, int val); -__device__ unsigned int atomicMax(unsigned int* address, unsigned int val); -__device__ unsigned long long int atomicMax(unsigned long long int* address, - unsigned long long int val); - - -// atomicCAS() -__device__ int atomicCAS(int* address, int compare, int val); -__device__ unsigned int atomicCAS(unsigned int* address, unsigned int compare, unsigned int val); -__device__ unsigned long long int atomicCAS(unsigned long long int* address, - unsigned long long int compare, - unsigned long long int val); - - -// atomicAnd() -__device__ int atomicAnd(int* address, int val); -__device__ unsigned int atomicAnd(unsigned int* address, unsigned int val); -__device__ unsigned long long int atomicAnd(unsigned long long int* address, - unsigned long long int val); - - -// atomicOr() -__device__ int atomicOr(int* address, int val); -__device__ unsigned int atomicOr(unsigned int* address, unsigned int val); -__device__ unsigned long long int atomicOr(unsigned long long int* address, - unsigned long long int val); - - -// atomicXor() -__device__ int atomicXor(int* address, int val); -__device__ unsigned int atomicXor(unsigned int* address, unsigned int val); -__device__ unsigned long long int atomicXor(unsigned long long int* address, - unsigned long long int val); - -// atomicInc() -__device__ unsigned int atomicInc(unsigned int* address, unsigned int val); - - -// atomicDec() -__device__ unsigned int atomicDec(unsigned int* address, unsigned int val); - // warp vote function __all __any __ballot __device__ int __all(int input); __device__ int __any(int input); diff --git a/include/hip/hip_runtime.h b/include/hip/hip_runtime.h index 157fc88a43..937ba61ecf 100644 --- a/include/hip/hip_runtime.h +++ b/include/hip/hip_runtime.h @@ -64,4 +64,4 @@ THE SOFTWARE. #include #include -#endif +#endif \ No newline at end of file diff --git a/src/device_util.cpp b/src/device_util.cpp index 5ce014b2b9..11f992a510 100644 --- a/src/device_util.cpp +++ b/src/device_util.cpp @@ -761,129 +761,6 @@ __device__ clock_t clock() { return (clock_t)hc::__cycle_u64(); }; // abort __device__ void abort() { return hc::abort(); } -// atomicAdd() -__device__ int atomicAdd(int* address, int val) { return hc::atomic_fetch_add(address, val); } -__device__ unsigned int atomicAdd(unsigned int* address, unsigned int val) { - return hc::atomic_fetch_add(address, val); -} -__device__ unsigned long long int atomicAdd(unsigned long long int* address, - unsigned long long int val) { - return (long long int)hc::atomic_fetch_add((uint64_t*)address, (uint64_t)val); -} -__device__ float atomicAdd(float* address, float val) { return hc::atomic_fetch_add(address, val); } - -// atomicSub() -__device__ int atomicSub(int* address, int val) { return hc::atomic_fetch_sub(address, val); } -__device__ unsigned int atomicSub(unsigned int* address, unsigned int val) { - return hc::atomic_fetch_sub(address, val); -} - -// atomicExch() -__device__ int atomicExch(int* address, int val) { return hc::atomic_exchange(address, val); } -__device__ unsigned int atomicExch(unsigned int* address, unsigned int val) { - return hc::atomic_exchange(address, val); -} -__device__ unsigned long long int atomicExch(unsigned long long int* address, - unsigned long long int val) { - return (long long int)hc::atomic_exchange((uint64_t*)address, (uint64_t)val); -} -__device__ float atomicExch(float* address, float val) { return hc::atomic_exchange(address, val); } - -// atomicMin() -__device__ int atomicMin(int* address, int val) { return hc::atomic_fetch_min(address, val); } -__device__ unsigned int atomicMin(unsigned int* address, unsigned int val) { - return hc::atomic_fetch_min(address, val); -} -__device__ unsigned long long int atomicMin(unsigned long long int* address, - unsigned long long int val) { - return (long long int)hc::atomic_fetch_min((uint64_t*)address, (uint64_t)val); -} - -// atomicMax() -__device__ int atomicMax(int* address, int val) { return hc::atomic_fetch_max(address, val); } -__device__ unsigned int atomicMax(unsigned int* address, unsigned int val) { - return hc::atomic_fetch_max(address, val); -} -__device__ unsigned long long int atomicMax(unsigned long long int* address, - unsigned long long int val) { - return (long long int)hc::atomic_fetch_max((uint64_t*)address, (uint64_t)val); -} - -// atomicCAS() -template -__device__ T atomicCAS_impl(T* address, T compare, T val) { - // the implementation assumes the atomic is lock-free and - // has the same size as the non-atmoic equivalent type - static_assert(sizeof(T) == sizeof(std::atomic), - "size mismatch between atomic and non-atomic types"); - - union { - T* address; - std::atomic* atomic_address; - } u; - u.address = address; - - T expected = compare; - - // hcc should generate a system scope atomic CAS - std::atomic_compare_exchange_weak_explicit( - u.atomic_address, &expected, val, std::memory_order_acq_rel, std::memory_order_relaxed); - return expected; -} - -__device__ int atomicCAS(int* address, int compare, int val) { - return atomicCAS_impl(address, compare, val); -} -__device__ unsigned int atomicCAS(unsigned int* address, unsigned int compare, unsigned int val) { - return atomicCAS_impl(address, compare, val); -} -__device__ unsigned long long int atomicCAS(unsigned long long int* address, - unsigned long long int compare, - unsigned long long int val) { - return atomicCAS_impl(address, compare, val); -} - -// atomicAnd() -__device__ int atomicAnd(int* address, int val) { return hc::atomic_fetch_and(address, val); } -__device__ unsigned int atomicAnd(unsigned int* address, unsigned int val) { - return hc::atomic_fetch_and(address, val); -} -__device__ unsigned long long int atomicAnd(unsigned long long int* address, - unsigned long long int val) { - return (long long int)hc::atomic_fetch_and((uint64_t*)address, (uint64_t)val); -} - -// atomicOr() -__device__ int atomicOr(int* address, int val) { return hc::atomic_fetch_or(address, val); } -__device__ unsigned int atomicOr(unsigned int* address, unsigned int val) { - return hc::atomic_fetch_or(address, val); -} -__device__ unsigned long long int atomicOr(unsigned long long int* address, - unsigned long long int val) { - return (long long int)hc::atomic_fetch_or((uint64_t*)address, (uint64_t)val); -} - -// atomicXor() -__device__ int atomicXor(int* address, int val) { return hc::atomic_fetch_xor(address, val); } -__device__ unsigned int atomicXor(unsigned int* address, unsigned int val) { - return hc::atomic_fetch_xor(address, val); -} -__device__ unsigned long long int atomicXor(unsigned long long int* address, - unsigned long long int val) { - return (long long int)hc::atomic_fetch_xor((uint64_t*)address, (uint64_t)val); -} - -// atomicInc -__device__ unsigned int atomicInc(unsigned int* address, unsigned int val) { - return hc::__atomic_wrapinc(address, val); -} - -// atomicDec -__device__ unsigned int atomicDec(unsigned int* address, unsigned int val) { - return hc::__atomic_wrapdec(address, val); -} - - // warp vote function __all __any __ballot __device__ int __all(int input) { return hc::__all(input); } diff --git a/tests/src/deviceLib/hipSimpleAtomicsTest.cpp b/tests/src/deviceLib/hipSimpleAtomicsTest.cpp index ce1d0372ff..129d7c1aa8 100644 --- a/tests/src/deviceLib/hipSimpleAtomicsTest.cpp +++ b/tests/src/deviceLib/hipSimpleAtomicsTest.cpp @@ -23,134 +23,37 @@ THE SOFTWARE. * HIT_END */ -// includes, system -#include -#include -#include -#include - - // Includes HIP Runtime #include "hip/hip_runtime.h" #include +// includes, system +#include +#include +#include +#include +#include +#include + #define EXIT_WAIVED 2 const char* sampleName = "hipSimpleAtomicsTest"; +using namespace std; + //////////////////////////////////////////////////////////////////////////////// // Auto-Verification Code bool testResult = true; //////////////////////////////////////////////////////////////////////////////// -// Declaration, forward -void runTest(int argc, char** argv); +bool computeGoldBitwise(...) { + return true; +} -#define min(a, b) (a) < (b) ? (a) : (b) -#define max(a, b) (a) > (b) ? (a) : (b) - -int computeGold(int* gpuData, const int len) { - int val = 0; - - for (int i = 0; i < len; ++i) { - val += 10; - } - - if (val != gpuData[0]) { - printf("atomicAdd failed\n"); - return false; - } - - val = 0; - - for (int i = 0; i < len; ++i) { - val -= 10; - } - - if (val != gpuData[1]) { - printf("atomicSub failed\n"); - return false; - } - - bool found = false; - - for (int i = 0; i < len; ++i) { - // third element should be a member of [0, len) - if (i == gpuData[2]) { - found = true; - break; - } - } - - if (!found) { - printf("atomicExch failed\n"); - return false; - } - - val = -(1 << 8); - - for (int i = 0; i < len; ++i) { - // fourth element should be len-1 - val = max(val, i); - } - - if (val != gpuData[3]) { - printf("atomicMax failed\n"); - return false; - } - - val = 1 << 8; - - for (int i = 0; i < len; ++i) { - val = min(val, i); - } - - if (val != gpuData[4]) { - printf("atomicMin failed\n"); - return false; - } - - int limit = 17; - val = 0; - - for (int i = 0; i < len; ++i) { - val = (val >= limit) ? 0 : val + 1; - } - - if (val != gpuData[5]) { - printf("atomicInc failed\n"); - return false; - } - - limit = 137; - val = 0; - - for (int i = 0; i < len; ++i) { - val = ((val == 0) || (val > limit)) ? limit : val - 1; - } - - if (val != gpuData[6]) { - printf("atomicDec failed\n"); - return false; - } - - found = false; - - for (int i = 0; i < len; ++i) { - // eighth element should be a member of [0, len) - if (i == gpuData[7]) { - found = true; - break; - } - } - - if (!found) { - printf("atomicCAS failed\n"); - return false; - } - - val = 0xff; +template{}>::type* = nullptr> +bool computeGoldBitwise(T* gpuData, int len) { + T val = 0xff; for (int i = 0; i < len; ++i) { // 9th element should be 1 @@ -189,22 +92,142 @@ int computeGold(int* gpuData, const int len) { return true; } -__global__ void testKernel(hipLaunchParm lp, int* g_odata) { +template +bool computeGold(T* gpuData, int len) { + T val = 0; + + for (int i = 0; i < len; ++i) { + val += 10; + } + + if (val != gpuData[0]) { + printf("atomicAdd failed\n"); + return false; + } + + val = 0; + + for (int i = 0; i < len; ++i) { + val -= 10; + } + + if (val != gpuData[1]) { + printf("atomicSub failed\n"); + return false; + } + + bool found = false; + + for (T i = 0; i < len; ++i) { + // third element should be a member of [0, len) + if (i == gpuData[2]) { + found = true; + break; + } + } + + if (!found) { + printf("atomicExch failed\n"); + return false; + } + + val = -(1 << 8); + + for (T i = 0; i < len; ++i) { + // fourth element should be len-1 + val = max(val, i); + } + + if (val != gpuData[3]) { + printf("atomicMax failed\n"); + return false; + } + + val = 1 << 8; + + for (T i = 0; i < len; ++i) { + val = min(val, i); + } + + if (val != gpuData[4]) { + printf("atomicMin failed\n"); + return false; + } + + int limit = 17; + val = 0; + + for (int i = 0; i < len; ++i) { + val = (val >= limit) ? 0 : val + 1; + } + + if (val != gpuData[5]) { + printf("atomicInc failed\n"); + return false; + } + + limit = 137; + val = 0; + + for (int i = 0; i < len; ++i) { + val = ((val == 0) || (val > limit)) ? limit : val - 1; + } + + if (val != gpuData[6]) { + printf("atomicDec failed\n"); + return false; + } + + found = false; + + for (T i = 0; i < len; ++i) { + // eighth element should be a member of [0, len) + if (i == gpuData[7]) { + found = true; + break; + } + } + if (!found) { + printf("atomicCAS failed\n"); + return false; + } + + return computeGoldBitwise(gpuData, len); +} + +__device__ +void testKernelExch(...) {} + +template{}>::type* = nullptr> +__device__ +void testKernelExch(T* g_odata) { // access thread id - const unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x; - - // Test various atomic instructions - - // Arithmetic atomic instructions - - // Atomic addition - atomicAdd(&g_odata[0], 10); - - // Atomic subtraction (final should be 0) - atomicSub(&g_odata[1], 10); + const T tid = blockDim.x * blockIdx.x + threadIdx.x; // Atomic exchange atomicExch(&g_odata[2], tid); +} + +__device__ +void testKernelSub(...) {} + +template< + typename T, + typename enable_if< + is_same{} || is_same{}>::type* = nullptr> +void testKernelSub(T* g_odata) { + // Atomic subtraction (final should be 0) + atomicSub(&g_odata[1], 10); +} + +__device__ +void testKernelIntegral(...) {} + +template{}>::type* = nullptr> +__device__ +void testKernelIntegral(T* g_odata) { + // access thread id + const T tid = blockDim.x * blockIdx.x + threadIdx.x; // Atomic maximum atomicMax(&g_odata[3], tid); @@ -231,20 +254,21 @@ __global__ void testKernel(hipLaunchParm lp, int* g_odata) { // Atomic XOR atomicXor(&g_odata[10], tid); + + testKernelSub(g_odata); } +template +__global__ void testKernel(T* g_odata) { + // Atomic addition + atomicAdd(&g_odata[0], 10); -int main(int argc, char** argv) { - printf("%s starting...\n", sampleName); - - runTest(argc, argv); - - hipDeviceReset(); - printf("%s completed, returned %s\n", sampleName, testResult ? "OK" : "ERROR!"); - exit(testResult ? EXIT_SUCCESS : EXIT_FAILURE); + testKernelIntegral(g_odata); + testKernelExch(g_odata); } -void runTest(int argc, char** argv) { +template +void runTest() { hipDeviceProp_t deviceProp; deviceProp.major = 0; deviceProp.minor = 0; @@ -262,10 +286,10 @@ void runTest(int argc, char** argv) { unsigned int numThreads = 256; unsigned int numBlocks = 64; unsigned int numData = 11; - unsigned int memSize = sizeof(int) * numData; + unsigned int memSize = sizeof(T) * numData; // allocate mem for the result on host side - int* hOData = (int*)malloc(memSize); + T* hOData = (T*)malloc(memSize); // initialize the memory for (unsigned int i = 0; i < numData; i++) hOData[i] = 0; @@ -274,13 +298,14 @@ void runTest(int argc, char** argv) { hOData[8] = hOData[10] = 0xff; // allocate device memory for result - int* dOData; + T* dOData; hipMalloc((void**)&dOData, memSize); // copy host memory to device to initialize to zero hipMemcpy(dOData, hOData, memSize, hipMemcpyHostToDevice); // execute the kernel - hipLaunchKernel(testKernel, dim3(numBlocks), dim3(numThreads), 0, 0, dOData); + hipLaunchKernelGGL( + testKernel, dim3(numBlocks), dim3(numThreads), 0, 0, dOData); // Copy result from device to host hipMemcpy(hOData, dOData, memSize, hipMemcpyDeviceToHost); @@ -294,3 +319,18 @@ void runTest(int argc, char** argv) { passed(); } + + +int main(int argc, char** argv) { + printf("%s starting...\n", sampleName); + + runTest(); + runTest(); + runTest(); + runTest(); + runTest(); + + hipDeviceReset(); + printf("%s completed, returned %s\n", sampleName, testResult ? "OK" : "ERROR!"); + exit(testResult ? EXIT_SUCCESS : EXIT_FAILURE); +} \ No newline at end of file From 59adb5e52a015a09332f698e5eb82b98789fdb0e Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Sat, 2 Jun 2018 17:46:37 +0100 Subject: [PATCH 02/18] Add missing __device__ for forward declares. --- include/hip/hcc_detail/hip_atomic.h | 10 +++++++--- 1 file changed, 7 insertions(+), 3 deletions(-) diff --git a/include/hip/hcc_detail/hip_atomic.h b/include/hip/hcc_detail/hip_atomic.h index 4af1794ba0..1304cf74a9 100644 --- a/include/hip/hcc_detail/hip_atomic.h +++ b/include/hip/hcc_detail/hip_atomic.h @@ -176,7 +176,9 @@ __device__ inline unsigned int atomicInc(unsigned int* address, unsigned int val) { - extern unsigned int __builtin_amdgcn_atomic_inc( + __device__ + extern + unsigned int __builtin_amdgcn_atomic_inc( unsigned int*, unsigned int, unsigned int, @@ -191,7 +193,9 @@ __device__ inline unsigned int atomicDec(unsigned int* address, unsigned int val) { - extern unsigned int __builtin_amdgcn_atomic_dec( + __device__ + extern + unsigned int __builtin_amdgcn_atomic_dec( unsigned int*, unsigned int, unsigned int, @@ -262,4 +266,4 @@ unsigned long long atomicXor( return __atomic_fetch_xor(address, val, __ATOMIC_RELAXED); } -// TODO: add scoped atomics i.e. atomic{*}_system && atomic{*}_block. \ No newline at end of file +// TODO: add scoped atomics i.e. atomic{*}_system && atomic{*}_block. From 23f5feaf1315917019c9ae9fdf2911b882c73560 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Sun, 3 Jun 2018 03:03:55 +0100 Subject: [PATCH 03/18] Fix hideous typos. --- include/hip/hcc_detail/hip_atomic.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/include/hip/hcc_detail/hip_atomic.h b/include/hip/hcc_detail/hip_atomic.h index 1304cf74a9..ce4b7c9008 100644 --- a/include/hip/hcc_detail/hip_atomic.h +++ b/include/hip/hcc_detail/hip_atomic.h @@ -57,7 +57,7 @@ __device__ inline float atomicAdd(float* address, float val) { - unsigned int* uaddr{reinterpret_cast(uaddr)}; + unsigned int* uaddr{reinterpret_cast(address)}; unsigned int old{__atomic_load_n(uaddr, __ATOMIC_RELAXED)}; unsigned int r; @@ -72,7 +72,7 @@ __device__ inline double atomicAdd(double* address, double val) { - unsigned long long* uaddr{reinterpret_cast(uaddr)}; + unsigned long long* uaddr{reinterpret_cast(address)}; unsigned long long old{__atomic_load_n(uaddr, __ATOMIC_RELAXED)}; unsigned long long r; From 9141037105bb1c1841423f272262a801cefbf498 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Tue, 5 Jun 2018 08:46:57 -0400 Subject: [PATCH 04/18] Fix channel_descriptor.h about vector 3 for gcc --- include/hip/hcc_detail/channel_descriptor.h | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/include/hip/hcc_detail/channel_descriptor.h b/include/hip/hcc_detail/channel_descriptor.h index 5fb0faeee1..de290fafc1 100644 --- a/include/hip/hcc_detail/channel_descriptor.h +++ b/include/hip/hcc_detail/channel_descriptor.h @@ -92,6 +92,7 @@ inline hipChannelFormatDesc hipCreateChannelDesc() { return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindSigned); } +#ifndef __GNUC__ // vector3 is the same as vector4 template <> inline hipChannelFormatDesc hipCreateChannelDesc() { int e = (int)sizeof(unsigned char) * 8; @@ -103,6 +104,7 @@ inline hipChannelFormatDesc hipCreateChannelDesc() { int e = (int)sizeof(signed char) * 8; return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindSigned); } +#endif template <> inline hipChannelFormatDesc hipCreateChannelDesc() { @@ -152,6 +154,7 @@ inline hipChannelFormatDesc hipCreateChannelDesc() { return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindSigned); } +#ifndef __GNUC__ template <> inline hipChannelFormatDesc hipCreateChannelDesc() { int e = (int)sizeof(unsigned short) * 8; @@ -163,6 +166,7 @@ inline hipChannelFormatDesc hipCreateChannelDesc() { int e = (int)sizeof(signed short) * 8; return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindSigned); } +#endif template <> inline hipChannelFormatDesc hipCreateChannelDesc() { @@ -212,6 +216,7 @@ inline hipChannelFormatDesc hipCreateChannelDesc() { return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindSigned); } +#ifndef __GNUC__ template <> inline hipChannelFormatDesc hipCreateChannelDesc() { int e = (int)sizeof(unsigned int) * 8; @@ -223,6 +228,7 @@ inline hipChannelFormatDesc hipCreateChannelDesc() { int e = (int)sizeof(signed int) * 8; return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindSigned); } +#endif template <> inline hipChannelFormatDesc hipCreateChannelDesc() { @@ -254,11 +260,13 @@ inline hipChannelFormatDesc hipCreateChannelDesc() { return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindFloat); } +#ifndef __GNUC__ template <> inline hipChannelFormatDesc hipCreateChannelDesc() { int e = (int)sizeof(float) * 8; return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindFloat); } +#endif template <> inline hipChannelFormatDesc hipCreateChannelDesc() { @@ -302,6 +310,7 @@ inline hipChannelFormatDesc hipCreateChannelDesc() { return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindSigned); } +#ifndef __GNUC__ template <> inline hipChannelFormatDesc hipCreateChannelDesc() { int e = (int)sizeof(unsigned long) * 8; @@ -313,6 +322,7 @@ inline hipChannelFormatDesc hipCreateChannelDesc() { int e = (int)sizeof(signed long) * 8; return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindSigned); } +#endif template <> inline hipChannelFormatDesc hipCreateChannelDesc() { From d726ba1d9ac5ea19559ec481f171d5e1e434d6a6 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Sat, 2 Jun 2018 18:01:09 -0400 Subject: [PATCH 05/18] Add more function declarations for hip-clang --- include/hip/hcc_detail/hip_runtime.h | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 1a6b0f7dda..7c1f2a4d16 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -98,6 +98,9 @@ struct Empty_launch_parm {}; #include "grid_launch_GGL.hpp" #endif // GENERIC_GRID_LAUNCH +#endif // HCC + +#if __HCC_OR_HIP_CLANG__ extern int HIP_TRACE_API; #ifdef __cplusplus @@ -106,15 +109,16 @@ extern int HIP_TRACE_API; #include #include #include +#if __HCC___ #include #include +#endif // __HCC__ // TODO-HCC remove old definitions ; ~1602 hcc supports __HCC_ACCELERATOR__ define. #if defined(__KALMAR_ACCELERATOR__) && !defined(__HCC_ACCELERATOR__) #define __HCC_ACCELERATOR__ __KALMAR_ACCELERATOR__ #endif - // TODO-HCC add a dummy implementation of assert, need to replace with a proper kernel exit call. #if __HIP_DEVICE_COMPILE__ == 1 #undef assert @@ -179,10 +183,6 @@ extern int HIP_TRACE_API; #define __HCC_C__ #endif -#endif // defined __HCC__ - -#if __HCC_OR_HIP_CLANG__ - // TODO - hipify-clang - change to use the function call. //#define warpSize hc::__wavesize() static constexpr int warpSize = 64; From 17e3093f0ebe08b43f34c8e88ef277d70a37d285 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Fri, 8 Jun 2018 07:44:02 -0400 Subject: [PATCH 06/18] Add __device__ to device functions in hip_fp16_math_fwd.h --- include/hip/hcc_detail/hip_fp16_math_fwd.h | 72 +++++++++++----------- 1 file changed, 36 insertions(+), 36 deletions(-) diff --git a/include/hip/hcc_detail/hip_fp16_math_fwd.h b/include/hip/hcc_detail/hip_fp16_math_fwd.h index 81f9fe4761..fccbcfbfdc 100644 --- a/include/hip/hcc_detail/hip_fp16_math_fwd.h +++ b/include/hip/hcc_detail/hip_fp16_math_fwd.h @@ -28,49 +28,49 @@ THE SOFTWARE. extern "C" { - __attribute__((const)) _Float16 __ocml_ceil_f16(_Float16); - _Float16 __ocml_cos_f16(_Float16); - __attribute__((pure)) _Float16 __ocml_exp_f16(_Float16); - __attribute__((pure)) _Float16 __ocml_exp10_f16(_Float16); - __attribute__((pure)) _Float16 __ocml_exp2_f16(_Float16); - __attribute__((const)) _Float16 __ocml_floor_f16(_Float16); - __attribute__((const)) + __device__ __attribute__((const)) _Float16 __ocml_ceil_f16(_Float16); + __device__ _Float16 __ocml_cos_f16(_Float16); + __device__ __attribute__((pure)) _Float16 __ocml_exp_f16(_Float16); + __device__ __attribute__((pure)) _Float16 __ocml_exp10_f16(_Float16); + __device__ __attribute__((pure)) _Float16 __ocml_exp2_f16(_Float16); + __device__ __attribute__((const)) _Float16 __ocml_floor_f16(_Float16); + __device__ __attribute__((const)) _Float16 __ocml_fma_f16(_Float16, _Float16, _Float16); - __attribute__((const)) int __ocml_isinf_f16(_Float16); - __attribute__((const)) int __ocml_isnan_f16(_Float16); - __attribute__((pure)) _Float16 __ocml_log_f16(_Float16); - __attribute__((pure)) _Float16 __ocml_log10_f16(_Float16); - __attribute__((pure)) _Float16 __ocml_log2_f16(_Float16); - __attribute__((const)) _Float16 __llvm_amdgcn_rcp_f16(_Float16); - __attribute__((const)) _Float16 __ocml_rint_f16(_Float16); - __attribute__((const)) _Float16 __ocml_rsqrt_f16(_Float16); - _Float16 __ocml_sin_f16(_Float16); - __attribute__((const)) _Float16 __ocml_sqrt_f16(_Float16); - __attribute__((const)) _Float16 __ocml_trunc_f16(_Float16); + __device__ __attribute__((const)) int __ocml_isinf_f16(_Float16); + __device__ __attribute__((const)) int __ocml_isnan_f16(_Float16); + __device__ __attribute__((pure)) _Float16 __ocml_log_f16(_Float16); + __device__ __attribute__((pure)) _Float16 __ocml_log10_f16(_Float16); + __device__ __attribute__((pure)) _Float16 __ocml_log2_f16(_Float16); + __device__ __attribute__((const)) _Float16 __llvm_amdgcn_rcp_f16(_Float16); + __device__ __attribute__((const)) _Float16 __ocml_rint_f16(_Float16); + __device__ __attribute__((const)) _Float16 __ocml_rsqrt_f16(_Float16); + __device__ _Float16 __ocml_sin_f16(_Float16); + __device__ __attribute__((const)) _Float16 __ocml_sqrt_f16(_Float16); + __device__ __attribute__((const)) _Float16 __ocml_trunc_f16(_Float16); typedef _Float16 __2f16 __attribute__((ext_vector_type(2))); typedef short __2i16 __attribute__((ext_vector_type(2))); - __attribute__((const)) __2f16 __ocml_ceil_2f16(__2f16); - __2f16 __ocml_cos_2f16(__2f16); - __attribute__((pure)) __2f16 __ocml_exp_2f16(__2f16); - __attribute__((pure)) __2f16 __ocml_exp10_2f16(__2f16); - __attribute__((pure)) __2f16 __ocml_exp2_2f16(__2f16); - __attribute__((const)) __2f16 __ocml_floor_2f16(__2f16); - __attribute__((const)) __2f16 __ocml_fma_2f16(__2f16, __2f16, __2f16); - __attribute__((const)) __2i16 __ocml_isinf_2f16(__2f16); - __attribute__((const)) __2i16 __ocml_isnan_2f16(__2f16); - __attribute__((pure)) __2f16 __ocml_log_2f16(__2f16); - __attribute__((pure)) __2f16 __ocml_log10_2f16(__2f16); - __attribute__((pure)) __2f16 __ocml_log2_2f16(__2f16); - inline + __device__ __attribute__((const)) __2f16 __ocml_ceil_2f16(__2f16); + __device__ __2f16 __ocml_cos_2f16(__2f16); + __device__ __attribute__((pure)) __2f16 __ocml_exp_2f16(__2f16); + __device__ __attribute__((pure)) __2f16 __ocml_exp10_2f16(__2f16); + __device__ __attribute__((pure)) __2f16 __ocml_exp2_2f16(__2f16); + __device__ __attribute__((const)) __2f16 __ocml_floor_2f16(__2f16); + __device__ __attribute__((const)) __2f16 __ocml_fma_2f16(__2f16, __2f16, __2f16); + __device__ __attribute__((const)) __2i16 __ocml_isinf_2f16(__2f16); + __device__ __attribute__((const)) __2i16 __ocml_isnan_2f16(__2f16); + __device__ __attribute__((pure)) __2f16 __ocml_log_2f16(__2f16); + __device__ __attribute__((pure)) __2f16 __ocml_log10_2f16(__2f16); + __device__ __attribute__((pure)) __2f16 __ocml_log2_2f16(__2f16); + __device__ inline __2f16 __llvm_amdgcn_rcp_2f16(__2f16 x) // Not currently exposed by ROCDL. { return __2f16{__llvm_amdgcn_rcp_f16(x.x), __llvm_amdgcn_rcp_f16(x.y)}; } - __attribute__((const)) __2f16 __ocml_rint_2f16(__2f16); - __attribute__((const)) __2f16 __ocml_rsqrt_2f16(__2f16); - __2f16 __ocml_sin_2f16(__2f16); - __attribute__((const)) __2f16 __ocml_sqrt_2f16(__2f16); - __attribute__((const)) __2f16 __ocml_trunc_2f16(__2f16); + __device__ __attribute__((const)) __2f16 __ocml_rint_2f16(__2f16); + __device__ __attribute__((const)) __2f16 __ocml_rsqrt_2f16(__2f16); + __device__ __2f16 __ocml_sin_2f16(__2f16); + __device__ __attribute__((const)) __2f16 __ocml_sqrt_2f16(__2f16); + __device__ __attribute__((const)) __2f16 __ocml_trunc_2f16(__2f16); } \ No newline at end of file From 5eeb57b0a6b2113bf5ec5fdcb333f6834522698c Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Tue, 5 Jun 2018 08:46:20 -0400 Subject: [PATCH 07/18] Add missing macro MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT to hip_complex.h --- include/hip/hcc_detail/hip_complex.h | 3 +++ 1 file changed, 3 insertions(+) diff --git a/include/hip/hcc_detail/hip_complex.h b/include/hip/hcc_detail/hip_complex.h index 973d5f564b..22808ed576 100644 --- a/include/hip/hcc_detail/hip_complex.h +++ b/include/hip/hcc_detail/hip_complex.h @@ -94,6 +94,9 @@ THE SOFTWARE. ret.y = lhs.y * rhs; \ return ret; \ } +#define MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ComplexT, T) \ + ComplexT(T val) : x(val), y(val) {} \ + ComplexT(T val1, T val2) : x(val1), y(val2) {} #endif From 325cf3ccf0ec626bed2dccfacef6176e3e40cc35 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Thu, 31 May 2018 13:58:44 -0400 Subject: [PATCH 08/18] Include cmath instead of math.h in hip_complex.h --- include/hip/hcc_detail/hip_complex.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/hip/hcc_detail/hip_complex.h b/include/hip/hcc_detail/hip_complex.h index 22808ed576..3dc7a2ed56 100644 --- a/include/hip/hcc_detail/hip_complex.h +++ b/include/hip/hcc_detail/hip_complex.h @@ -24,7 +24,7 @@ THE SOFTWARE. #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H #include "hip/hcc_detail/hip_vector_types.h" -#include +#include #if __cplusplus #define COMPLEX_ADD_OP_OVERLOAD(type) \ From 2523c39a3706fc8d95d196a1eb4179497dda4a7e Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Wed, 6 Jun 2018 11:15:54 -0400 Subject: [PATCH 09/18] Includes or by __cplusplus in hip_complex.h --- include/hip/hcc_detail/hip_complex.h | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/include/hip/hcc_detail/hip_complex.h b/include/hip/hcc_detail/hip_complex.h index 3dc7a2ed56..72f139ba84 100644 --- a/include/hip/hcc_detail/hip_complex.h +++ b/include/hip/hcc_detail/hip_complex.h @@ -24,7 +24,15 @@ THE SOFTWARE. #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H #include "hip/hcc_detail/hip_vector_types.h" + +// TODO: Clang has a bug which allows device functions to call std functions +// when std functions are introduced into default namespace by using statement. +// math.h may be included after this bug is fixed. +#if __cplusplus #include +#else +#include "math.h" +#endif #if __cplusplus #define COMPLEX_ADD_OP_OVERLOAD(type) \ From 7a5605d006862c423710171f33641be7f54432ac Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Sun, 17 Jun 2018 12:18:05 -0400 Subject: [PATCH 10/18] Add missing __device__ __host__ to complex constructor Also add missing typedef value_type --- include/hip/hcc_detail/hip_complex.h | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/include/hip/hcc_detail/hip_complex.h b/include/hip/hcc_detail/hip_complex.h index 72f139ba84..783833dc7e 100644 --- a/include/hip/hcc_detail/hip_complex.h +++ b/include/hip/hcc_detail/hip_complex.h @@ -103,14 +103,15 @@ THE SOFTWARE. return ret; \ } #define MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ComplexT, T) \ - ComplexT(T val) : x(val), y(val) {} \ - ComplexT(T val1, T val2) : x(val1), y(val2) {} + __device__ __host__ ComplexT(T val) : x(val), y(val) {} \ + __device__ __host__ ComplexT(T val1, T val2) : x(val1), y(val2) {} #endif struct hipFloatComplex { #ifdef __cplusplus public: + typedef float value_type; __device__ __host__ hipFloatComplex() : x(0.0f), y(0.0f) {} __device__ __host__ hipFloatComplex(float x) : x(x), y(0.0f) {} __device__ __host__ hipFloatComplex(float x, float y) : x(x), y(y) {} @@ -130,6 +131,7 @@ struct hipFloatComplex { struct hipDoubleComplex { #ifdef __cplusplus public: + typedef double value_type; __device__ __host__ hipDoubleComplex() : x(0.0f), y(0.0f) {} __device__ __host__ hipDoubleComplex(double x) : x(x), y(0.0f) {} __device__ __host__ hipDoubleComplex(double x, double y) : x(x), y(y) {} From 9181fbb0b7e5f7db4df486b8600ea78092c4ffb1 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Mon, 18 Jun 2018 11:57:57 -0400 Subject: [PATCH 11/18] Add abs/real/imag functions for hipFloatComplex/hipDoubleComplex --- include/hip/hcc_detail/hip_complex.h | 9 ++++++++- 1 file changed, 8 insertions(+), 1 deletion(-) diff --git a/include/hip/hcc_detail/hip_complex.h b/include/hip/hcc_detail/hip_complex.h index 783833dc7e..7f4dd840b7 100644 --- a/include/hip/hcc_detail/hip_complex.h +++ b/include/hip/hcc_detail/hip_complex.h @@ -234,7 +234,6 @@ __device__ __host__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hi __device__ __host__ static inline float hipCabsf(hipFloatComplex z) { return sqrtf(hipCsqabsf(z)); } - __device__ __host__ static inline double hipCreal(hipDoubleComplex z) { return z.x; } __device__ __host__ static inline double hipCimag(hipDoubleComplex z) { return z.y; } @@ -314,4 +313,12 @@ __device__ __host__ static inline hipDoubleComplex hipCfma(hipDoubleComplex p, h return make_hipDoubleComplex(real, imag); } +#define __DEFINE_HIP_COMPLEX_FUN(func) \ +__device__ __host__ inline float func(const hipFloatComplex& z) { return hipC##func##f(z); } \ +__device__ __host__ inline double func(const hipDoubleComplex& z) { return hipC##func(z); } + +__DEFINE_HIP_COMPLEX_FUN(abs) +__DEFINE_HIP_COMPLEX_FUN(real) +__DEFINE_HIP_COMPLEX_FUN(imag) + #endif From 4f0aefbd6320881804feadf38c2c00f09c021c95 Mon Sep 17 00:00:00 2001 From: pradeepisro Date: Tue, 19 Jun 2018 12:26:18 +0530 Subject: [PATCH 12/18] added changes to FindHIP.cmake which would allow us to build using ninja #467 --- cmake/FindHIP.cmake | 8 +++++++- 1 file changed, 7 insertions(+), 1 deletion(-) diff --git a/cmake/FindHIP.cmake b/cmake/FindHIP.cmake index a45fb1b44d..e540ac35a6 100644 --- a/cmake/FindHIP.cmake +++ b/cmake/FindHIP.cmake @@ -471,7 +471,13 @@ macro(HIP_PREPARE_TARGET_COMMANDS _target _format _generated_files _source_files INPUT "${custom_target_script_pregen}" ) set(main_dep DEPENDS ${source_file}) - set(verbose_output "$(VERBOSE)") + if(CMAKE_GENERATOR MATCHES "Makefiles") + set(verbose_output "$(VERBOSE)") + elseif(HIP_VERBOSE_BUILD) + set(verbose_output ON) + else() + set(verbose_output OFF) + endif() # Create up the comment string file(RELATIVE_PATH generated_file_relative_path "${CMAKE_BINARY_DIR}" "${generated_file}") From f6938961b9011dd42a393fed34a868fd9bf5260f Mon Sep 17 00:00:00 2001 From: pradeepisro Date: Tue, 19 Jun 2018 13:50:24 +0530 Subject: [PATCH 13/18] added option HIP_BUILD_VERBOSE to enable verbosity in HIP build --- cmake/FindHIP.cmake | 1 + 1 file changed, 1 insertion(+) diff --git a/cmake/FindHIP.cmake b/cmake/FindHIP.cmake index e540ac35a6..b761356c9a 100644 --- a/cmake/FindHIP.cmake +++ b/cmake/FindHIP.cmake @@ -20,6 +20,7 @@ foreach(config ${_hip_configuration_types}) mark_as_advanced(HIP_HIPCC_FLAGS_${config_upper} HIP_HCC_FLAGS_${config_upper} HIP_NVCC_FLAGS_${config_upper}) endforeach() option(HIP_HOST_COMPILATION_CPP "Host code compilation mode" ON) +option(HIP_VERBOSE_BUILD "Print out the commands run while compiling the HIP source file. With the Makefile generator this defaults to VERBOSE variable specified on the command line, but can be forced on with this option." OFF) mark_as_advanced(HIP_HOST_COMPILATION_CPP) ############################################################################### From 84da72dae8b204cd4c89e78b2d99219e33011290 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Mon, 18 Jun 2018 14:35:48 -0400 Subject: [PATCH 14/18] Add conj, operator-,==,!= for hipFloatComplex/hipDoubleComplex --- include/hip/hcc_detail/hip_complex.h | 44 ++++++++++++++++++++++++---- 1 file changed, 38 insertions(+), 6 deletions(-) diff --git a/include/hip/hcc_detail/hip_complex.h b/include/hip/hcc_detail/hip_complex.h index 7f4dd840b7..48a2852a5a 100644 --- a/include/hip/hcc_detail/hip_complex.h +++ b/include/hip/hcc_detail/hip_complex.h @@ -35,6 +35,24 @@ THE SOFTWARE. #endif #if __cplusplus +#define COMPLEX_NEG_OP_OVERLOAD(type) \ + __device__ __host__ static inline type operator-(const type& op) { \ + type ret; \ + ret.x = -op.x; \ + ret.y = -op.y; \ + return ret; \ + } + +#define COMPLEX_EQ_OP_OVERLOAD(type) \ + __device__ __host__ static inline bool operator==(const type& lhs, const type& rhs) { \ + return lhs.x == rhs.x && lhs.y == rhs.y; \ + } + +#define COMPLEX_NE_OP_OVERLOAD(type) \ + __device__ __host__ static inline bool operator!=(const type& lhs, const type& rhs) { \ + return !(lhs == rhs); \ + } + #define COMPLEX_ADD_OP_OVERLOAD(type) \ __device__ __host__ static inline type operator+(const type& lhs, const type& rhs) { \ type ret; \ @@ -150,6 +168,9 @@ struct hipDoubleComplex { #if __cplusplus +COMPLEX_NEG_OP_OVERLOAD(hipFloatComplex) +COMPLEX_EQ_OP_OVERLOAD(hipFloatComplex) +COMPLEX_NE_OP_OVERLOAD(hipFloatComplex) COMPLEX_ADD_OP_OVERLOAD(hipFloatComplex) COMPLEX_SUB_OP_OVERLOAD(hipFloatComplex) COMPLEX_MUL_OP_OVERLOAD(hipFloatComplex) @@ -169,6 +190,9 @@ COMPLEX_SCALAR_PRODUCT(hipFloatComplex, double) COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed long long) COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned long long) +COMPLEX_NEG_OP_OVERLOAD(hipDoubleComplex) +COMPLEX_EQ_OP_OVERLOAD(hipDoubleComplex) +COMPLEX_NE_OP_OVERLOAD(hipDoubleComplex) COMPLEX_ADD_OP_OVERLOAD(hipDoubleComplex) COMPLEX_SUB_OP_OVERLOAD(hipDoubleComplex) COMPLEX_MUL_OP_OVERLOAD(hipDoubleComplex) @@ -313,12 +337,20 @@ __device__ __host__ static inline hipDoubleComplex hipCfma(hipDoubleComplex p, h return make_hipDoubleComplex(real, imag); } -#define __DEFINE_HIP_COMPLEX_FUN(func) \ -__device__ __host__ inline float func(const hipFloatComplex& z) { return hipC##func##f(z); } \ -__device__ __host__ inline double func(const hipDoubleComplex& z) { return hipC##func(z); } +// Complex functions returning real numbers. +#define __DEFINE_HIP_COMPLEX_REAL_FUN(func, hipFun) \ +__device__ __host__ inline float func(const hipFloatComplex& z) { return hipFun##f(z); } \ +__device__ __host__ inline double func(const hipDoubleComplex& z) { return hipFun(z); } -__DEFINE_HIP_COMPLEX_FUN(abs) -__DEFINE_HIP_COMPLEX_FUN(real) -__DEFINE_HIP_COMPLEX_FUN(imag) +__DEFINE_HIP_COMPLEX_REAL_FUN(abs, hipCabs) +__DEFINE_HIP_COMPLEX_REAL_FUN(real, hipCreal) +__DEFINE_HIP_COMPLEX_REAL_FUN(imag, hipCimag) + +// Complex functions returning complex numbers. +#define __DEFINE_HIP_COMPLEX_FUN(func, hipFun) \ +__device__ __host__ inline hipFloatComplex func(const hipFloatComplex& z) { return hipFun##f(z); } \ +__device__ __host__ inline hipDoubleComplex func(const hipDoubleComplex& z) { return hipFun(z); } + +__DEFINE_HIP_COMPLEX_FUN(conj, hipConj) #endif From 73023e0a54127e00d986508250fe05b1a5bbbd82 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Thu, 21 Jun 2018 18:12:55 +0000 Subject: [PATCH 15/18] HIPCC - Use clang if clang++ is not compiled --- bin/hipcc | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/bin/hipcc b/bin/hipcc index 03f35b27fc..1be33d55e1 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -115,8 +115,13 @@ if ($HIP_PLATFORM eq "clang") { $ROCM_PATH=$ENV{'ROCM_PATH'} // "/opt/rocm"; $HIPCC="$HIP_CLANG_PATH/clang++"; $HIPCXXFLAGS .= "-std=c++11 -I$HIP_PATH/include"; - $HIPLDFLAGS = "--hip-link --hip-device-lib-path=$DEVICE_LIB_PATH -L$HIP_PATH/lib -lhip_hcc"; + + # If $HIPCC clang++ is not compiled, use clang instead + if ( ! -e $HIPCC ) { + $HIPCC="$HIP_CLANG_PATH/clang"; + $HIPLDFLAGS = "--driver-mode=g++ " . $HIPLDFLAGS; + } } elsif ($HIP_PLATFORM eq "hcc") { $HSA_PATH=$ENV{'HSA_PATH'} // "/opt/rocm/hsa"; From 6c7a64efa2e4169a6425009f0d7354e23cb2b486 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Mon, 25 Jun 2018 19:16:27 +0100 Subject: [PATCH 16/18] Removes use of unimplemented OCML functionality. --- include/hip/hcc_detail/math_functions.h | 18 ++++++++--------- include/hip/hcc_detail/math_fwd.h | 26 ------------------------- 2 files changed, 9 insertions(+), 35 deletions(-) diff --git a/include/hip/hcc_detail/math_functions.h b/include/hip/hcc_detail/math_functions.h index a9288dc204..8f7c885fc7 100644 --- a/include/hip/hcc_detail/math_functions.h +++ b/include/hip/hcc_detail/math_functions.h @@ -508,19 +508,19 @@ inline float __fadd_rz(float x, float y) { return __ocml_add_rtz_f32(x, y); } __device__ inline -float __fdiv_rd(float x, float y) { return __ocml_div_rtp_f32(x, y); } +float __fdiv_rd(float x, float y) { return x / y; } __device__ inline -float __fdiv_rn(float x, float y) { return __ocml_div_rte_f32(x, y); } +float __fdiv_rn(float x, float y) { return x / y; } __device__ inline -float __fdiv_ru(float x, float y) { return __ocml_div_rtn_f32(x, y); } +float __fdiv_ru(float x, float y) { return x / y; } __device__ inline -float __fdiv_rz(float x, float y) { return __ocml_div_rtz_f32(x, y); } +float __fdiv_rz(float x, float y) { return x / y; } __device__ inline -float __fdividef(float x, float y) { return __ocml_div_rte_f32(x, y); } +float __fdividef(float x, float y) { return x / y; } __device__ inline float __fmaf_rd(float x, float y, float z) @@ -1028,16 +1028,16 @@ inline double __dadd_rz(double x, double y) { return __ocml_add_rtz_f64(x, y); } __device__ inline -double __ddiv_rd(double x, double y) { return __ocml_div_rtp_f64(x, y); } +double __ddiv_rd(double x, double y) { return x / y; } __device__ inline -double __ddiv_rn(double x, double y) { return __ocml_div_rte_f64(x, y); } +double __ddiv_rn(double x, double y) { return x / y; } __device__ inline -double __ddiv_ru(double x, double y) { return __ocml_div_rtn_f64(x, y); } +double __ddiv_ru(double x, double y) { return x / y; } __device__ inline -double __ddiv_rz(double x, double y) { return __ocml_div_rtz_f64(x, y); } +double __ddiv_rz(double x, double y) { return x / y; } __device__ inline double __dmul_rd(double x, double y) { return __ocml_mul_rtp_f64(x, y); } diff --git a/include/hip/hcc_detail/math_fwd.h b/include/hip/hcc_detail/math_fwd.h index ac5d6e7dc6..102714e117 100644 --- a/include/hip/hcc_detail/math_fwd.h +++ b/include/hip/hcc_detail/math_fwd.h @@ -69,8 +69,6 @@ float __ocml_cosh_f32(float); __device__ float __ocml_cospi_f32(float); __device__ -float __ocml_div_rtz_f32(float, float); -__device__ float __ocml_i0_f32(float); __device__ float __ocml_i1_f32(float); @@ -290,18 +288,6 @@ __attribute__((const)) float __ocml_mul_rtz_f32(float, float); __device__ __attribute__((const)) -float __ocml_div_rte_f32(float, float); -__device__ -__attribute__((const)) -float __ocml_div_rtn_f32(float, float); -__device__ -__attribute__((const)) -float __ocml_div_rtp_f32(float, float); -__device__ -__attribute__((const)) -float __ocml_div_rtz_f32(float, float); -__device__ -__attribute__((const)) float __ocml_sqrt_rte_f32(float); __device__ __attribute__((const)) @@ -598,18 +584,6 @@ __attribute__((const)) double __ocml_mul_rtz_f64(double, double); __device__ __attribute__((const)) -double __ocml_div_rte_f64(double, double); -__device__ -__attribute__((const)) -double __ocml_div_rtn_f64(double, double); -__device__ -__attribute__((const)) -double __ocml_div_rtp_f64(double, double); -__device__ -__attribute__((const)) -double __ocml_div_rtz_f64(double, double); -__device__ -__attribute__((const)) double __ocml_sqrt_rte_f64(double); __device__ __attribute__((const)) From 02b160491d3526dfbe1036880bf5a0356cbe62a5 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Mon, 25 Jun 2018 15:33:41 -0400 Subject: [PATCH 17/18] Include host_defines.h in hip_fp16.h since it uses __host__ __device__ attributes --- include/hip/hcc_detail/hip_fp16.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/hip/hcc_detail/hip_fp16.h b/include/hip/hcc_detail/hip_fp16.h index 8657bc30a3..777113a9de 100644 --- a/include/hip/hcc_detail/hip_fp16.h +++ b/include/hip/hcc_detail/hip_fp16.h @@ -21,7 +21,7 @@ THE SOFTWARE. */ #pragma once - +#include "hip/hcc_detail/host_defines.h" #include #if defined(__cplusplus) #include From dde875f23ba37bc4247942654e69d576fbd89b29 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Thu, 28 Jun 2018 11:19:22 +0530 Subject: [PATCH 18/18] Fix typo --- include/hip/hcc_detail/hip_runtime.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 7c1f2a4d16..ca4e73eee0 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -109,7 +109,7 @@ extern int HIP_TRACE_API; #include #include #include -#if __HCC___ +#if __HCC__ #include #include #endif // __HCC__