From 357b6844fa9958a98bf9b5eba16fd4efa00765a2 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Fri, 21 Sep 2018 15:28:10 -0400 Subject: [PATCH 1/8] Improve hip_trig test case Hip_trig is failing but returning a passing value on HIP-Clang. Also update with debug output and free arrays used. --- tests/src/deviceLib/hip_trig.cpp | 43 +++++++++++++++++++++++++------- 1 file changed, 34 insertions(+), 9 deletions(-) diff --git a/tests/src/deviceLib/hip_trig.cpp b/tests/src/deviceLib/hip_trig.cpp index 3f200fcf8f..b7542c25a4 100644 --- a/tests/src/deviceLib/hip_trig.cpp +++ b/tests/src/deviceLib/hip_trig.cpp @@ -35,6 +35,8 @@ THE SOFTWARE. #define LEN 512 #define SIZE LEN << 2 +#define TEST_DEBUG (0) + __global__ void kernel_trig(hipLaunchParm lp, float* In, float* sin_d, float* cos_d, float* tan_d, float* sin_pd, float* cos_pd) { int tid = threadIdx.x + blockIdx.x * blockDim.x; @@ -47,6 +49,7 @@ __global__ void kernel_trig(hipLaunchParm lp, float* In, float* sin_d, float* co int main() { float *In, *sin_h, *cos_h, *tan_h, *sin_ph, *cos_ph; float *In_d, *sin_d, *cos_d, *tan_d, *sin_pd, *cos_pd; + int errors = 0; In = new float[LEN]; sin_h = new float[LEN]; cos_h = new float[LEN]; @@ -61,14 +64,16 @@ int main() { sin_ph[i] = 0.0f; cos_ph[i] = 0.0f; } - hipMalloc((void**)&In_d, SIZE); - hipMalloc((void**)&sin_d, SIZE); - hipMalloc((void**)&cos_d, SIZE); - hipMalloc((void**)&tan_d, SIZE); - hipMalloc((void**)&sin_pd, SIZE); - hipMalloc((void**)&cos_pd, SIZE); + HIP_ASSERT(hipMalloc((void**)&In_d, SIZE)); + HIP_ASSERT(hipMalloc((void**)&sin_d, SIZE)); + HIP_ASSERT(hipMalloc((void**)&cos_d, SIZE)); + HIP_ASSERT(hipMalloc((void**)&tan_d, SIZE)); + HIP_ASSERT(hipMalloc((void**)&sin_pd, SIZE)); + HIP_ASSERT(hipMalloc((void**)&cos_pd, SIZE)); + hipMemcpy(In_d, In, SIZE, hipMemcpyHostToDevice); - hipLaunchKernel(kernel_trig, dim3(LEN, 1, 1), dim3(1, 1, 1), 0, 0, In_d, sin_d, cos_d, tan_d, + hipLaunchKernel(kernel_trig, dim3(LEN, 1, 1), dim3(1, 1, 1), 0, 0, + In_d, sin_d, cos_d, tan_d, sin_pd, cos_pd); hipMemcpy(sin_h, sin_d, SIZE, hipMemcpyDeviceToHost); hipMemcpy(cos_h, cos_d, SIZE, hipMemcpyDeviceToHost); @@ -77,8 +82,28 @@ int main() { hipMemcpy(cos_ph, cos_pd, SIZE, hipMemcpyDeviceToHost); for (int i = 0; i < LEN; i++) { if (sin_h[i] != sin_ph[i] || cos_h[i] != cos_ph[i] || tan_h[i] * cos_h[i] != sin_h[i]) { - std::cout << "Failed!" << std::endl; + errors++; +#if TEST_DEBUG + std::cout << "Check Failed!" << std::endl; + std::cout << " sin_h: " << sin_h[i] << " sin_ph: " << sin_ph[i] << "\n" + << " cos_h: " << cos_h[i] << " cos_ph:" << cos_ph[i] << "\n" + << " tan_h * cos_h: " << tan_h[i] * cos_h[i] << " sin_h[i]: " << sin_h[i] << "\n"; +#endif } } - passed(); + + HIP_ASSERT(hipFree(In_d)); + HIP_ASSERT(hipFree(sin_d)); + HIP_ASSERT(hipFree(cos_d)); + HIP_ASSERT(hipFree(tan_d)); + HIP_ASSERT(hipFree(sin_pd)); + HIP_ASSERT(hipFree(cos_pd)); + + if (errors != 0) { + std::cout << "hip_trig FAILED!" << std::endl; + return -1; + } else { + std::cout << "hip_trig PASSED!" << std::endl; + } + return errors; } From 3d3820272a45a9be58d0c77ee4de18c4c620eb24 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Mon, 24 Sep 2018 15:01:24 +0000 Subject: [PATCH 2/8] Fix missing HIP_ASSERT in hip_trig --- tests/src/deviceLib/hip_trig.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/tests/src/deviceLib/hip_trig.cpp b/tests/src/deviceLib/hip_trig.cpp index b7542c25a4..29c24cf5b4 100644 --- a/tests/src/deviceLib/hip_trig.cpp +++ b/tests/src/deviceLib/hip_trig.cpp @@ -32,6 +32,8 @@ THE SOFTWARE. #include "test_common.h" #include +#define HIP_ASSERT(x) (assert((x) == hipSuccess)) + #define LEN 512 #define SIZE LEN << 2 From 4e06ab07e272fd3e4d5a6c6983e7fed5ee27f46c Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 25 Sep 2018 16:50:40 +0300 Subject: [PATCH 3/8] [HIPIFY][docs] Update CUDNN_API_supported_by_HIP.md --- docs/markdown/CUDNN_API_supported_by_HIP.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/docs/markdown/CUDNN_API_supported_by_HIP.md b/docs/markdown/CUDNN_API_supported_by_HIP.md index ec640fecdd..ffdeb881e5 100644 --- a/docs/markdown/CUDNN_API_supported_by_HIP.md +++ b/docs/markdown/CUDNN_API_supported_by_HIP.md @@ -239,8 +239,8 @@ |`cudnnSetTensor` |`hipdnnSetTensor` | |`cudnnScaleTensor` |`hipdnnScaleTensor` | |`cudnnCreateFilterDescriptor` |`hipdnnCreateFilterDescriptor` | -|`cudnnSetFilter4dDescriptor` | | -|`cudnnGetFilter4dDescriptor` | | +|`cudnnSetFilter4dDescriptor` |`hipdnnSetFilter4dDescriptor` | +|`cudnnGetFilter4dDescriptor` |`hipdnnGetFilter4dDescriptor` | |`cudnnSetFilterNdDescriptor` |`hipdnnSetFilterNdDescriptor` | |`cudnnGetFilterNdDescriptor` |`hipdnnGetFilterNdDescriptor` | |`cudnnDestroyFilterDescriptor` |`hipdnnDestroyFilterDescriptor` | From d279c7a1dde54b407903240aaf95f918e5769d00 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 25 Sep 2018 18:46:23 +0300 Subject: [PATCH 4/8] [HIPIFY][BLAS] Add support of hipblasGemmEx and corresponding types TODO (hipBLAS/HIP): rename hipblasDatatype_t to hipDataType_t and move it from hipBLAS to HIP, as Data types are used not only in BLAS library. --- docs/markdown/CUBLAS_API_supported_by_HIP.md | 8 +- ..._Runtime_API_functions_supported_by_HIP.md | 96 +++++++++++-------- hipify-clang/src/CUDA2HipMap.cpp | 47 +++++---- 3 files changed, 82 insertions(+), 69 deletions(-) diff --git a/docs/markdown/CUBLAS_API_supported_by_HIP.md b/docs/markdown/CUBLAS_API_supported_by_HIP.md index 932092612b..750703f4e1 100644 --- a/docs/markdown/CUBLAS_API_supported_by_HIP.md +++ b/docs/markdown/CUBLAS_API_supported_by_HIP.md @@ -35,9 +35,9 @@ | enum |***`cublasAtomicsMode_t`*** | | | 0 |*`CUBLAS_ATOMICS_NOT_ALLOWED`* | | | 1 |*`CUBLAS_ATOMICS_ALLOWED`* | | -| enum |***`cublasAtomicsMode_t`*** | | -| -1 |*`CUBLAS_GEMM_DFALT`* | | -| -1 |*`CUBLAS_GEMM_DEFAULT`* | | +| enum |***`cublasGemmAlgo_t`*** |***`hipblasGemmAlgo_t`*** | +| -1 |*`CUBLAS_GEMM_DFALT`* |*`HIPBLAS_GEMM_DEFAULT`* | 160 | +| -1 |*`CUBLAS_GEMM_DEFAULT`* |*`HIPBLAS_GEMM_DEFAULT`* | 160 | | 0 |*`CUBLAS_GEMM_ALGO0`* | | | 1 |*`CUBLAS_GEMM_ALGO1`* | | | 2 |*`CUBLAS_GEMM_ALGO2`* | | @@ -391,7 +391,7 @@ |`cublasZgemm3m` | | |`cublasHgemm` |`hipblasHgemm` | |`cublasSgemmEx` | | -|`cublasGemmEx` | | +|`cublasGemmEx` |`hipblasGemmEx` | |`cublasCgemmEx` | | |`cublasUint8gemmBias` | | |`cublasSsyrk` | | diff --git a/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md b/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md index dca2683b12..7f87981000 100644 --- a/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md +++ b/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md @@ -93,10 +93,10 @@ ## **7. Occupancy** -| **CUDA** | **HIP** | -|-----------------------------------------------------------|-------------------------------| +| **CUDA** | **HIP** | +|-----------------------------------------------------------|-----------------------------------------------| | `cudaOccupancyMaxActiveBlocksPerMultiprocessor` | `hipOccupancyMaxActiveBlocksPerMultiprocessor`| -| `cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags` | | +| `cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags` | | ## **8. Execution Control [deprecated since 7.0]** @@ -326,13 +326,13 @@ ## **25. Texture Object Management** -| **CUDA** | **HIP** | -|-----------------------------------------------------------|-------------------------------| -| `cudaCreateTextureObject` |`hipCreateTextureObject` | -| `cudaDestroyTextureObject` |`hipDestroyTextureObject` | -| `cudaGetTextureObjectResourceDesc` |`hipGetTextureObjectResourceDesc` | +| **CUDA** | **HIP** | +|-----------------------------------------------------------|--------------------------------------| +| `cudaCreateTextureObject` |`hipCreateTextureObject` | +| `cudaDestroyTextureObject` |`hipDestroyTextureObject` | +| `cudaGetTextureObjectResourceDesc` |`hipGetTextureObjectResourceDesc` | | `cudaGetTextureObjectResourceViewDesc` |`hipGetTextureObjectResourceViewDesc` | -| `cudaGetTextureObjectTextureDesc` |`hipGetTextureObjectTextureDesc` | +| `cudaGetTextureObjectTextureDesc` |`hipGetTextureObjectTextureDesc` | ## **26. Surface Object Management** @@ -352,36 +352,36 @@ ## **28. C++ API Routines** *(7.0 contains, 7.5 doesn’t)* -| **CUDA** | **HIP** | -|-----------------------------------------------------------|-------------------------------| -| `cudaBindSurfaceToArray` | | -| `cudaBindTexture` | `hipBindTexture` | -| `cudaBindTexture2D` | | -| `cudaBindTextureToArray` | | -| `cudaBindTextureToMipmappedArray` | | -| `cudaCreateChannelDesc` | `hipCreateChannelDesc` | -| `cudaFuncGetAttributes` | | -| `cudaFuncSetCacheConfig` | | -| `cudaGetSymbolAddress` | | -| `cudaGetSymbolSize` | | -| `cudaGetTextureAlignmentOffset` | | -| `cudaLaunch` | | -| `cudaLaunchKernel` | | -| `cudaMallocHost` | | -| `cudaMallocManaged` | | -| `cudaMemcpyFromSymbol` | | -| `cudaMemcpyFromSymbolAsync` | | -| `cudaMemcpyToSymbol` | | -| `cudaMemcpyToSymbolAsync` | | +| **CUDA** | **HIP** | +|-----------------------------------------------------------|------------------------------------------------| +| `cudaBindSurfaceToArray` | | +| `cudaBindTexture` | `hipBindTexture` | +| `cudaBindTexture2D` | | +| `cudaBindTextureToArray` | | +| `cudaBindTextureToMipmappedArray` | | +| `cudaCreateChannelDesc` | `hipCreateChannelDesc` | +| `cudaFuncGetAttributes` | | +| `cudaFuncSetCacheConfig` | | +| `cudaGetSymbolAddress` | | +| `cudaGetSymbolSize` | | +| `cudaGetTextureAlignmentOffset` | | +| `cudaLaunch` | | +| `cudaLaunchKernel` | | +| `cudaMallocHost` | | +| `cudaMallocManaged` | | +| `cudaMemcpyFromSymbol` | | +| `cudaMemcpyFromSymbolAsync` | | +| `cudaMemcpyToSymbol` | | +| `cudaMemcpyToSymbolAsync` | | | `cudaOccupancyMaxActiveBlocksPerMultiprocessor` | `hipOccupancyMaxActiveBlocksPerMultiprocessor` | -| `cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags` | | -| `cudaOccupancyMaxPotentialBlockSize` | `hipOccupancyMaxPotentialBlockSize` | -| `cudaOccupancyMaxPotentialBlockSizeVariableSMem` | | -| `cudaOccupancyMaxPotentialBlockSizeVariableSMemWithFlags` | | -| `cudaOccupancyMaxPotentialBlockSizeWithFlags` | | -| `cudaSetupArgument` | | -| `cudaStreamAttachMemAsync` | | -| `cudaUnbindTexture` | `hipUnbindTexture` | +| `cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags` | | +| `cudaOccupancyMaxPotentialBlockSize` | `hipOccupancyMaxPotentialBlockSize` | +| `cudaOccupancyMaxPotentialBlockSizeVariableSMem` | | +| `cudaOccupancyMaxPotentialBlockSizeVariableSMemWithFlags` | | +| `cudaOccupancyMaxPotentialBlockSizeWithFlags` | | +| `cudaSetupArgument` | | +| `cudaStreamAttachMemAsync` | | +| `cudaUnbindTexture` | `hipUnbindTexture` | ## **30. Profiler Control** @@ -395,8 +395,8 @@ ## **31. Data types** -| **type** | **CUDA** | **HIP** | -|-------------:|-----------------------------------------------|------------------------------------------------------| +| **type** | **CUDA** | **HIP** |**HIP value** (if differs) | +|-------------:|-----------------------------------------------|------------------------------------------------------|---------------------------| | struct | `cudaChannelFormatDesc` | `hipChannelFormatDesc` | | struct | `cudaDeviceProp` | `hipDeviceProp_t` | | struct | `cudaExtent` | `hipExtent` | @@ -790,3 +790,19 @@ | define | `cudaTextureType1DLayered` | `hipTextureType1DLayered` | | define | `cudaTextureType2DLayered` | `hipTextureType2DLayered` | | define | `cudaTextureTypeCubemapLayered` | `hipTextureTypeCubemapLayered` | +| enum |***`cudaDataType_t`*** |***`hipblasDatatype_t`*** | +| enum |***`cudaDataType`*** |***`hipblasDatatype_t`*** | +| 2 |*`CUDA_R_16F`* |*`HIPBLAS_R_16F`* | 150 | +| 6 |*`CUDA_C_16F`* |*`HIPBLAS_C_16F`* | 153 | +| 0 |*`CUDA_R_32F`* |*`HIPBLAS_R_32F`* | 151 | +| 4 |*`CUDA_C_32F`* |*`HIPBLAS_C_32F`* | 154 | +| 1 |*`CUDA_R_64F`* |*`HIPBLAS_R_64F`* | 152 | +| 5 |*`CUDA_C_64F`* |*`HIPBLAS_C_64F`* | 155 | +| 3 |*`CUDA_R_8I`* | | +| 7 |*`CUDA_C_8I`* | | +| 8 |*`CUDA_R_8U`* | | +| 9 |*`CUDA_C_8U`* | | +| 10 |*`CUDA_R_32I`* | | +| 11 |*`CUDA_C_32I`* | | +| 12 |*`CUDA_R_32U`* | | +| 13 |*`CUDA_C_32U`* | | diff --git a/hipify-clang/src/CUDA2HipMap.cpp b/hipify-clang/src/CUDA2HipMap.cpp index ca53738abb..d8859dfaa3 100644 --- a/hipify-clang/src/CUDA2HipMap.cpp +++ b/hipify-clang/src/CUDA2HipMap.cpp @@ -1389,22 +1389,23 @@ const std::map CUDA_IDENTIFIER_MAP{ /////////////////////////////// CUDA RT API /////////////////////////////// // Data types - {"cudaDataType_t", {"hipDataType_t", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}}, - {"cudaDataType", {"hipDataType_t", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}}, - {"CUDA_R_16F", {"hipR16F", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}}, - {"CUDA_C_16F", {"hipC16F", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}}, - {"CUDA_R_32F", {"hipR32F", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}}, - {"CUDA_C_32F", {"hipC32F", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}}, - {"CUDA_R_64F", {"hipR64F", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}}, - {"CUDA_C_64F", {"hipC64F", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}}, - {"CUDA_R_8I", {"hipR8I", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}}, - {"CUDA_C_8I", {"hipC8I", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}}, - {"CUDA_R_8U", {"hipR8U", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}}, - {"CUDA_C_8U", {"hipC8U", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}}, - {"CUDA_R_32I", {"hipR32I", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}}, - {"CUDA_C_32I", {"hipC32I", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}}, - {"CUDA_R_32U", {"hipR32U", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}}, - {"CUDA_C_32U", {"hipC32U", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}}, + // TODO: rename hipblasDatatype_t to hipDataType_t and move from hipBLAS to HIP + {"cudaDataType_t", {"hipblasDatatype_t", CONV_TYPE, API_RUNTIME}}, + {"cudaDataType", {"hipblasDatatype_t", CONV_TYPE, API_RUNTIME}}, + {"CUDA_R_16F", {"HIPBLAS_R_16F", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 2 // 150 + {"CUDA_C_16F", {"HIPBLAS_C_16F", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 6 // 153 + {"CUDA_R_32F", {"HIPBLAS_R_32F", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 0 // 151 + {"CUDA_C_32F", {"HIPBLAS_C_32F", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 4 // 154 + {"CUDA_R_64F", {"HIPBLAS_R_64F", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 1 // 152 + {"CUDA_C_64F", {"HIPBLAS_C_64F", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 5 // 155 + {"CUDA_R_8I", {"HIPBLAS_R_8I", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 3 // + {"CUDA_C_8I", {"HIPBLAS_C_8I", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 7 // + {"CUDA_R_8U", {"HIPBLAS_R_8U", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 8 // + {"CUDA_C_8U", {"HIPBLAS_C_8U", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 9 // + {"CUDA_R_32I", {"HIPBLAS_R_32I", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 10 // + {"CUDA_C_32I", {"HIPBLAS_C_32I", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 11 // + {"CUDA_R_32U", {"HIPBLAS_R_32U", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 12 // + {"CUDA_C_32U", {"HIPBLAS_C_32U", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 13 // // Library property types // IMPORTANT: no cuda prefix @@ -2115,11 +2116,7 @@ const std::map CUDA_IDENTIFIER_MAP{ {"CUBLAS_ATOMICS_ALLOWED", {"HIPBLAS_ATOMICS_ALLOWED", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // Blas Data Type - {"cublasDataType_t", {"hipblasDataType_t", CONV_TYPE, API_BLAS, HIP_UNSUPPORTED}}, - {"CUBLAS_DATA_FLOAT", {"HIPBLAS_DATA_FLOAT", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, - {"CUBLAS_DATA_DOUBLE", {"HIPBLAS_DATA_DOUBLE", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, - {"CUBLAS_DATA_HALF", {"HIPBLAS_DATA_HALF", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, - {"CUBLAS_DATA_INT8", {"HIPBLAS_DATA_INT8", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, + {"cublasDataType_t", {"hipblasDatatype_t", CONV_TYPE, API_BLAS}}, // Blas Math mode/tensor operation {"cublasMath_t", {"hipblasMath_t", CONV_TYPE, API_BLAS, HIP_UNSUPPORTED}}, @@ -2127,9 +2124,9 @@ const std::map CUDA_IDENTIFIER_MAP{ {"CUBLAS_TENSOR_OP_MATH", {"HIPBLAS_TENSOR_OP_MATH", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // Blass different GEMM algorithms - {"cublasGemmAlgo_t", {"hipblasGemmAlgo_t", CONV_TYPE, API_BLAS, HIP_UNSUPPORTED}}, - {"CUBLAS_GEMM_DFALT", {"HIPBLAS_GEMM_DFALT", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // -1 - {"CUBLAS_GEMM_DEFAULT", {"HIPBLAS_GEMM_DEFAULT", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // -1 + {"cublasGemmAlgo_t", {"hipblasGemmAlgo_t", CONV_TYPE, API_BLAS}}, + {"CUBLAS_GEMM_DFALT", {"HIPBLAS_GEMM_DEFAULT", CONV_NUMERIC_LITERAL, API_BLAS}}, // -1 // 160 + {"CUBLAS_GEMM_DEFAULT", {"HIPBLAS_GEMM_DEFAULT", CONV_NUMERIC_LITERAL, API_BLAS}}, // -1 // 160 {"CUBLAS_GEMM_ALGO0", {"HIPBLAS_GEMM_ALGO0", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 0 {"CUBLAS_GEMM_ALGO1", {"HIPBLAS_GEMM_ALGO1", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 1 {"CUBLAS_GEMM_ALGO2", {"HIPBLAS_GEMM_ALGO2", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 2 @@ -2687,7 +2684,7 @@ const std::map CUDA_IDENTIFIER_MAP{ //IO in FP16 / FP32, computation in float {"cublasSgemmEx", {"hipblasSgemmEx", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, - {"cublasGemmEx", {"hipblasGemmEx", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + {"cublasGemmEx", {"hipblasGemmEx", CONV_MATH_FUNC, API_BLAS}}, {"cublasGemmBatchedEx", {"hipblasGemmBatchedEx", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasGemmStridedBatchedEx", {"hipblasGemmStridedBatchedEx", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, // IO in Int8 complex/cuComplex, computation in cuComplex From 5179a72cda7dfebc45c167755e5d439974f7da92 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Tue, 25 Sep 2018 15:58:36 +0000 Subject: [PATCH 5/8] Use trig functions from ocml instead --- tests/src/deviceLib/hip_trig.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/tests/src/deviceLib/hip_trig.cpp b/tests/src/deviceLib/hip_trig.cpp index 29c24cf5b4..7b076065fa 100644 --- a/tests/src/deviceLib/hip_trig.cpp +++ b/tests/src/deviceLib/hip_trig.cpp @@ -42,10 +42,10 @@ THE SOFTWARE. __global__ void kernel_trig(hipLaunchParm lp, float* In, float* sin_d, float* cos_d, float* tan_d, float* sin_pd, float* cos_pd) { int tid = threadIdx.x + blockIdx.x * blockDim.x; - sin_d[tid] = __sinf(In[tid]); - cos_d[tid] = __cosf(In[tid]); - tan_d[tid] = __tanf(In[tid]); - __sincosf(In[tid], &sin_pd[tid], &cos_pd[tid]); + sin_d[tid] = sinf(In[tid]); + cos_d[tid] = cosf(In[tid]); + tan_d[tid] = tanf(In[tid]); + sincosf(In[tid], &sin_pd[tid], &cos_pd[tid]); } int main() { From 90f57d452a6b75d3c7a70c41072809f57792319e Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 26 Sep 2018 12:47:36 +0530 Subject: [PATCH 6/8] Return hipSuccess when sizeBytes=0 in hipMemset --- src/hip_memory.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp index dc5390f014..bda6ad2650 100644 --- a/src/hip_memory.cpp +++ b/src/hip_memory.cpp @@ -1543,6 +1543,8 @@ hipError_t ihipMemset(void* dst, int value, size_t sizeBytes, hipStream_t strea { hipError_t e = hipSuccess; + if (sizeBytes == 0) return e; + if (stream && (dst != NULL)) { if(copyDataType == ihipMemsetDataTypeChar){ if ((sizeBytes & 0x3) == 0) { From 221ef1db1f343479639da79bd26b4cb8e9530a2f Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 26 Sep 2018 17:01:59 +0300 Subject: [PATCH 7/8] [HIPIFY][doc] Update README.md due to new LLVM 7.0.0 and CUDA 10.0 releases. --- hipify-clang/README.md | 15 +++++++++++---- 1 file changed, 11 insertions(+), 4 deletions(-) diff --git a/hipify-clang/README.md b/hipify-clang/README.md index fdb37d2e2d..d89a50ffbe 100644 --- a/hipify-clang/README.md +++ b/hipify-clang/README.md @@ -29,11 +29,16 @@ ##Dependencies `hipify-clang` requires: -1. LLVM+CLANG of at least version 3.8.0, latest stable and recommended release is 6.0.1. -2. CUDA at least version 7.5, latest supported release is 9.0. +1. LLVM+CLANG of at least version 3.8.0, latest stable and recommended releases: + + 6.0.1 (linux and windows), -| **LLVM release version** | **CUDA latest supported version** | -|:------------------------:|:---------------------------------:| + 7.0.0 (linux only). + +2. CUDA at least version 7.5, latest supported release is 9.2. + +| **LLVM release version** | **CUDA latest supported version** |**Comments**| +|:------------------------:|:---------------------------------:|:-------- -:| | 3.8.0 | 7.5 | | 3.8.1 | 7.5 | | 3.9.0 | 7.5 | @@ -45,6 +50,8 @@ | 5.0.2 | 8.0 | | 6.0.0 | 9.0 | | 6.0.1 | 9.0 | +| 7.0.0 | 9.2 | linux only | +| | 10.0 | not yet supported | In most cases, you can get a suitable version of LLVM+CLANG with your package manager. From c943885a65d60909021cec9bdd9d245995de07a1 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 26 Sep 2018 17:26:25 +0300 Subject: [PATCH 8/8] [HIPIFY][docs] Fix typos in Readme.md --- hipify-clang/README.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/hipify-clang/README.md b/hipify-clang/README.md index 061250a807..02bc9cfed0 100644 --- a/hipify-clang/README.md +++ b/hipify-clang/README.md @@ -38,8 +38,8 @@ 2. CUDA at least version 7.5, latest supported release is 9.2. -| **LLVM release version** | **CUDA latest supported version** |**Comments**| -|:------------------------:|:---------------------------------:|:-------- -:| +| **LLVM release version** | **CUDA latest supported version** | **Comments** | +|:------------------------:|:---------------------------------:|:------------:| | 3.8.0 | 7.5 | | 3.8.1 | 7.5 | | 3.9.0 | 7.5 |