Log1p #377
Open
Log1p #377
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
HOST Implementation
| * \retval RPP_SUCCESS Successful completion. | ||
| * \retval RPP_ERROR* Unsuccessful completion. | ||
| */ | ||
| RppStatus rppt_log1p_gpu(RppPtr_t srcPtr, RpptGenericDescPtr srcGenericDescPtr, RppPtr_t dstPtr, RpptGenericDescPtr dstGenericDescPtr, Rpp32u *roiTensor, rppHandle_t rppHandle); |
There was a problem hiding this comment.
please check the gpu declaration, it is repeated twice
src/include/cpu/rpp_cpu_simd.hpp
Outdated
| @@ -148,6 +149,8 @@ const __m128i xmm_pxMask04To07 = _mm_setr_epi8(4, 0x80, 0x80, 0x80, 5, 0x80, 0x8 | |||
| const __m128i xmm_pxMask08To11 = _mm_setr_epi8(8, 0x80, 0x80, 0x80, 9, 0x80, 0x80, 0x80, 10, 0x80, 0x80, 0x80, 11, 0x80, 0x80, 0x80); | |||
| const __m128i xmm_pxMask12To15 = _mm_setr_epi8(12, 0x80, 0x80, 0x80, 13, 0x80, 0x80, 0x80, 14, 0x80, 0x80, 0x80, 15, 0x80, 0x80, 0x80); | |||
|
|
|||
|
|
|||
There was a problem hiding this comment.
please remove these extra spaces
src/include/cpu/rpp_cpu_simd.hpp
Outdated
| @@ -2693,6 +2713,8 @@ static inline __m256 log_ps(__m256 x) | |||
| __m256 one = *(__m256 *)&avx_p1; | |||
| __m256 invalid_mask = _mm256_cmp_ps(x, avx_p0, _CMP_LE_OQ); | |||
|
|
|||
| // x = _mm256_add_ps(x, one); | |||
There was a problem hiding this comment.
please remove the commented-out lines, if these are not necessary
| @@ -464,6 +464,18 @@ __device__ __forceinline__ float rpp_hip_unpack3(int src) | |||
| { | |||
| return (float)(schar)((src >> 24) & 0xFF); | |||
| } | |||
| // Un-Packing from I16s | |||
|
|
|||
| else if ((srcGenericDescPtr->dataType == RpptDataType::I8) && (dstGenericDescPtr->dataType == RpptDataType::I8)) return RPP_ERROR_INVALID_DST_DATATYPE; | ||
| else if ((srcGenericDescPtr->dataType == RpptDataType::I16) && (dstGenericDescPtr->dataType == RpptDataType::F32)) | ||
| { | ||
| log1p_generic_host_tensor(static_cast<Rpp16s *>(srcPtr) + srcGenericDescPtr->offsetInBytes, |
| @@ -559,4 +594,66 @@ RppStatus rppt_log_gpu(RppPtr_t srcPtr, | |||
| #endif // backend | |||
| } | |||
|
|
|||
| RppStatus rppt_log1p_gpu(RppPtr_t srcPtr, | |||
| set_generic_descriptor(srcDescriptorPtrND, nDim, offSetInBytes, bitDepth, batchSize, roiTensor); | ||
| set_generic_descriptor(dstDescriptorPtrND, nDim, offSetInBytes, bitDepth, batchSize, roiTensor); | ||
| set_generic_descriptor(dstDescriptorPtrND, nDim, offSetInBytes, 2, batchSize, roiTensor); |
There was a problem hiding this comment.
I guess you have hardcoded the value for your use case, please revert these changes
| @@ -129,7 +142,7 @@ void fill_roi_values(Rpp32u nDim, Rpp32u batchSize, Rpp32u *roiTensor, bool qaMo | |||
| } | |||
| case 4: | |||
| { | |||
| std::array<Rpp32u, 8> roi = {0, 0, 0, 0, 1, 128, 128, 128}; | |||
| std::array<Rpp32u, 8> roi = {0, 0, 0, 0, 128, 128, 128, 4}; | |||
| @@ -369,8 +384,14 @@ void compare_output(Rpp32f *outputF32, Rpp32u nDim, Rpp32u batchSize, Rpp32u buf | |||
| for(int j = 0; j < sampleLength; j++) | |||
| { | |||
| bool invalid_comparision = ((out[j] == 0.0f) && (ref[j] != 0.0f)); | |||
| if(!invalid_comparision && abs(out[j] - ref[j]) < 1e-4) | |||
| cnt++; | |||
| if(!invalid_comparision ) | |||
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Host and HIP implementation of log1p , taken reference from log
In HOST in AVX2, for 3d and 4d combined the dimension to vectorize.
In HIP for log in 3d implementation as there is usage of multiple kernel launches, used Nd for 3d as well in log1p
And some changes were hardcoded in the test suite to make it able to run for I16 datatype, which are not meant to be pushed but just to evaluate the QA test.