From bc32a77fbeaea6dd37fb4acfff1bde86d7cc59ac Mon Sep 17 00:00:00 2001 From: dlangbe Date: Fri, 26 Jul 2024 20:04:29 +0000 Subject: [PATCH] Added seeding for random fills --- test/01_contraction/contraction_test.cpp | 39 +++++++++++--------- test/02_permutation/permutation_resource.cpp | 6 ++- test/03_reduction/reduction_resource.cpp | 10 +++-- test/device/common.hpp | 9 +++-- test/utils.hpp | 5 ++- 5 files changed, 39 insertions(+), 30 deletions(-) diff --git a/test/01_contraction/contraction_test.cpp b/test/01_contraction/contraction_test.cpp index 078fc79c..0ead586c 100644 --- a/test/01_contraction/contraction_test.cpp +++ b/test/01_contraction/contraction_test.cpp @@ -257,14 +257,17 @@ namespace hiptensor auto resource = getResource(); resource->resizeStorage(lengths, elementBytes); + uint32_t seed = static_cast(std::time(nullptr)); + if(ADataType == HIP_R_16F && BDataType == HIP_R_16F && DDataType == HIP_R_16F) { // Initialize matrix data on device - fillLaunchKernel<_Float16>((_Float16*)resource->deviceA().get(), elementsA); - fillLaunchKernel<_Float16>((_Float16*)resource->deviceB().get(), elementsB); + fillLaunchKernel<_Float16>((_Float16*)resource->deviceA().get(), elementsA, seed - 1); + fillLaunchKernel<_Float16>((_Float16*)resource->deviceB().get(), elementsB, seed); if(CDataType == HIP_R_16F) { - fillLaunchKernel<_Float16>((_Float16*)resource->deviceC().get(), elementsCD); + printf("Filling C\n\n"); + fillLaunchKernel<_Float16>((_Float16*)resource->deviceC().get(), elementsCD, seed + 1); } fillValLaunchKernel<_Float16>((_Float16*)resource->deviceD().get(), elementsCD, @@ -273,12 +276,12 @@ namespace hiptensor else if(ADataType == HIP_R_16BF && BDataType == HIP_R_16BF && DDataType == HIP_R_16BF) { // Initialize matrix data on device - fillLaunchKernel((hip_bfloat16*)resource->deviceA().get(), elementsA); - fillLaunchKernel((hip_bfloat16*)resource->deviceB().get(), elementsB); + fillLaunchKernel((hip_bfloat16*)resource->deviceA().get(), elementsA, seed - 1); + fillLaunchKernel((hip_bfloat16*)resource->deviceB().get(), elementsB, seed); if(CDataType == HIP_R_16BF) { fillLaunchKernel((hip_bfloat16*)resource->deviceC().get(), - elementsCD); + elementsCD, seed + 1); } fillValLaunchKernel( (hip_bfloat16*)resource->deviceD().get(), @@ -288,11 +291,11 @@ namespace hiptensor else if(ADataType == HIP_R_32F && BDataType == HIP_R_32F && DDataType == HIP_R_32F) { // Initialize matrix data on device - fillLaunchKernel((float*)resource->deviceA().get(), elementsA); - fillLaunchKernel((float*)resource->deviceB().get(), elementsB); + fillLaunchKernel((float*)resource->deviceA().get(), elementsA, seed - 1); + fillLaunchKernel((float*)resource->deviceB().get(), elementsB, seed); if(CDataType == HIP_R_32F) { - fillLaunchKernel((float*)resource->deviceC().get(), elementsCD); + fillLaunchKernel((float*)resource->deviceC().get(), elementsCD, seed + 1); } fillValLaunchKernel((float*)resource->deviceD().get(), elementsCD, @@ -301,11 +304,11 @@ namespace hiptensor else if(ADataType == HIP_R_64F && BDataType == HIP_R_64F && DDataType == HIP_R_64F) { // Initialize matrix data on device - fillLaunchKernel((double*)resource->deviceA().get(), elementsA); - fillLaunchKernel((double*)resource->deviceB().get(), elementsB); + fillLaunchKernel((double*)resource->deviceA().get(), elementsA, seed - 1); + fillLaunchKernel((double*)resource->deviceB().get(), elementsB, seed); if(CDataType == HIP_R_64F) { - fillLaunchKernel((double*)resource->deviceC().get(), elementsCD); + fillLaunchKernel((double*)resource->deviceC().get(), elementsCD, seed + 1); } fillValLaunchKernel((double*)resource->deviceD().get(), elementsCD, @@ -315,13 +318,13 @@ namespace hiptensor { // Initialize matrix data on device fillLaunchKernel((hipFloatComplex*)resource->deviceA().get(), - elementsA); + elementsA, seed - 1); fillLaunchKernel((hipFloatComplex*)resource->deviceB().get(), - elementsB); + elementsB, seed); if(CDataType == HIP_C_32F) { fillLaunchKernel((hipFloatComplex*)resource->deviceC().get(), - elementsCD); + elementsCD, seed + 1); } fillValLaunchKernel( (hipFloatComplex*)resource->deviceD().get(), @@ -332,13 +335,13 @@ namespace hiptensor { // Initialize matrix data on device fillLaunchKernel((hipDoubleComplex*)resource->deviceA().get(), - elementsA); + elementsA, seed - 1); fillLaunchKernel((hipDoubleComplex*)resource->deviceB().get(), - elementsB); + elementsB, seed); if(CDataType == HIP_C_64F) { fillLaunchKernel((hipDoubleComplex*)resource->deviceC().get(), - elementsCD); + elementsCD, seed + 1); } fillValLaunchKernel( (hipDoubleComplex*)resource->deviceD().get(), diff --git a/test/02_permutation/permutation_resource.cpp b/test/02_permutation/permutation_resource.cpp index 6acd7577..5a03aa37 100644 --- a/test/02_permutation/permutation_resource.cpp +++ b/test/02_permutation/permutation_resource.cpp @@ -96,13 +96,15 @@ namespace hiptensor void PermutationResource::fillRandToA() { + uint32_t seed = static_cast(std::time(nullptr)); + if(mCurrentDataType == HIP_R_32F) { - fillLaunchKernel((float*)deviceA().get(), mCurrentMatrixElement); + fillLaunchKernel((float*)deviceA().get(), mCurrentMatrixElement, seed); } else { - fillLaunchKernel<_Float16>((_Float16*)deviceA().get(), mCurrentMatrixElement); + fillLaunchKernel<_Float16>((_Float16*)deviceA().get(), mCurrentMatrixElement, seed); } Base::copyData(hostA(), deviceA(), getCurrentMatrixMemorySize()); } diff --git a/test/03_reduction/reduction_resource.cpp b/test/03_reduction/reduction_resource.cpp index 22e1fb32..f215800e 100644 --- a/test/03_reduction/reduction_resource.cpp +++ b/test/03_reduction/reduction_resource.cpp @@ -127,21 +127,23 @@ namespace hiptensor void ReductionResource::fillRand(HostPtrT& hostBuf, DevicePtrT& deviceBuf, size_t elementCount) { + uint32_t seed = static_cast(std::time(nullptr)); + if(mCurrentDataType == HIP_R_16F) { - fillLaunchKernel((float16_t*)deviceBuf.get(), elementCount); + fillLaunchKernel((float16_t*)deviceBuf.get(), elementCount, seed); } else if(mCurrentDataType == HIP_R_16BF) { - fillLaunchKernel((bfloat16_t*)deviceBuf.get(), elementCount); + fillLaunchKernel((bfloat16_t*)deviceBuf.get(), elementCount, seed); } else if(mCurrentDataType == HIP_R_32F) { - fillLaunchKernel((float32_t*)deviceBuf.get(), elementCount); + fillLaunchKernel((float32_t*)deviceBuf.get(), elementCount, seed); } else if(mCurrentDataType == HIP_R_64F) { - fillLaunchKernel((float64_t*)deviceBuf.get(), elementCount); + fillLaunchKernel((float64_t*)deviceBuf.get(), elementCount, seed); } Base::copyData(hostBuf, deviceBuf, elementCount); } diff --git a/test/device/common.hpp b/test/device/common.hpp index 73410eea..66b57231 100644 --- a/test/device/common.hpp +++ b/test/device/common.hpp @@ -82,25 +82,26 @@ __device__ inline float gen_random_float(unsigned input) // fill kernel for 'elementSize' elements template -__global__ void fillKernel(DataType* data, uint32_t elementSize) +__global__ void fillKernel(DataType* data, uint32_t elementSize, uint32_t seed) { uint32_t index = (blockIdx.x * blockDim.x + threadIdx.x); + uint32_t seededIndex = static_cast(uint64_t(index + seed) % UINT_MAX); if(index < elementSize) { if constexpr(std::is_same_v) { - auto value = gen_random_float(index); + auto value = gen_random_float(seededIndex); data[index] = make_hipFloatComplex(value, value); } else if constexpr(std::is_same_v) { - auto value = static_cast(gen_random_float(index)); + auto value = static_cast(gen_random_float(seededIndex)); data[index] = make_hipDoubleComplex(value, value); } else { - auto value = gen_random_float(index); + auto value = gen_random_float(seededIndex); data[index] = static_cast(value); } } diff --git a/test/utils.hpp b/test/utils.hpp index 8edb8551..324c7b06 100644 --- a/test/utils.hpp +++ b/test/utils.hpp @@ -164,7 +164,7 @@ auto getProduct(const Container& container, // fill kernel for 'elementSize' elements template -__host__ static inline void fillLaunchKernel(DataType* data, uint32_t elementSize) +__host__ static inline void fillLaunchKernel(DataType* data, uint32_t elementSize, uint32_t seed) { auto blockDim = dim3(1024, 1, 1); auto gridDim = dim3(ceilDiv(elementSize, blockDim.x), 1, 1); @@ -174,7 +174,8 @@ __host__ static inline void fillLaunchKernel(DataType* data, uint32_t elementSiz 0, 0, data, - elementSize); + elementSize, + seed); } // fill kernel wrapper for 'elementSize' elements with a specific value