From c3db7b929a4d55431eebe0ccee28a522a00f0006 Mon Sep 17 00:00:00 2001 From: Sameer Sahasrabuddhe Date: Thu, 25 Apr 2019 20:41:16 +0530 Subject: [PATCH 01/17] minor cleanup: eliminate repetition [ROCm/hip commit: abb93757078a7aab18c3af80f318f76c90011a6d] --- projects/hip/src/hip_hcc.cpp | 50 +++--------------------------------- 1 file changed, 3 insertions(+), 47 deletions(-) diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index ac285e6a83..dd08f2bb2e 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -1591,63 +1591,19 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, grid_launch_parm* lp, const char* kernelNameStr) { - stream = ihipSyncAndResolveStream(stream); - lp->grid_dim.x = grid; - lp->grid_dim.y = 1; - lp->grid_dim.z = 1; - lp->group_dim.x = block.x; - lp->group_dim.y = block.y; - lp->group_dim.z = block.z; - lp->barrier_bit = barrier_bit_queue_default; - lp->launch_fence = -1; - - auto crit = stream->lockopen_preKernelCommand(); - lp->av = &(crit->_av); - lp->cf = nullptr; - ihipPrintKernelLaunch(kernelNameStr, lp, stream); - return (stream); + return ihipPreLaunchKernel(stream, dim3(grid), block, lp, kernelNameStr); } hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, grid_launch_parm* lp, const char* kernelNameStr) { - stream = ihipSyncAndResolveStream(stream); - lp->grid_dim.x = grid.x; - lp->grid_dim.y = grid.y; - lp->grid_dim.z = grid.z; - lp->group_dim.x = block; - lp->group_dim.y = 1; - lp->group_dim.z = 1; - lp->barrier_bit = barrier_bit_queue_default; - lp->launch_fence = -1; - - auto crit = stream->lockopen_preKernelCommand(); - lp->av = &(crit->_av); - lp->cf = nullptr; - ihipPrintKernelLaunch(kernelNameStr, lp, stream); - return (stream); + return ihipPreLaunchKernel(stream, grid, dim3(block), lp, kernelNameStr); } hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, grid_launch_parm* lp, const char* kernelNameStr) { - stream = ihipSyncAndResolveStream(stream); - lp->grid_dim.x = grid; - lp->grid_dim.y = 1; - lp->grid_dim.z = 1; - lp->group_dim.x = block; - lp->group_dim.y = 1; - lp->group_dim.z = 1; - lp->barrier_bit = barrier_bit_queue_default; - lp->launch_fence = -1; - - auto crit = stream->lockopen_preKernelCommand(); - lp->av = &(crit->_av); - lp->cf = nullptr; - - - ihipPrintKernelLaunch(kernelNameStr, lp, stream); - return (stream); + return ihipPreLaunchKernel(stream, dim3(grid), dim3(block), lp, kernelNameStr); } From 2b16e01c667ada126d6c7ed088aaee507359e41d Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 29 Apr 2019 15:41:08 +0300 Subject: [PATCH 02/17] [HIPIFY][doc] Update Readme.md: latest cuDNN 7.5.1.10 is supported + tested with CUDA 9.0, 9,2, 10.0 and 10.1 [ROCm/hip commit: c0705f892beb0e8c7b0811c98e6bb676a1b89be0] --- projects/hip/hipify-clang/README.md | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/projects/hip/hipify-clang/README.md b/projects/hip/hipify-clang/README.md index 0daaace2f0..86e05a9567 100644 --- a/projects/hip/hipify-clang/README.md +++ b/projects/hip/hipify-clang/README.md @@ -142,9 +142,9 @@ To run it: * Path to cuDNN should be specified by the `CUDA_DNN_ROOT_DIR` option: - - Linux: `-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.0-v7.5.0.56` + - Linux: `-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.0-v7.5.1.10` - - Windows: `-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-9.0-windows10-x64-v7.1` + - Windows: `-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-9.0-windows10-x64-v7.5.1.10` 5. Ensure [`python`](https://www.python.org/downloads) of minimum required version 2.7 is installed. @@ -178,9 +178,9 @@ To run it: On Linux the following configurations are tested: -Ubuntu 14: LLVM 5.0.0 - 6.0.1, CUDA 7.0 - 9.0, cudnn-5.0.5 - cudnn-7.5.0.56 +Ubuntu 14: LLVM 5.0.0 - 6.0.1, CUDA 7.0 - 9.0, cudnn-5.0.5 - cudnn-7.5.1.10 -Ubuntu 16-18: LLVM 8.0.0, CUDA 8.0 - 10.0, cudnn-5.1.10 - cudnn-7.5.0.56 +Ubuntu 16-18: LLVM 8.0.0, CUDA 8.0 - 10.0, cudnn-5.1.10 - cudnn-7.5.1.10 Build system for the above configurations: @@ -195,7 +195,7 @@ cmake -DCMAKE_INSTALL_PREFIX=../dist \ -DCMAKE_PREFIX_PATH=/srv/git/LLVM/8.0.0/dist \ -DCUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda-10.0 \ - -DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.0-v7.5.0.56 \ + -DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.0-v7.5.1.10 \ -DLLVM_EXTERNAL_LIT=/srv/git/LLVM/8.0.0/build/bin/llvm-lit \ .. ``` @@ -307,9 +307,9 @@ On Windows 10 the following configurations are tested: LLVM 5.0.0 - 5.0.2, CUDA 8.0, cudnn-5.1.10 - cudnn-7.1.4.18 -LLVM 6.0.0 - 6.0.1, CUDA 9.0, cudnn-7.0.5.15 - cudnn-7.5.0.56 +LLVM 6.0.0 - 6.0.1, CUDA 9.0, cudnn-7.0.5.15 - cudnn-7.5.1.10 -LLVM 7.0.0 - 8.0.0 (with patch*), CUDA 7.5 - 10.0, cudnn-7.0.5.15 - cudnn-7.5.0.56 +LLVM 7.0.0 - 8.0.0 (with patch*), CUDA 7.5 - 10.0, cudnn-7.0.5.15 - cudnn-7.5.1.10 Build system for the above configurations: @@ -326,7 +326,7 @@ cmake -DCMAKE_PREFIX_PATH=f:/LLVM/6.0.1/dist \ -DCUDA_TOOLKIT_ROOT_DIR="c:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0" \ -DCUDA_SDK_ROOT_DIR="c:/ProgramData/NVIDIA Corporation/CUDA Samples/v9.0" \ - -DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-9.0-windows10-x64-v7.1 \ + -DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-9.0-windows10-x64-v7.5.1.10 \ -DLLVM_EXTERNAL_LIT=f:/LLVM/6.0.1/build/Release/bin/llvm-lit.py \ -Thost=x64 .. From 448be14118aacf5c280947336deabedf0e969e30 Mon Sep 17 00:00:00 2001 From: Michael LIAO Date: Mon, 29 Apr 2019 14:17:11 -0400 Subject: [PATCH 03/17] [hipcc] Repace `--rpath` with `--rpath-link` - With `--rpath`, app's RUNPATH is set to the temporary library path during build. Instead, `--rpath-link` should be used to leave RUNPATH untouched. User should set LD_LIBRARY_PATH or system ld.so.conf to proper setting. [ROCm/hip commit: 7f8b80103d3163dd7435b40fbea9eb3949f7c647] --- projects/hip/bin/hipcc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index 6f020662de..fadc74c137 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/bin/hipcc @@ -191,7 +191,7 @@ if ($HIP_PLATFORM eq "clang") { $HIPCXXFLAGS .= " -std=c++11 -isystem $HIP_CLANG_INCLUDE_PATH"; $HIPLDFLAGS .= " -L$HIP_LIB_PATH"; if (not $isWindows) { - $HIPLDFLAGS .= " -Wl,--rpath=$HIP_LIB_PATH"; + $HIPLDFLAGS .= " -Wl,--rpath-link=$HIP_LIB_PATH"; $HIPLDFLAGS .= " -lhip_hcc"; } else { $HIPLDFLAGS .= " -lamdhip64"; From cc4de2bc28919d25105232e187e2835ddfbf6639 Mon Sep 17 00:00:00 2001 From: Michael LIAO Date: Mon, 29 Apr 2019 15:07:11 -0400 Subject: [PATCH 04/17] [devfunc] Re-implement ballot using AMDGCN builtins - As the signature of `amdgcn.icmp` is changed for next-gen chip, using clang builtins is portable way to hide that details. [ROCm/hip commit: a64637da2c59b0a0db4b47777dd8009b1c9d4ccb] --- projects/hip/include/hip/hcc_detail/device_functions.h | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/projects/hip/include/hip/hcc_detail/device_functions.h b/projects/hip/include/hip/hcc_detail/device_functions.h index 044ca1e634..83514ae079 100644 --- a/projects/hip/include/hip/hcc_detail/device_functions.h +++ b/projects/hip/include/hip/hcc_detail/device_functions.h @@ -736,13 +736,21 @@ int __any(int predicate) { __device__ inline unsigned long long int __ballot(int predicate) { +#if defined(__HCC__) return __llvm_amdgcn_icmp_i32(predicate, 0, ICMP_NE); +#else + return __builtin_amdgcn_uicmp(predicate, 0, ICMP_NE); +#endif } __device__ inline unsigned long long int __ballot64(int predicate) { +#if defined(__HCC__) return __llvm_amdgcn_icmp_i32(predicate, 0, ICMP_NE); +#else + return __builtin_amdgcn_uicmp(predicate, 0, ICMP_NE); +#endif } // hip.amdgcn.bc - lanemask From e637e72364656b97a846c538d53b27a36c3290cc Mon Sep 17 00:00:00 2001 From: Michael LIAO Date: Fri, 26 Apr 2019 14:51:25 -0400 Subject: [PATCH 05/17] [Device Function] Fix implementation of `__bitinsert_u64` - It's a common mistake by assuming 1 << shamt would be promoted to 64-bit, if shamt is a 64-bit integer. That's not the case. Replace that left shift to a 64-bit one to ensure it won't fall into undefined behavior. - Fix the host-side implementation as well for device function testing. [ROCm/hip commit: 9bd2d5746d31a1c5d2578319e934120a879e2506] --- projects/hip/include/hip/hcc_detail/device_functions.h | 4 ++-- projects/hip/tests/src/deviceLib/hip_bitinsert.cpp | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/device_functions.h b/projects/hip/include/hip/hcc_detail/device_functions.h index 044ca1e634..e75c0b6526 100644 --- a/projects/hip/include/hip/hcc_detail/device_functions.h +++ b/projects/hip/include/hip/hcc_detail/device_functions.h @@ -103,7 +103,7 @@ __device__ static inline unsigned int __bitinsert_u32(unsigned int src0, unsigne __device__ static inline uint64_t __bitinsert_u64(uint64_t src0, uint64_t src1, unsigned int src2, unsigned int src3) { uint64_t offset = src2 & 63; uint64_t width = src3 & 63; - uint64_t mask = (1 << width) - 1; + uint64_t mask = (1ULL << width) - 1; return ((src0 & ~(mask << offset)) | ((src1 & mask) << offset)); } @@ -753,7 +753,7 @@ int64_t __lanemask_gt() int32_t activelane = __ockl_activelane_u32(); int64_t ballot = __ballot64(1); if (activelane != 63) { - int64_t tmp = (~0UL) << (activelane + 1); + int64_t tmp = (~0ULL) << (activelane + 1); return tmp & ballot; } return 0; diff --git a/projects/hip/tests/src/deviceLib/hip_bitinsert.cpp b/projects/hip/tests/src/deviceLib/hip_bitinsert.cpp index bf00f4143a..bf1146b91a 100644 --- a/projects/hip/tests/src/deviceLib/hip_bitinsert.cpp +++ b/projects/hip/tests/src/deviceLib/hip_bitinsert.cpp @@ -46,7 +46,7 @@ T bit_insert(T src0, T src1, unsigned int src2, unsigned int src3) { unsigned int bits = sizeof(T) * 8; T offset = src2 & (bits - 1); T width = src3 & (bits - 1); - T mask = (1 << width) - 1; + T mask = (((T)1) << width) - 1; return ((src0 & ~(mask << offset)) | ((src1 & mask) << offset)); } From 1f57ee6125cb90dd05b53b53dc668bd351316ab9 Mon Sep 17 00:00:00 2001 From: emankov Date: Tue, 30 Apr 2019 17:37:34 +0300 Subject: [PATCH 06/17] [HIPIFY][tests] Add cuSPARSE CSR-BCSR-SPMV-conversions example [ROCm/hip commit: e3082f514242f3f289b9d7df6fdc37621a5cdf9d] --- .../libraries/cuSPARSE/cuSPARSE_12.cu | 411 ++++++++++++++++++ 1 file changed, 411 insertions(+) create mode 100644 projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_12.cu diff --git a/projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_12.cu b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_12.cu new file mode 100644 index 0000000000..23254750f4 --- /dev/null +++ b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_12.cu @@ -0,0 +1,411 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// CHECK: #include +// CHECK: #include +// CHECK: #include +#include +#include +#include +#include +#include +#include + +#define Min(x,y) ((x)<(y)?(x):(y)) +#define Max(x,y) ((x)>(y)?(x):(y)) +#define Abs(x) ((x)>(0)?(x):-(x)) +// CHECK: static void CudaCheckCore(hipError_t code, const char *file, int line) { +static void CudaCheckCore(cudaError_t code, const char *file, int line) { + // CHECK: if (code != hipSuccess) { + if (code != cudaSuccess) { + // CHECK: fprintf(stderr,"Cuda Error %d : %s %s %d\n", code, hipGetErrorString(code), file, line); + fprintf(stderr,"Cuda Error %d : %s %s %d\n", code, cudaGetErrorString(code), file, line); + exit(code); + } +} + +#define CudaCheck( test ) { CudaCheckCore((test), __FILE__, __LINE__); } +// CHECK: #define CudaCheckAfterCall() { CudaCheckCore((hipGetLastError()), __FILE__, __LINE__); } +#define CudaCheckAfterCall() { CudaCheckCore((cudaGetLastError()), __FILE__, __LINE__); } + +// CHECK: static const char * cusparseGetErrorString(hipsparseStatus_t error) { +static const char * cusparseGetErrorString(cusparseStatus_t error) { + switch (error) { + // CHECK: case HIPSPARSE_STATUS_SUCCESS: + case CUSPARSE_STATUS_SUCCESS: + return "The operation completed successfully."; + // CHECK: case HIPSPARSE_STATUS_NOT_INITIALIZED: + case CUSPARSE_STATUS_NOT_INITIALIZED: + return "The cuSPARSE library was not initialized. This is usually caused by the lack of a prior call, an error in the CUDA Runtime API called by the cuSPARSE routine, or an error in the hardware setup.\n" \ + "To correct: call cusparseCreate() prior to the function call; and check that the hardware, an appropriate version of the driver, and the cuSPARSE library are correctly installed."; + // CHECK: case HIPSPARSE_STATUS_ALLOC_FAILED: + case CUSPARSE_STATUS_ALLOC_FAILED: + return "Resource allocation failed inside the cuSPARSE library. This is usually caused by a cudaMalloc() failure.\n"\ + "To correct: prior to the function call, deallocate previously allocated memory as much as possible."; + // CHECK: case HIPSPARSE_STATUS_INVALID_VALUE: + case CUSPARSE_STATUS_INVALID_VALUE: + return "An unsupported value or parameter was passed to the function (a negative vector size, for example).\n"\ + "To correct: ensure that all the parameters being passed have valid values."; + // CHECK: case HIPSPARSE_STATUS_ARCH_MISMATCH: + case CUSPARSE_STATUS_ARCH_MISMATCH: + return "The function requires a feature absent from the device architecture; usually caused by the lack of support for atomic operations or double precision.\n"\ + "To correct: compile and run the application on a device with appropriate compute capability, which is 1.1 for 32-bit atomic operations and 1.3 for double precision."; + // CHECK: case HIPSPARSE_STATUS_MAPPING_ERROR: + case CUSPARSE_STATUS_MAPPING_ERROR: + return "An access to GPU memory space failed, which is usually caused by a failure to bind a texture.\n"\ + "To correct: prior to the function call, unbind any previously bound textures."; + // CHECK: case HIPSPARSE_STATUS_EXECUTION_FAILED: + case CUSPARSE_STATUS_EXECUTION_FAILED: + return "The GPU program failed to execute. This is often caused by a launch failure of the kernel on the GPU, which can be caused by multiple reasons.\n"\ + "To correct: check that the hardware, an appropriate version of the driver, and the cuSPARSE library are correctly installed."; + // CHECK: case HIPSPARSE_STATUS_INTERNAL_ERROR: + case CUSPARSE_STATUS_INTERNAL_ERROR: + return "An internal cuSPARSE operation failed. This error is usually caused by a cudaMemcpyAsync() failure.\n"\ + "To correct: check that the hardware, an appropriate version of the driver, and the cuSPARSE library are correctly installed. Also, check that the memory passed as a parameter to the routine is not being deallocated prior to the routine’s completion."; + // CHECK: case HIPSPARSE_STATUS_MATRIX_TYPE_NOT_SUPPORTED: + // CHECK: "To correct: check that the fields in hipsparseMatDescr_t descrA were set correctly."; + case CUSPARSE_STATUS_MATRIX_TYPE_NOT_SUPPORTED: + return "The matrix type is not supported by this function. This is usually caused by passing an invalid matrix descriptor to the function.\n"\ + "To correct: check that the fields in cusparseMatDescr_t descrA were set correctly."; + } + return ""; +} + +// CHECK: static void CudaSparseCheckCore(hipsparseStatus_t code, const char *file, int line) { +static void CudaSparseCheckCore(cusparseStatus_t code, const char *file, int line) { + // CHECK: if (code != HIPSPARSE_STATUS_SUCCESS) { + if (code != CUSPARSE_STATUS_SUCCESS) { + fprintf(stderr,"Cuda Error %d : %s %s %d\n", code, cusparseGetErrorString(code), file, line); + exit(code); + } +} + +#define CudaSparseCheck( test ) { CudaSparseCheckCore((test), __FILE__, __LINE__); } + +// Alloc and copy +template +ObjectType* allocAndCopy(const ObjectType src[], const int size) { + ObjectType* dest = NULL; + // CHECK: CudaCheck( hipMalloc(&dest,size*sizeof(ObjectType)) ); + CudaCheck( cudaMalloc(&dest,size*sizeof(ObjectType)) ); + // CHECK: CudaCheck( hipMemcpy(dest, src, size*sizeof(ObjectType), hipMemcpyHostToDevice ) ); + CudaCheck( cudaMemcpy(dest, src, size*sizeof(ObjectType), cudaMemcpyHostToDevice ) ); + return dest; +} + +template +ObjectType* alloc(const int size) { + ObjectType* dest = NULL; + // CHECK: CudaCheck( hipMalloc(&dest,size*sizeof(ObjectType)) ); + CudaCheck( cudaMalloc(&dest,size*sizeof(ObjectType)) ); + return dest; +} + +template +ObjectType* allocAndCopyPart(const ObjectType src[], const int size, const int allocSize) { + ObjectType* dest = NULL; + assert(size <= allocSize); + // CHECK: CudaCheck( hipMalloc(&dest,allocSize*sizeof(ObjectType)) ); + // CHECK: CudaCheck( hipMemcpy(dest, src, size*sizeof(ObjectType), hipMemcpyHostToDevice ) ); + // CHECK: CudaCheck( hipMemset(&dest[size],0,(allocSize-size)*sizeof(ObjectType)) ); + CudaCheck( cudaMalloc(&dest,allocSize*sizeof(ObjectType)) ); + CudaCheck( cudaMemcpy(dest, src, size*sizeof(ObjectType), cudaMemcpyHostToDevice ) ); + CudaCheck( cudaMemset(&dest[size],0,(allocSize-size)*sizeof(ObjectType)) ); + return dest; +} + +// COO part +#include + +struct Ijv { + int i, j; + double v; +}; + +bool IjvComp(const Ijv& v1, const Ijv& v2) { + return v1.i < v2.i || (v1.i == v2.i && v1.j < v2.j); +} + +struct COOArrays { + int m; + int nnz; + double *val;/*values(NNZ)*/ + int *rowind;/*i(NNZ)*/ + int *colind;/*j(NNZ)*/ + + COOArrays() { + val = NULL; + rowind = NULL; + colind = NULL; + } + + ~COOArrays() { + delete[] val; + delete[] rowind; + delete[] colind; + } + + void sortToRowMajor() { + Ijv* ijvs = new Ijv[nnz]; + for(int idxCopy = 0 ; idxCopy < nnz ; ++idxCopy){ + ijvs[idxCopy].i = rowind[idxCopy]; + ijvs[idxCopy].j = colind[idxCopy]; + ijvs[idxCopy].v = val[idxCopy]; + } + std::sort(ijvs, ijvs+nnz, IjvComp); + for(int idxCopy = 0 ; idxCopy < nnz ; ++idxCopy){ + rowind[idxCopy] = ijvs[idxCopy].i; + colind[idxCopy] = ijvs[idxCopy].j; + val[idxCopy] = ijvs[idxCopy].v; + } + delete[] ijvs; + } +}; + +void compute_COO(COOArrays& coo, double *x , double *y ) { + for(int idxVal = 0 ; idxVal < coo.nnz ; ++idxVal){ + y[coo.rowind[idxVal]] += x[coo.colind[idxVal]] * coo.val[idxVal]; + } +} + +// COO part +struct CRSArrays { + int m; //< the dim of the matrix + int nnz;//< the number of nnz (== ia[m]) + double *cu_csrValA; //< the values (of size NNZ) + int *cu_csrRowPtrA;//< the usual rowptr (of size m+1) + int *cu_csrColIndA;//< the colidx of each NNZ (of size nnz) + // CHECK: hipStream_t streamId; + // CHECK: hipsparseHandle_t cusparseHandle; + cudaStream_t streamId; + cusparseHandle_t cusparseHandle; + + CRSArrays() { + cu_csrValA = NULL; + cu_csrRowPtrA = NULL; + cu_csrColIndA = NULL; + // Create sparse handle (needed to call sparse functions + streamId = 0; + // CHECK-NOT: hipsparseHandle = 0; + cusparseHandle = 0; + // CHECK: CudaSparseCheck(hipsparseCreate(&cusparseHandle)); + // CHECK: CudaSparseCheck(hipsparseSetStream(cusparseHandle, streamId)); + CudaSparseCheck(cusparseCreate(&cusparseHandle)); + CudaSparseCheck(cusparseSetStream(cusparseHandle, streamId)); + } + + ~CRSArrays() { + // CHECK: CudaCheck(hipFree(cu_csrValA)); + // CHECK: CudaCheck(hipFree(cu_csrRowPtrA)); + // CHECK: CudaCheck(hipFree(cu_csrColIndA)); + CudaCheck(cudaFree(cu_csrValA)); + CudaCheck(cudaFree(cu_csrRowPtrA)); + CudaCheck(cudaFree(cu_csrColIndA)); + // Destroy sparse handle + // CHECK: CudaSparseCheck(hipsparseDestroy(cusparseHandle)); + CudaSparseCheck(cusparseDestroy(cusparseHandle)); + } +}; + +void COO_to_CRS(COOArrays& coo, CRSArrays* crs) { + // We need COO to be sorted by row (and column) + coo.sortToRowMajor(); + crs->m = coo.m; + crs->nnz = coo.nnz; + // Convert COO to CSR (it is just for the rows idx) + crs->cu_csrRowPtrA = alloc(coo.m+1); + { + int* cu_cooRowIndA = allocAndCopy(coo.rowind, coo.nnz); + // CHECK: CudaSparseCheck(hipsparseXcoo2csr(crs->cusparseHandle, cu_cooRowIndA, + // CHECK: coo.nnz, coo.m, crs->cu_csrRowPtrA, HIPSPARSE_INDEX_BASE_ZERO)); + CudaSparseCheck(cusparseXcoo2csr(crs->cusparseHandle, cu_cooRowIndA, + coo.nnz, coo.m, crs->cu_csrRowPtrA, CUSPARSE_INDEX_BASE_ZERO)); + // CHECK: CudaCheck(hipFree(cu_cooRowIndA)); + CudaCheck(cudaFree(cu_cooRowIndA)); + } + // Copy cols idx and values that are unchanged + crs->cu_csrValA = allocAndCopy(coo.val, coo.nnz); + crs->cu_csrColIndA = allocAndCopy(coo.colind, coo.nnz); +} + +double compute_CRS( CRSArrays& crs, double *x , double *y) { + // For blas 2 gemv y = alpha.x.A + Beta.y + const double alpha = 1.0; + const double beta = 0.0; + // Copy input + double* cu_x = allocAndCopy(x, crs.m); + double* cu_y = allocAndCopy(y, crs.m); + // Init matrix properties + // CHECK: hipsparseMatDescr_t descr = 0; + cusparseMatDescr_t descr = 0; + // CHECK: CudaSparseCheck(hipsparseCreateMatDescr(&descr)); + CudaSparseCheck(cusparseCreateMatDescr(&descr)); + // CHECK: hipsparseSetMatType(descr,HIPSPARSE_MATRIX_TYPE_GENERAL); + cusparseSetMatType(descr,CUSPARSE_MATRIX_TYPE_GENERAL); + // CHECK: hipsparseSetMatIndexBase(descr,HIPSPARSE_INDEX_BASE_ZERO); + cusparseSetMatIndexBase(descr,CUSPARSE_INDEX_BASE_ZERO); + // Compute gemv + float gemvComputeTume = 0; + { + // CHECK: hipEvent_t startTime, stopTime; + // CHECK: hipEventCreate(&startTime); + // CHECK: hipEventCreate(&stopTime); + // CHECK: hipEventRecord(startTime, crs.streamId); + cudaEvent_t startTime, stopTime; + cudaEventCreate(&startTime); + cudaEventCreate(&stopTime); + cudaEventRecord(startTime, crs.streamId); + // CHECK: CudaSparseCheck(hipsparseDcsrmv(crs.cusparseHandle, HIPSPARSE_OPERATION_NON_TRANSPOSE, + CudaSparseCheck(cusparseDcsrmv(crs.cusparseHandle, CUSPARSE_OPERATION_NON_TRANSPOSE, + crs.m, crs.m, crs.nnz, &alpha, + descr, crs.cu_csrValA, crs.cu_csrRowPtrA, + crs.cu_csrColIndA, cu_x, &beta, cu_y)); + // CHECK: hipEventRecord(stopTime, crs.streamId); + // CHECK: hipEventSynchronize(stopTime); + // CHECK: hipEventElapsedTime(&gemvComputeTume, startTime, stopTime); + cudaEventRecord(stopTime, crs.streamId); + cudaEventSynchronize(stopTime); + cudaEventElapsedTime(&gemvComputeTume, startTime, stopTime); + gemvComputeTume /=1000.0; + } + // Get back result + // CHECK: CudaCheck( hipMemcpy(y, cu_y, crs.m*sizeof(double), hipMemcpyDeviceToHost ) ); + CudaCheck( cudaMemcpy(y, cu_y, crs.m*sizeof(double), cudaMemcpyDeviceToHost ) ); + // Dealloc vectors + // CHECK: CudaCheck(hipFree(cu_x)); + // CHECK: CudaCheck(hipFree(cu_y)); + CudaCheck(cudaFree(cu_x)); + CudaCheck(cudaFree(cu_y)); + return gemvComputeTume; +} + +// BCSR part +struct BCRSArrays { + int m; + int nnz; + int nbBlocks; + int nbBlockRow; + int blockSize; + int* cu_bsrRowPtrC; + int* cu_bsrColIndC; + double* cu_bsrValC; + // CHECK: hipStream_t streamId; + cudaStream_t streamId; + // CHECK: hipsparseHandle_t cusparseHandle; + cusparseHandle_t cusparseHandle; + + BCRSArrays() { + cu_bsrRowPtrC = NULL; + cu_bsrColIndC = NULL; + cu_bsrValC = NULL; + // Create sparse handle (needed to call sparse functions + streamId = 0; + // CHECK: CudaSparseCheck(hipsparseCreate(&cusparseHandle)); + // CHECK: CudaSparseCheck(hipsparseSetStream(cusparseHandle, streamId)); + CudaSparseCheck(cusparseCreate(&cusparseHandle)); + CudaSparseCheck(cusparseSetStream(cusparseHandle, streamId)); + } + + ~BCRSArrays() { + // CHECK: CudaCheck(hipFree(cu_bsrRowPtrC)); + // CHECK: CudaCheck(hipFree(cu_bsrColIndC)); + // CHECK: CudaCheck(hipFree(cu_bsrValC)); + CudaCheck(cudaFree(cu_bsrRowPtrC)); + CudaCheck(cudaFree(cu_bsrColIndC)); + CudaCheck(cudaFree(cu_bsrValC)); + // Destroy sparse handle + // CHECK: CudaSparseCheck(hipsparseDestroy(cusparseHandle)); + CudaSparseCheck(cusparseDestroy(cusparseHandle)); + } +}; + +void CRS_to_BCRS(CRSArrays& csr, BCRSArrays* bcrs, const int blockSize) { + bcrs->m = csr.m; + bcrs->nnz = csr.nnz; + bcrs->blockSize = blockSize; + bcrs->nbBlockRow = (csr.m + blockSize-1)/blockSize; + // CHECK: hipMalloc((void**)&bcrs->cu_bsrRowPtrC, sizeof(int) *(bcrs->nbBlockRow+1)); + cudaMalloc((void**)&bcrs->cu_bsrRowPtrC, sizeof(int) *(bcrs->nbBlockRow+1)); + // CHECK: hipsparseMatDescr_t descr = 0; + cusparseMatDescr_t descr = 0; + // CHECK: CudaSparseCheck(hipsparseCreateMatDescr(&descr)); + // CHECK: hipsparseSetMatType(descr,HIPSPARSE_MATRIX_TYPE_GENERAL); + // CHECK: hipsparseSetMatIndexBase(descr,HIPSPARSE_INDEX_BASE_ZERO); + CudaSparseCheck(cusparseCreateMatDescr(&descr)); + cusparseSetMatType(descr,CUSPARSE_MATRIX_TYPE_GENERAL); + cusparseSetMatIndexBase(descr,CUSPARSE_INDEX_BASE_ZERO); + int nbNnzBlocks; + // NOTE: cusparseXcsr2bsrNnz and CUSPARSE_DIRECTION_COLUMN (of type cusparseDirection_t) are yet unsupported by HIP + // CHECK-NOT: hipsparseXcsr2bsrNnz(bcrs->cusparseHandle, HIPSPARSE_DIRECTION_COLUMN, csr.m, csr.m, descr, csr.cu_csrRowPtrA, csr.cu_csrColIndA, + cusparseXcsr2bsrNnz(bcrs->cusparseHandle, CUSPARSE_DIRECTION_COLUMN, csr.m, csr.m, descr, csr.cu_csrRowPtrA, csr.cu_csrColIndA, + blockSize, descr, bcrs->cu_bsrRowPtrC, &nbNnzBlocks); + { + int firstBlockIdx, lastBlockIdx; + // CHECK: hipMemcpy(&lastBlockIdx, bcrs->cu_bsrRowPtrC+bcrs->nbBlockRow, sizeof(int), hipMemcpyDeviceToHost); + // CHECK: hipMemcpy(&firstBlockIdx, bcrs->cu_bsrRowPtrC, sizeof(int), hipMemcpyDeviceToHost); + cudaMemcpy(&lastBlockIdx, bcrs->cu_bsrRowPtrC+bcrs->nbBlockRow, sizeof(int), cudaMemcpyDeviceToHost); + cudaMemcpy(&firstBlockIdx, bcrs->cu_bsrRowPtrC, sizeof(int), cudaMemcpyDeviceToHost); + assert(firstBlockIdx == 0); // we are in base 0 + assert(nbNnzBlocks == lastBlockIdx - firstBlockIdx); + } + bcrs->nbBlocks = nbNnzBlocks; + // CHECK: CudaCheck(hipMalloc((void**)&bcrs->cu_bsrColIndC, sizeof(int)*nbNnzBlocks)); + // CHECK: CudaCheck(hipMalloc((void**)&bcrs->cu_bsrValC, sizeof(double)*(blockSize*blockSize)*nbNnzBlocks)); + CudaCheck(cudaMalloc((void**)&bcrs->cu_bsrColIndC, sizeof(int)*nbNnzBlocks)); + CudaCheck(cudaMalloc((void**)&bcrs->cu_bsrValC, sizeof(double)*(blockSize*blockSize)*nbNnzBlocks)); + // NOTE: cusparseDcsr2bsr and CUSPARSE_DIRECTION_COLUMN (of type cusparseDirection_t) are yet unsupported by HIP + // CHECK-NOT: hipsparseDcsr2bsr(bcrs->cusparseHandle, HIPSPARSE_DIRECTION_COLUMN, + cusparseDcsr2bsr(bcrs->cusparseHandle, CUSPARSE_DIRECTION_COLUMN, + csr.m, csr.m, descr, csr.cu_csrValA, csr.cu_csrRowPtrA, csr.cu_csrColIndA, blockSize, descr, bcrs->cu_bsrValC, bcrs->cu_bsrRowPtrC, bcrs->cu_bsrColIndC); +} + +double compute_BSR(BCRSArrays& bcsr, double *x , double *y){ + // For blas 2 gemv y = alpha.x.A + Beta.y + const double alpha = 1.0; + const double beta = 0.0; + // Copy input + const int sizeMultipleBlockSize = ((bcsr.m+bcsr.blockSize-1)/bcsr.blockSize)*bcsr.blockSize; + double* cu_x = allocAndCopyPart(x, bcsr.m, sizeMultipleBlockSize); + double* cu_y = allocAndCopyPart(y, bcsr.m, sizeMultipleBlockSize); + // Init matrix properties + // CHECK: hipsparseMatDescr_t descr = 0; + // CHECK: CudaSparseCheck(hipsparseCreateMatDescr(&descr)); + // CHECK: hipsparseSetMatType(descr,HIPSPARSE_MATRIX_TYPE_GENERAL); + // CHECK: hipsparseSetMatIndexBase(descr,HIPSPARSE_INDEX_BASE_ZERO); + cusparseMatDescr_t descr = 0; + CudaSparseCheck(cusparseCreateMatDescr(&descr)); + cusparseSetMatType(descr,CUSPARSE_MATRIX_TYPE_GENERAL); + cusparseSetMatIndexBase(descr,CUSPARSE_INDEX_BASE_ZERO); + // Compute gemv + float gemvComputeTume = 0; + { + // CHECK: hipEvent_t startTime, stopTime; + // CHECK: hipEventCreate(&startTime); + // CHECK: hipEventCreate(&stopTime); + // CHECK: hipEventRecord(startTime, bcsr.streamId); + cudaEvent_t startTime, stopTime; + cudaEventCreate(&startTime); + cudaEventCreate(&stopTime); + cudaEventRecord(startTime, bcsr.streamId); + // NOTE: cusparseDbsrmv and CUSPARSE_DIRECTION_COLUMN (of type cusparseDirection_t) are yet unsupported by HIP + // CHECK: cusparseDbsrmv(bcsr.cusparseHandle, CUSPARSE_DIRECTION_COLUMN, HIPSPARSE_OPERATION_NON_TRANSPOSE, + cusparseDbsrmv(bcsr.cusparseHandle, CUSPARSE_DIRECTION_COLUMN, CUSPARSE_OPERATION_NON_TRANSPOSE, + bcsr.nbBlockRow, bcsr.m, bcsr.nbBlocks, &alpha, descr, + bcsr.cu_bsrValC, bcsr.cu_bsrRowPtrC, bcsr.cu_bsrColIndC, bcsr.blockSize, + cu_x, &beta, cu_y); + // CHECK: hipEventRecord(stopTime, bcsr.streamId); + // CHECK: hipEventSynchronize(stopTime); + // CHECK: hipEventElapsedTime(&gemvComputeTume, startTime, stopTime); + cudaEventRecord(stopTime, bcsr.streamId); + cudaEventSynchronize(stopTime); + cudaEventElapsedTime(&gemvComputeTume, startTime, stopTime); + gemvComputeTume /=1000.0; + } + // Get back result + // CHECK: CudaCheck( hipMemcpy(y, cu_y, bcsr.m*sizeof(double), hipMemcpyDeviceToHost ) ); + CudaCheck( cudaMemcpy(y, cu_y, bcsr.m*sizeof(double), cudaMemcpyDeviceToHost ) ); + // Dealloc vectors + // CHECK: CudaCheck(hipFree(cu_x)); + // CHECK: CudaCheck(hipFree(cu_y)); + CudaCheck(cudaFree(cu_x)); + CudaCheck(cudaFree(cu_y)); + return gemvComputeTume; +} From 2c2b897902e58ff351d740dfbe9b87f03c77e61a Mon Sep 17 00:00:00 2001 From: Michael LIAO Date: Wed, 1 May 2019 22:26:35 -0400 Subject: [PATCH 07/17] [devfunc] Add necessary `__device__` and `__host__` attributes. - Minor clean up to keep consistent function declaration. [ROCm/hip commit: 44692978fe35f8467b62eb9eae02f2582c93ae4b] --- .../include/hip/hcc_detail/hip_vector_types.h | 156 ++++++++---------- 1 file changed, 65 insertions(+), 91 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_vector_types.h b/projects/hip/include/hip/hcc_detail/hip_vector_types.h index 1df6385fa6..d19a0e0e3f 100644 --- a/projects/hip/include/hip/hcc_detail/hip_vector_types.h +++ b/projects/hip/include/hip/hcc_detail/hip_vector_types.h @@ -108,13 +108,13 @@ THE SOFTWARE. using HIP_vector_base::data; using typename HIP_vector_base::Native_vec_; - __host__ __device__ + inline __host__ __device__ HIP_vector_type() = default; template< typename U, typename std::enable_if< std::is_convertible{}>::type* = nullptr> - __host__ __device__ + inline __host__ __device__ HIP_vector_type(U x) noexcept { for (auto i = 0u; i != rank; ++i) data[i] = x; @@ -123,52 +123,52 @@ THE SOFTWARE. typename... Us, typename std::enable_if< (rank > 1) && sizeof...(Us) == rank>::type* = nullptr> - __host__ __device__ + inline __host__ __device__ HIP_vector_type(Us... xs) noexcept { data = Native_vec_{static_cast(xs)...}; } - __host__ __device__ + inline __host__ __device__ HIP_vector_type(const HIP_vector_type&) = default; - __host__ __device__ + inline __host__ __device__ HIP_vector_type(HIP_vector_type&&) = default; - __host__ __device__ + inline __host__ __device__ ~HIP_vector_type() = default; - __host__ __device__ + inline __host__ __device__ HIP_vector_type& operator=(const HIP_vector_type&) = default; - __host__ __device__ + inline __host__ __device__ HIP_vector_type& operator=(HIP_vector_type&&) = default; // Operators - __host__ __device__ + inline __host__ __device__ HIP_vector_type& operator++() noexcept { return *this += HIP_vector_type{1}; } - __host__ __device__ + inline __host__ __device__ HIP_vector_type operator++(int) noexcept { auto tmp(*this); ++*this; return tmp; } - __host__ __device__ + inline __host__ __device__ HIP_vector_type& operator--() noexcept { return *this -= HIP_vector_type{1}; } - __host__ __device__ + inline __host__ __device__ HIP_vector_type operator--(int) noexcept { auto tmp(*this); --*this; return tmp; } - __host__ __device__ + inline __host__ __device__ HIP_vector_type& operator+=(const HIP_vector_type& x) noexcept { data += x.data; return *this; } - __host__ __device__ + inline __host__ __device__ HIP_vector_type& operator-=(const HIP_vector_type& x) noexcept { data -= x.data; @@ -178,18 +178,18 @@ THE SOFTWARE. typename U, typename std::enable_if< std::is_convertible{}>::type* = nullptr> - __host__ __device__ + inline __host__ __device__ HIP_vector_type& operator-=(U x) noexcept { return *this -= HIP_vector_type{x}; } - __host__ __device__ + inline __host__ __device__ HIP_vector_type& operator*=(const HIP_vector_type& x) noexcept { data *= x.data; return *this; } - __host__ __device__ + inline __host__ __device__ HIP_vector_type& operator/=(const HIP_vector_type& x) noexcept { data /= x.data; @@ -199,7 +199,7 @@ THE SOFTWARE. template< typename U = T, typename std::enable_if{}>::type* = nullptr> - __host__ __device__ + inline __host__ __device__ HIP_vector_type operator-() noexcept { auto tmp(*this); @@ -210,7 +210,7 @@ THE SOFTWARE. template< typename U = T, typename std::enable_if{}>::type* = nullptr> - __host__ __device__ + inline __host__ __device__ HIP_vector_type operator~() noexcept { HIP_vector_type r{*this}; @@ -220,7 +220,7 @@ THE SOFTWARE. template< typename U = T, typename std::enable_if{}>::type* = nullptr> - __host__ __device__ + inline __host__ __device__ HIP_vector_type& operator%=(const HIP_vector_type& x) noexcept { data %= x.data; @@ -229,7 +229,7 @@ THE SOFTWARE. template< typename U = T, typename std::enable_if{}>::type* = nullptr> - __host__ __device__ + inline __host__ __device__ HIP_vector_type& operator^=(const HIP_vector_type& x) noexcept { data ^= x.data; @@ -238,7 +238,7 @@ THE SOFTWARE. template< typename U = T, typename std::enable_if{}>::type* = nullptr> - __host__ __device__ + inline __host__ __device__ HIP_vector_type& operator|=(const HIP_vector_type& x) noexcept { data |= x.data; @@ -247,7 +247,7 @@ THE SOFTWARE. template< typename U = T, typename std::enable_if{}>::type* = nullptr> - __host__ __device__ + inline __host__ __device__ HIP_vector_type& operator&=(const HIP_vector_type& x) noexcept { data &= x.data; @@ -256,7 +256,7 @@ THE SOFTWARE. template< typename U = T, typename std::enable_if{}>::type* = nullptr> - __host__ __device__ + inline __host__ __device__ HIP_vector_type& operator>>=(const HIP_vector_type& x) noexcept { data >>= x.data; @@ -265,7 +265,7 @@ THE SOFTWARE. template< typename U = T, typename std::enable_if{}>::type* = nullptr> - __host__ __device__ + inline __host__ __device__ HIP_vector_type& operator<<=(const HIP_vector_type& x) noexcept { data <<= x.data; @@ -275,24 +275,21 @@ THE SOFTWARE. template - __host__ __device__ - inline + inline __host__ __device__ HIP_vector_type operator+( const HIP_vector_type& x, const HIP_vector_type& y) noexcept { return HIP_vector_type{x} += y; } template - __host__ __device__ - inline + inline __host__ __device__ HIP_vector_type operator+( const HIP_vector_type& x, U y) noexcept { return HIP_vector_type{x} += y; } template - __host__ __device__ - inline + inline __host__ __device__ HIP_vector_type operator+( U x, const HIP_vector_type& y) noexcept { @@ -300,24 +297,21 @@ THE SOFTWARE. } template - __host__ __device__ - inline + inline __host__ __device__ HIP_vector_type operator-( const HIP_vector_type& x, const HIP_vector_type& y) noexcept { return HIP_vector_type{x} -= y; } template - __host__ __device__ - inline + inline __host__ __device__ HIP_vector_type operator-( const HIP_vector_type& x, U y) noexcept { return HIP_vector_type{x} -= y; } template - __host__ __device__ - inline + inline __host__ __device__ HIP_vector_type operator-( U x, const HIP_vector_type& y) noexcept { @@ -325,24 +319,21 @@ THE SOFTWARE. } template - __host__ __device__ - inline + inline __host__ __device__ HIP_vector_type operator*( const HIP_vector_type& x, const HIP_vector_type& y) noexcept { return HIP_vector_type{x} *= y; } template - __host__ __device__ - inline + inline __host__ __device__ HIP_vector_type operator*( const HIP_vector_type& x, U y) noexcept { return HIP_vector_type{x} *= y; } template - __host__ __device__ - inline + inline __host__ __device__ HIP_vector_type operator*( U x, const HIP_vector_type& y) noexcept { @@ -350,24 +341,21 @@ THE SOFTWARE. } template - __host__ __device__ - inline + inline __host__ __device__ HIP_vector_type operator/( const HIP_vector_type& x, const HIP_vector_type& y) noexcept { return HIP_vector_type{x} /= y; } template - __host__ __device__ - inline + inline __host__ __device__ HIP_vector_type operator/( const HIP_vector_type& x, U y) noexcept { return HIP_vector_type{x} /= y; } template - __host__ __device__ - inline + inline __host__ __device__ HIP_vector_type operator/( U x, const HIP_vector_type& y) noexcept { @@ -375,8 +363,7 @@ THE SOFTWARE. } template - __host__ __device__ - inline + inline __host__ __device__ bool operator==( const HIP_vector_type& x, const HIP_vector_type& y) noexcept { @@ -385,38 +372,33 @@ THE SOFTWARE. return true; } template - __host__ __device__ - inline + inline __host__ __device__ bool operator==(const HIP_vector_type& x, U y) noexcept { return x == HIP_vector_type{y}; } template - __host__ __device__ - inline + inline __host__ __device__ bool operator==(U x, const HIP_vector_type& y) noexcept { return HIP_vector_type{x} == y; } template - __host__ __device__ - inline + inline __host__ __device__ bool operator!=( const HIP_vector_type& x, const HIP_vector_type& y) noexcept { return !(x == y); } template - __host__ __device__ - inline + inline __host__ __device__ bool operator!=(const HIP_vector_type& x, U y) noexcept { return !(x == y); } template - __host__ __device__ - inline + inline __host__ __device__ bool operator!=(U x, const HIP_vector_type& y) noexcept { return !(x == y); @@ -426,7 +408,7 @@ THE SOFTWARE. typename T, unsigned int n, typename std::enable_if{}>* = nullptr> - inline + inline __host__ __device__ HIP_vector_type operator%( const HIP_vector_type& x, const HIP_vector_type& y) noexcept { @@ -437,7 +419,7 @@ THE SOFTWARE. unsigned int n, typename U, typename std::enable_if{}>* = nullptr> - inline + inline __host__ __device__ HIP_vector_type operator%( const HIP_vector_type& x, U y) noexcept { @@ -448,7 +430,7 @@ THE SOFTWARE. unsigned int n, typename U, typename std::enable_if{}>* = nullptr> - inline + inline __host__ __device__ HIP_vector_type operator%( U x, const HIP_vector_type& y) noexcept { @@ -459,7 +441,7 @@ THE SOFTWARE. typename T, unsigned int n, typename std::enable_if{}>* = nullptr> - inline + inline __host__ __device__ HIP_vector_type operator^( const HIP_vector_type& x, const HIP_vector_type& y) noexcept { @@ -470,7 +452,7 @@ THE SOFTWARE. unsigned int n, typename U, typename std::enable_if{}>* = nullptr> - inline + inline __host__ __device__ HIP_vector_type operator^( const HIP_vector_type& x, U y) noexcept { @@ -481,7 +463,7 @@ THE SOFTWARE. unsigned int n, typename U, typename std::enable_if{}>* = nullptr> - inline + inline __host__ __device__ HIP_vector_type operator^( U x, const HIP_vector_type& y) noexcept { @@ -492,7 +474,7 @@ THE SOFTWARE. typename T, unsigned int n, typename std::enable_if{}>* = nullptr> - inline + inline __host__ __device__ HIP_vector_type operator|( const HIP_vector_type& x, const HIP_vector_type& y) noexcept { @@ -503,7 +485,7 @@ THE SOFTWARE. unsigned int n, typename U, typename std::enable_if{}>* = nullptr> - inline + inline __host__ __device__ HIP_vector_type operator|( const HIP_vector_type& x, U y) noexcept { @@ -514,7 +496,7 @@ THE SOFTWARE. unsigned int n, typename U, typename std::enable_if{}>* = nullptr> - inline + inline __host__ __device__ HIP_vector_type operator|( U x, const HIP_vector_type& y) noexcept { @@ -525,7 +507,7 @@ THE SOFTWARE. typename T, unsigned int n, typename std::enable_if{}>* = nullptr> - inline + inline __host__ __device__ HIP_vector_type operator&( const HIP_vector_type& x, const HIP_vector_type& y) noexcept { @@ -536,7 +518,7 @@ THE SOFTWARE. unsigned int n, typename U, typename std::enable_if{}>* = nullptr> - inline + inline __host__ __device__ HIP_vector_type operator&( const HIP_vector_type& x, U y) noexcept { @@ -547,7 +529,7 @@ THE SOFTWARE. unsigned int n, typename U, typename std::enable_if{}>* = nullptr> - inline + inline __host__ __device__ HIP_vector_type operator&( U x, const HIP_vector_type& y) noexcept { @@ -558,7 +540,7 @@ THE SOFTWARE. typename T, unsigned int n, typename std::enable_if{}>* = nullptr> - inline + inline __host__ __device__ HIP_vector_type operator>>( const HIP_vector_type& x, const HIP_vector_type& y) noexcept { @@ -569,7 +551,7 @@ THE SOFTWARE. unsigned int n, typename U, typename std::enable_if{}>* = nullptr> - inline + inline __host__ __device__ HIP_vector_type operator>>( const HIP_vector_type& x, U y) noexcept { @@ -580,7 +562,7 @@ THE SOFTWARE. unsigned int n, typename U, typename std::enable_if{}>* = nullptr> - inline + inline __host__ __device__ HIP_vector_type operator>>( U x, const HIP_vector_type& y) noexcept { @@ -591,7 +573,7 @@ THE SOFTWARE. typename T, unsigned int n, typename std::enable_if{}>* = nullptr> - inline + inline __host__ __device__ HIP_vector_type operator<<( const HIP_vector_type& x, const HIP_vector_type& y) noexcept { @@ -602,7 +584,7 @@ THE SOFTWARE. unsigned int n, typename U, typename std::enable_if{}>* = nullptr> - inline + inline __host__ __device__ HIP_vector_type operator<<( const HIP_vector_type& x, U y) noexcept { @@ -613,7 +595,7 @@ THE SOFTWARE. unsigned int n, typename U, typename std::enable_if{}>* = nullptr> - inline + inline __host__ __device__ HIP_vector_type operator<<( U x, const HIP_vector_type& y) noexcept { @@ -685,27 +667,19 @@ __MAKE_VECTOR_TYPE__(float, float); __MAKE_VECTOR_TYPE__(double, double); #define DECLOP_MAKE_ONE_COMPONENT(comp, type) \ - __device__ __host__ \ - static \ - inline \ + static inline __device__ __host__ \ type make_##type(comp x) { type r = {x}; return r; } #define DECLOP_MAKE_TWO_COMPONENT(comp, type) \ - __device__ __host__ \ - static \ - inline \ + static inline __device__ __host__ \ type make_##type(comp x, comp y) { type r = {x, y}; return r; } #define DECLOP_MAKE_THREE_COMPONENT(comp, type) \ - __device__ __host__ \ - static \ - inline \ + static inline __device__ __host__ \ type make_##type(comp x, comp y, comp z) { type r = {x, y, z}; return r; } #define DECLOP_MAKE_FOUR_COMPONENT(comp, type) \ - __device__ __host__ \ - static \ - inline \ + static inline __device__ __host__ \ type make_##type(comp x, comp y, comp z, comp w) { \ type r = {x, y, z, w}; \ return r; \ From b7ae0b6442f7f15bc90c314d31ce4e0ac193d685 Mon Sep 17 00:00:00 2001 From: Michael LIAO Date: Thu, 2 May 2019 12:34:07 -0400 Subject: [PATCH 08/17] [test] Remove unused common routines. - Remove used common routines, which variadic functions, not supported feature, are used. [ROCm/hip commit: 16a27e06ee1c10b3937ea7fda91845a56c26d5cc] --- .../tests/src/deviceLib/vector_test_common.h | 42 ------------------- 1 file changed, 42 deletions(-) diff --git a/projects/hip/tests/src/deviceLib/vector_test_common.h b/projects/hip/tests/src/deviceLib/vector_test_common.h index fac5ab84a1..3fabcd7d04 100644 --- a/projects/hip/tests/src/deviceLib/vector_test_common.h +++ b/projects/hip/tests/src/deviceLib/vector_test_common.h @@ -25,45 +25,3 @@ THE SOFTWARE. template using Enable_if_t = typename std::enable_if::type; - -__host__ __device__ -std::false_type is_vec4(...); -__host__ __device__ -std::false_type is_vec3(...); -__host__ __device__ -std::false_type is_vec2(...); -__host__ __device__ -std::false_type is_vec1(...); - -template -__host__ __device__ -auto is_vec4(const T&) -> decltype(std::declval().xyzw, std::true_type{}); -template< - typename T, Enable_if_t())){}>* = nullptr> -__host__ __device__ -auto is_vec3(const T&) -> decltype(std::declval().xyz, std::true_type{}); -template< - typename T, - Enable_if_t< - !decltype(is_vec4(std::declval())){} && - !decltype(is_vec3(std::declval())){}>* = nullptr> -__host__ __device__ -auto is_vec2(const T&) -> decltype(std::declval().xy, std::true_type{}); -template< - typename T, - Enable_if_t< - !decltype(is_vec4(std::declval())){} && - !decltype(is_vec3(std::declval())){} && - !decltype(is_vec2(std::declval())){}>* = nullptr> -__host__ __device__ -auto is_vec1(const T&) -> decltype(std::declval().x, std::true_type{}); - -template -__host__ __device__ -constexpr -bool is_vec() { - return (dimension == 1) ? decltype(is_vec1(std::declval())){} : - ((dimension == 2) ? decltype(is_vec2(std::declval())){} : - ((dimension == 3) ? decltype(is_vec3(std::declval())){} : - decltype(is_vec4(std::declval())){})); -} \ No newline at end of file From 45eb04ea5c81153ec08e862534c5b14bd24c1c61 Mon Sep 17 00:00:00 2001 From: Konstantin Pyzhov Date: Fri, 3 May 2019 05:39:45 -0400 Subject: [PATCH 09/17] Refined hipSetDevice test. [ROCm/hip commit: 069fb7e486a9a0c3b2457f3c7aed15ee7db6f47c] --- .../src/Negative/Device/hipSetDevice.cpp | 22 ++++++++++++++----- 1 file changed, 16 insertions(+), 6 deletions(-) diff --git a/projects/hip/tests/src/Negative/Device/hipSetDevice.cpp b/projects/hip/tests/src/Negative/Device/hipSetDevice.cpp index 6e906ce8eb..ec0549e760 100644 --- a/projects/hip/tests/src/Negative/Device/hipSetDevice.cpp +++ b/projects/hip/tests/src/Negative/Device/hipSetDevice.cpp @@ -1,9 +1,19 @@ -#include "hip/hip_runtime.h" -#include "hip/hip_runtime_api.h" -#include "hipDeviceUtil.h" +#include "test_common.h" int main() { - HIP_CHECK(hipSetDevice(0), hipSetDevice); - HIP_CHECK(hipSetDevice(1026), hipSetDevice); - HIP_CHECK(hipSetDevice(-1), hipSetDevice); + int numDevices = 0; + + HIPCHECK_API(hipGetDeviceCount(&numDevices), hipSuccess); + if (numDevices > 0) { + for (int deviceId = 0; deviceId < numDevices; deviceId++) { + HIPCHECK_API(hipSetDevice(deviceId), hipSuccess); + } + HIPCHECK_API(hipSetDevice(numDevices), hipErrorInvalidDevice); + HIPCHECK_API(hipSetDevice(-1), hipErrorInvalidDevice); + } + else { + failed("Error: failed to find any compatible devices."); + } + + passed(); } From ae391b32111384a6e669058a0f2d178d83ea956e Mon Sep 17 00:00:00 2001 From: Konstantin Pyzhov Date: Fri, 3 May 2019 05:43:30 -0400 Subject: [PATCH 10/17] hipMultiThreadDevice test: Reduced maximum number of created HIP streams on Windows. [ROCm/hip commit: e29c9d50a6e3fef14c69912e33211ea3d8e2bb8f] --- .../src/runtimeApi/multiThread/hipMultiThreadDevice.cpp | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) diff --git a/projects/hip/tests/src/runtimeApi/multiThread/hipMultiThreadDevice.cpp b/projects/hip/tests/src/runtimeApi/multiThread/hipMultiThreadDevice.cpp index b95fb4b416..249fbeb792 100644 --- a/projects/hip/tests/src/runtimeApi/multiThread/hipMultiThreadDevice.cpp +++ b/projects/hip/tests/src/runtimeApi/multiThread/hipMultiThreadDevice.cpp @@ -9,6 +9,11 @@ #include "hip/hip_runtime_api.h" #include "test_common.h" +#ifdef _WIN32 +#define MAX_BURST_SIZE 40 +#else +#define MAX_BURST_SIZE 100 +#endif // Create a lot of streams and then destroy 'em. void createThenDestroyStreams(int iterations, int burstSize) { @@ -50,8 +55,8 @@ void waitStreams(int iterations) { // Some create many queue, some not many. // void multiThread_pyramid(bool serialize, int iters) { - printf("%s creating %d streams\n", __func__, iters * 100); - std::thread t1(createThenDestroyStreams, iters * 1, 100); + printf("%s creating %d streams\n", __func__, iters * MAX_BURST_SIZE); + std::thread t1(createThenDestroyStreams, iters * 1, MAX_BURST_SIZE); if (serialize) { t1.join(); printf("t1 done\n"); From cb55dc384a683371f6bbf69fb780c45cdfe72cdf Mon Sep 17 00:00:00 2001 From: Michael LIAO Date: Fri, 3 May 2019 14:56:00 -0400 Subject: [PATCH 11/17] [test] Use explicit cast for address space cast. [ROCm/hip commit: 2746d959064165a980d6c7f6a662f3797e302bf2] --- projects/hip/tests/src/deviceLib/hipAsynchronousStreams.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/tests/src/deviceLib/hipAsynchronousStreams.cpp b/projects/hip/tests/src/deviceLib/hipAsynchronousStreams.cpp index 17e6254154..b5f6d25793 100644 --- a/projects/hip/tests/src/deviceLib/hipAsynchronousStreams.cpp +++ b/projects/hip/tests/src/deviceLib/hipAsynchronousStreams.cpp @@ -117,7 +117,7 @@ __hip_as_write_block(__ockl_as_stream_t *stream, uchar service_id, uchar flags) { __ockl_as_stream_t ATTR_GLOBAL *gstream = - reinterpret_cast<__ockl_as_stream_t ATTR_GLOBAL *>(stream); + (__ockl_as_stream_t ATTR_GLOBAL *)(stream); return __ockl_as_write_block(gstream, service_id, connection_id, str, len, flags); } From d71327882c8282df67b6334dc738db8044779c0e Mon Sep 17 00:00:00 2001 From: Michael LIAO Date: Fri, 3 May 2019 15:55:40 -0400 Subject: [PATCH 12/17] [test] Add device variant of `std::declval`. - Current clang disallows any invocation of wrong-side functions even under context with type-inspection only. Work around that by adding a variant of `std::decl` with `__device__` attribute. [ROCm/hip commit: 4fbfb646f63050ed6cc7f748da1d34b526359802] --- .../tests/src/deviceLib/hipVectorTypesDevice.cpp | 13 ++++++++----- 1 file changed, 8 insertions(+), 5 deletions(-) diff --git a/projects/hip/tests/src/deviceLib/hipVectorTypesDevice.cpp b/projects/hip/tests/src/deviceLib/hipVectorTypesDevice.cpp index 74bf932b3a..a1db8de361 100644 --- a/projects/hip/tests/src/deviceLib/hipVectorTypesDevice.cpp +++ b/projects/hip/tests/src/deviceLib/hipVectorTypesDevice.cpp @@ -37,9 +37,12 @@ THE SOFTWARE. using namespace std; +template __device__ +typename std::add_rvalue_reference::type _declval() noexcept; + template< typename V, - Enable_if_t().x)>{}>* = nullptr> + Enable_if_t().x)>{}>* = nullptr> __device__ constexpr bool integer_unary_tests(const V&, const V&) { @@ -48,7 +51,7 @@ bool integer_unary_tests(const V&, const V&) { template< typename V, - Enable_if_t().x)>{}>* = nullptr> + Enable_if_t().x)>{}>* = nullptr> __device__ bool integer_unary_tests(V& f1, V& f2) { f1 %= f2; @@ -73,7 +76,7 @@ bool integer_unary_tests(V& f1, V& f2) { template< typename V, - Enable_if_t().x)>{}>* = nullptr> + Enable_if_t().x)>{}>* = nullptr> __device__ constexpr bool integer_binary_tests(const V&, const V&, const V&) { @@ -82,7 +85,7 @@ bool integer_binary_tests(const V&, const V&, const V&) { template< typename V, - Enable_if_t().x)>{}>* = nullptr> + Enable_if_t().x)>{}>* = nullptr> __device__ bool integer_binary_tests(V& f1, V& f2, V& f3) { f3 = f1 % f2; @@ -201,4 +204,4 @@ int main() { else { failed("Failed some vector test."); } -} \ No newline at end of file +} From f7eaff474b29b369618e014683d2d4532a67faa4 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 6 May 2019 15:54:31 +0530 Subject: [PATCH 13/17] Implement hipExtGetLinkTypeAndHopCount for ROCm devices Change-Id: Ie5bb4f640ac6d189c7fceeab22627a7494fd10bd [ROCm/hip commit: 035ef04e1f44db28af80b9b079ab0607774f0a84] --- .../include/hip/hcc_detail/hip_runtime_api.h | 14 ++++++ projects/hip/src/hip_device.cpp | 49 +++++++++++++++++++ 2 files changed, 63 insertions(+) diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index e8196af702..b6eb21e0ef 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -542,6 +542,20 @@ hipError_t hipSetDeviceFlags(unsigned flags); */ hipError_t hipChooseDevice(int* device, const hipDeviceProp_t* prop); +/** + * @brief Returns the link type and hop count between two devices + * + * @param [in] device1 Ordinal for device1 + * @param [in] device2 Ordinal for device2 + * @param [out] linktype Returns the link type (See hsa_amd_link_info_type_t) between the two devices + * @param [out] hopcount Returns the hop count between the two devices + * + * Queries and returns the HSA link type and the hop count between the two specified devices. + * + * @returns #hipSuccess, #hipInvalidDevice, #hipErrorRuntimeOther + */ +hipError_t hipExtGetLinkTypeAndHopCount(int device1, int device2, uint32_t* linktype, uint32_t* hopcount); + // end doxygen Device /** * @} diff --git a/projects/hip/src/hip_device.cpp b/projects/hip/src/hip_device.cpp index 246ce8cf6f..6570b73316 100644 --- a/projects/hip/src/hip_device.cpp +++ b/projects/hip/src/hip_device.cpp @@ -531,3 +531,52 @@ hipError_t hipChooseDevice(int* device, const hipDeviceProp_t* prop) { } return ihipLogStatus(e); } + +#define HSA_ERROR_CHECK(hsa_error_code) \ + if ((hsa_error_code != HSA_STATUS_SUCCESS) && (hsa_error_code != HSA_STATUS_INFO_BREAK)) { \ + return ihipLogStatus(hipErrorRuntimeOther); \ + } + +hipError_t hipExtGetLinkTypeAndHopCount(int device1, int device2, uint32_t* linktype, uint32_t* hopcount) { + HIP_INIT_API(hipExtGetLinkTypeAndHopCount, device1, device2, linktype, hopcount); + + if ((device1 < 0) || (device1 >= g_deviceCnt) || (device2 < 0) || (device2 >= g_deviceCnt)) { + return ihipLogStatus(hipErrorInvalidDevice); + } else { + auto device1Handle = ihipGetDevice(device1); + auto device2Handle = ihipGetDevice(device2); + + const auto& find_pool = [](hsa_amd_memory_pool_t pool, void* data) { + bool allowed; + hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALLOWED, &allowed); + if (allowed) { + hsa_amd_segment_t segment; + hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment); + if (HSA_AMD_SEGMENT_GLOBAL != segment) return HSA_STATUS_SUCCESS; + + uint32_t flags; + hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flags); + if (flags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED) { + *((hsa_amd_memory_pool_t*)data) = pool; + return HSA_STATUS_INFO_BREAK; + } + } + return HSA_STATUS_SUCCESS; + }; + + hsa_status_t err; + hsa_amd_memory_pool_t pool; + err = hsa_amd_agent_iterate_memory_pools(device2Handle->_hsaAgent, find_pool, (void*)&pool); + HSA_ERROR_CHECK(err); + + hsa_amd_memory_pool_link_info_t link_info; + err = hsa_amd_agent_memory_pool_get_info(device1Handle->_hsaAgent, pool, HSA_AMD_AGENT_MEMORY_POOL_INFO_LINK_INFO, &link_info); + HSA_ERROR_CHECK(err); + *linktype = link_info.link_type; + + err = hsa_amd_agent_memory_pool_get_info(device1Handle->_hsaAgent, pool, HSA_AMD_AGENT_MEMORY_POOL_INFO_NUM_LINK_HOPS, hopcount); + HSA_ERROR_CHECK(err); + + return ihipLogStatus(hipSuccess); + } +} From d700f67b79018f9b3aa33cf3a6362c4e1956fb36 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Tue, 7 May 2019 03:59:43 +0530 Subject: [PATCH 14/17] Fix symbol address issue on NVCC path [ROCm/hip commit: e906bd69cec791a9fe97d7a1be5e4b74bd6fa122] --- projects/hip/include/hip/nvcc_detail/hip_runtime.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime.h b/projects/hip/include/hip/nvcc_detail/hip_runtime.h index 19d740a1ee..3a18f11614 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime.h +++ b/projects/hip/include/hip/nvcc_detail/hip_runtime.h @@ -96,7 +96,7 @@ typedef int hipLaunchParm; #define hipGridDim_y gridDim.y #define hipGridDim_z gridDim.z -#define HIP_SYMBOL(X) X +#define HIP_SYMBOL(X) &X /** * extern __shared__ From 236834c679f567800525e903deb4b405c732f141 Mon Sep 17 00:00:00 2001 From: wkwchau Date: Tue, 7 May 2019 06:28:25 -0400 Subject: [PATCH 15/17] Return hipErrorInsufficientDriver status when CPU device not found (#1064) * Return hipErrorInsufficientDriver status when CPU device not found - no exception thrown * Return hipErrorInsufficientDriver status when CPU device not found [ROCm/hip commit: 29b3b46b425a039b9a4de6ff81a12605b8f7f648] --- projects/hip/src/hip_hcc.cpp | 15 ++++++++++++--- 1 file changed, 12 insertions(+), 3 deletions(-) diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index dd08f2bb2e..8011160cc1 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -130,6 +130,7 @@ std::vector g_hip_visible_devices; hsa_agent_t g_cpu_agent; hsa_agent_t* g_allAgents; // CPU agents + all the visible GPU agents. unsigned g_numLogicalThreads; +bool g_initDeviceFound = false; std::atomic g_lastShortTid(1); @@ -1405,7 +1406,8 @@ void ihipInit() { hsa_status_t err = hsa_iterate_agents(findCpuAgent, &g_cpu_agent); if (err != HSA_STATUS_INFO_BREAK) { // didn't find a CPU. - throw ihipException(hipErrorRuntimeOther); + g_initDeviceFound = false; + return; } g_deviceArray = new ihipDevice_t*[deviceCnt]; @@ -1440,14 +1442,21 @@ void ihipInit() { tprintf(DB_SYNC, "pid=%u %-30s g_numLogicalThreads=%u\n", getpid(), "", g_numLogicalThreads); + + g_initDeviceFound = true; } namespace hip_impl { hipError_t hip_init() { static std::once_flag hip_initialized; std::call_once(hip_initialized, ihipInit); - ihipCtxStackUpdate(); - return hipSuccess; + if (g_initDeviceFound) { + ihipCtxStackUpdate(); + return hipSuccess; + } + else { + return hipErrorInsufficientDriver; + } } } From 6f69a5d16e3a3d58e31f41d40178958a1b94cde3 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Tue, 7 May 2019 06:29:40 -0400 Subject: [PATCH 16/17] Add documentation for supported clang options (#1065) * Add documentation for supported clang options * Fix typo [ROCm/hip commit: 669d17707903f31f8675b13130c98d3b8eb77bf7] --- projects/hip/docs/markdown/clang_options.md | 606 ++++++++++++++++++ projects/hip/docs/markdown/clang_options.txt | 606 ++++++++++++++++++ .../hip/docs/markdown/gen_clang_option_doc.sh | 75 +++ .../docs/markdown/hip_programming_guide.md | 2 + 4 files changed, 1289 insertions(+) create mode 100644 projects/hip/docs/markdown/clang_options.md create mode 100644 projects/hip/docs/markdown/clang_options.txt create mode 100755 projects/hip/docs/markdown/gen_clang_option_doc.sh diff --git a/projects/hip/docs/markdown/clang_options.md b/projects/hip/docs/markdown/clang_options.md new file mode 100644 index 0000000000..7c2538426c --- /dev/null +++ b/projects/hip/docs/markdown/clang_options.md @@ -0,0 +1,606 @@ +# Support of Clang options + Clang versiion: clang version 9.0.0 (https://github.com/llvm-mirror/clang.git 50988b3ea26e892a74a2848e5a06b96f2472a05c) (https://github.com/RadeonOpenCompute/llvm.git 6483275d3d7603784c30e08927f2b315dfb95639) + +|Option|Support|Description| +|-------|------|-------| +|`-###`|Supported|`Print (but do not run) the commands to run for this compilation`| +|`--analyzer-output `|Supported|`Static analyzer report output format (html\|plist\|plist-multi-file\|plist-html\|text).`| +|`--analyze`|Supported|`Run the static analyzer`| +|`-arcmt-migrate-emit-errors`|Unsupported|`Emit ARC errors even if the migrator can fix them`| +|`-arcmt-migrate-report-output `|Unsupported|`Output path for the plist report`| +|`-B `|Supported|`Add to search path for binaries and object files used implicitly`| +|`-CC`|Supported|`Include comments from within macros in preprocessed output`| +|`-cfguard`|Supported|`Emit tables required for Windows Control Flow Guard.`| +|`-cl-denorms-are-zero`|Supported|`OpenCL only. Allow denormals to be flushed to zero.`| +|`-cl-fast-relaxed-math`|Supported|`OpenCL only. Sets -cl-finite-math-only and -cl-unsafe-math-optimizations, and defines __FAST_RELAXED_MATH__.`| +|`-cl-finite-math-only`|Supported|`OpenCL only. Allow floating-point optimizations that assume arguments and results are not NaNs or +-Inf.`| +|`-cl-fp32-correctly-rounded-divide-sqrt`|Supported|`OpenCL only. Specify that single precision floating-point divide and sqrt used in the program source are correctly rounded.`| +|`-cl-kernel-arg-info`|Supported|`OpenCL only. Generate kernel argument metadata.`| +|`-cl-mad-enable`|Supported|`OpenCL only. Allow use of less precise MAD computations in the generated binary.`| +|`-cl-no-signed-zeros`|Supported|`OpenCL only. Allow use of less precise no signed zeros computations in the generated binary.`| +|`-cl-opt-disable`|Supported|`OpenCL only. This option disables all optimizations. By default optimizations are enabled.`| +|`-cl-single-precision-constant`|Supported|`OpenCL only. Treat double precision floating-point constant as single precision constant.`| +|`-cl-std=`|Supported|`OpenCL language standard to compile for.`| +|`-cl-strict-aliasing`|Supported|`OpenCL only. This option is added for compatibility with OpenCL 1.0.`| +|`-cl-uniform-work-group-size`|Supported|`OpenCL only. Defines that the global work-size be a multiple of the work-group size specified to clEnqueueNDRangeKernel`| +|`-cl-unsafe-math-optimizations`|Supported|`OpenCL only. Allow unsafe floating-point optimizations. Also implies -cl-no-signed-zeros and -cl-mad-enable.`| +|`--config `|Supported|`Specifies configuration file`| +|`--cuda-compile-host-device`|Supported|`Compile CUDA code for both host and device (default). Has no effect on non-CUDA compilations.`| +|`--cuda-device-only`|Supported|`Compile CUDA code for device only`| +|`--cuda-gpu-arch=`|Supported|`CUDA GPU architecture (e.g. sm_35). May be specified more than once.`| +|`--cuda-host-only`|Supported|`Compile CUDA code for host only. Has no effect on non-CUDA compilations.`| +|`--cuda-include-ptx=`|Unsupported|`Include PTX for the following GPU architecture (e.g. sm_35) or 'all'. May be specified more than once.`| +|`--cuda-noopt-device-debug`|Unsupported|`Enable device-side debug info generation. Disables ptxas optimizations.`| +|`--cuda-path-ignore-env`|Unsupported|`Ignore environment variables to detect CUDA installation`| +|`--cuda-path=`|Unsupported|`CUDA installation path`| +|`-cxx-isystem `|Supported|`Add directory to the C++ SYSTEM include search path`| +|`-C`|Supported|`Include comments in preprocessed output`| +|`-c`|Supported|`Only run preprocess, compile, and assemble steps`| +|`-dD`|Supported|`Print macro definitions in -E mode in addition to normal output`| +|`-dependency-dot `|Supported|`Filename to write DOT-formatted header dependencies to`| +|`-dependency-file `|Supported|`Filename (or -) to write dependency output to`| +|`-dI`|Supported|`Print include directives in -E mode in addition to normal output`| +|`-dM`|Supported|`Print macro definitions in -E mode instead of normal output`| +|`-D `|Supported|`= Define to (or 1 if omitted)`| +|`-emit-ast`|Supported|`Emit Clang AST files for source inputs`| +|`-emit-llvm`|Supported|`Use the LLVM representation for assembler and object files`| +|`-enable-trivial-auto-var-init-zero-knowing-it-will-be-removed-from-clang`|Supported|`Trivial automatic variable initialization to zero is only here for benchmarks, it'll eventually be removed, and I'm OK with that because I'm only using it to benchmark`| +|`-E`||`Only run the preprocessor`| +|`-faddrsig`|Supported|`Emit an address-significance table`| +|`-faligned-allocation`|Supported|`Enable C++17 aligned allocation functions`| +|`-fallow-editor-placeholders`|Supported|`Treat editor placeholders as valid source code`| +|`-fansi-escape-codes`|Supported|`Use ANSI escape codes for diagnostics`| +|`-fapple-kext`|Unsupported|`Use Apple's kernel extensions ABI`| +|`-fapple-pragma-pack`|Unsupported|`Enable Apple gcc-compatible #pragma pack handling`| +|`-fapplication-extension`|Unsupported|`Restrict code to those available for App Extensions`| +|`-fblocks`|Supported|`Enable the 'blocks' language feature`| +|`-fborland-extensions`|Unsupported|`Accept non-standard constructs supported by the Borland compiler`| +|`-fbuild-session-file=`|Supported|`Use the last modification time of as the build session timestamp`| +|`-fbuild-session-timestamp=