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/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` |
diff --git a/hipify-clang/README.md b/hipify-clang/README.md
index 7f55dea800..02bc9cfed0 100644
--- a/hipify-clang/README.md
+++ b/hipify-clang/README.md
@@ -30,11 +30,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 |
@@ -46,6 +51,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.
diff --git a/hipify-clang/src/CUDA2HipMap.cpp b/hipify-clang/src/CUDA2HipMap.cpp
index 4bd7263bad..01f24a6d43 100644
--- a/hipify-clang/src/CUDA2HipMap.cpp
+++ b/hipify-clang/src/CUDA2HipMap.cpp
@@ -1417,22 +1417,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
@@ -2143,11 +2144,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}},
@@ -2155,9 +2152,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
@@ -2715,7 +2712,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
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) {
diff --git a/tests/src/deviceLib/hip_trig.cpp b/tests/src/deviceLib/hip_trig.cpp
index 3f200fcf8f..7b076065fa 100644
--- a/tests/src/deviceLib/hip_trig.cpp
+++ b/tests/src/deviceLib/hip_trig.cpp
@@ -32,21 +32,26 @@ THE SOFTWARE.
#include "test_common.h"
#include
+#define HIP_ASSERT(x) (assert((x) == hipSuccess))
+
#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;
- 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() {
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 +66,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 +84,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;
}