Skip to content

Commit

Permalink
Added seeding for random fills
Browse files Browse the repository at this point in the history
  • Loading branch information
@dlangbe
dlangbe committed Jul 26, 2024
1 parent fcdb964 commit bc32a77
Show file tree
Hide file tree
Showing 5 changed files with 39 additions and 30 deletions.
39 changes: 21 additions & 18 deletions test/01_contraction/contraction_test.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -257,14 +257,17 @@ namespace hiptensor
auto resource = getResource();
resource->resizeStorage(lengths, elementBytes);

uint32_t seed = static_cast<uint32_t>(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,
Expand All @@ -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>((hip_bfloat16*)resource->deviceA().get(), elementsA);
fillLaunchKernel<hip_bfloat16>((hip_bfloat16*)resource->deviceB().get(), elementsB);
fillLaunchKernel<hip_bfloat16>((hip_bfloat16*)resource->deviceA().get(), elementsA, seed - 1);
fillLaunchKernel<hip_bfloat16>((hip_bfloat16*)resource->deviceB().get(), elementsB, seed);
if(CDataType == HIP_R_16BF)
{
fillLaunchKernel<hip_bfloat16>((hip_bfloat16*)resource->deviceC().get(),
elementsCD);
elementsCD, seed + 1);
}
fillValLaunchKernel<hip_bfloat16>(
(hip_bfloat16*)resource->deviceD().get(),
Expand All @@ -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>((float*)resource->deviceA().get(), elementsA);
fillLaunchKernel<float>((float*)resource->deviceB().get(), elementsB);
fillLaunchKernel<float>((float*)resource->deviceA().get(), elementsA, seed - 1);
fillLaunchKernel<float>((float*)resource->deviceB().get(), elementsB, seed);
if(CDataType == HIP_R_32F)
{
fillLaunchKernel<float>((float*)resource->deviceC().get(), elementsCD);
fillLaunchKernel<float>((float*)resource->deviceC().get(), elementsCD, seed + 1);
}
fillValLaunchKernel<float>((float*)resource->deviceD().get(),
elementsCD,
Expand All @@ -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>((double*)resource->deviceA().get(), elementsA);
fillLaunchKernel<double>((double*)resource->deviceB().get(), elementsB);
fillLaunchKernel<double>((double*)resource->deviceA().get(), elementsA, seed - 1);
fillLaunchKernel<double>((double*)resource->deviceB().get(), elementsB, seed);
if(CDataType == HIP_R_64F)
{
fillLaunchKernel<double>((double*)resource->deviceC().get(), elementsCD);
fillLaunchKernel<double>((double*)resource->deviceC().get(), elementsCD, seed + 1);
}
fillValLaunchKernel<double>((double*)resource->deviceD().get(),
elementsCD,
Expand All @@ -315,13 +318,13 @@ namespace hiptensor
{
// Initialize matrix data on device
fillLaunchKernel<hipFloatComplex>((hipFloatComplex*)resource->deviceA().get(),
elementsA);
elementsA, seed - 1);
fillLaunchKernel<hipFloatComplex>((hipFloatComplex*)resource->deviceB().get(),
elementsB);
elementsB, seed);
if(CDataType == HIP_C_32F)
{
fillLaunchKernel<hipFloatComplex>((hipFloatComplex*)resource->deviceC().get(),
elementsCD);
elementsCD, seed + 1);
}
fillValLaunchKernel<hipFloatComplex>(
(hipFloatComplex*)resource->deviceD().get(),
Expand All @@ -332,13 +335,13 @@ namespace hiptensor
{
// Initialize matrix data on device
fillLaunchKernel<hipDoubleComplex>((hipDoubleComplex*)resource->deviceA().get(),
elementsA);
elementsA, seed - 1);
fillLaunchKernel<hipDoubleComplex>((hipDoubleComplex*)resource->deviceB().get(),
elementsB);
elementsB, seed);
if(CDataType == HIP_C_64F)
{
fillLaunchKernel<hipDoubleComplex>((hipDoubleComplex*)resource->deviceC().get(),
elementsCD);
elementsCD, seed + 1);
}
fillValLaunchKernel<hipDoubleComplex>(
(hipDoubleComplex*)resource->deviceD().get(),
Expand Down
6 changes: 4 additions & 2 deletions test/02_permutation/permutation_resource.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -96,13 +96,15 @@ namespace hiptensor

void PermutationResource::fillRandToA()
{
uint32_t seed = static_cast<uint32_t>(std::time(nullptr));

if(mCurrentDataType == HIP_R_32F)
{
fillLaunchKernel<float>((float*)deviceA().get(), mCurrentMatrixElement);
fillLaunchKernel<float>((float*)deviceA().get(), mCurrentMatrixElement, seed);
}
else
{
fillLaunchKernel<_Float16>((_Float16*)deviceA().get(), mCurrentMatrixElement);
fillLaunchKernel<_Float16>((_Float16*)deviceA().get(), mCurrentMatrixElement, seed);
}
Base::copyData(hostA(), deviceA(), getCurrentMatrixMemorySize());
}
Expand Down
10 changes: 6 additions & 4 deletions test/03_reduction/reduction_resource.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -127,21 +127,23 @@ namespace hiptensor

void ReductionResource::fillRand(HostPtrT& hostBuf, DevicePtrT& deviceBuf, size_t elementCount)
{
uint32_t seed = static_cast<uint32_t>(std::time(nullptr));

if(mCurrentDataType == HIP_R_16F)
{
fillLaunchKernel<float16_t>((float16_t*)deviceBuf.get(), elementCount);
fillLaunchKernel<float16_t>((float16_t*)deviceBuf.get(), elementCount, seed);
}
else if(mCurrentDataType == HIP_R_16BF)
{
fillLaunchKernel<bfloat16_t>((bfloat16_t*)deviceBuf.get(), elementCount);
fillLaunchKernel<bfloat16_t>((bfloat16_t*)deviceBuf.get(), elementCount, seed);
}
else if(mCurrentDataType == HIP_R_32F)
{
fillLaunchKernel<float32_t>((float32_t*)deviceBuf.get(), elementCount);
fillLaunchKernel<float32_t>((float32_t*)deviceBuf.get(), elementCount, seed);
}
else if(mCurrentDataType == HIP_R_64F)
{
fillLaunchKernel<float64_t>((float64_t*)deviceBuf.get(), elementCount);
fillLaunchKernel<float64_t>((float64_t*)deviceBuf.get(), elementCount, seed);
}
Base::copyData(hostBuf, deviceBuf, elementCount);
}
Expand Down
9 changes: 5 additions & 4 deletions test/device/common.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -82,25 +82,26 @@ __device__ inline float gen_random_float(unsigned input)

// fill kernel for 'elementSize' elements
template <typename DataType>
__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<uint32_t>(uint64_t(index + seed) % UINT_MAX);

if(index < elementSize)
{
if constexpr(std::is_same_v<DataType, hipFloatComplex>)
{
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<DataType, hipDoubleComplex>)
{
auto value = static_cast<double>(gen_random_float(index));
auto value = static_cast<double>(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<DataType>(value);
}
}
Expand Down
5 changes: 3 additions & 2 deletions test/utils.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -164,7 +164,7 @@ auto getProduct(const Container& container,

// fill kernel for 'elementSize' elements
template <typename DataType>
__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);
Expand All @@ -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
Expand Down

0 comments on commit bc32a77

Please sign in to comment.