Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

New fill kernel and validation tolerances #251

Merged
merged 4 commits into from
Aug 1, 2024
Merged
Show file tree
Hide file tree
Changes from 1 commit
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
39 changes: 21 additions & 18 deletions test/01_contraction/contraction_test.cpp
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));
dlangbe marked this conversation as resolved.
Show resolved Hide resolved

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");
dlangbe marked this conversation as resolved.
Show resolved Hide resolved
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
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
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
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
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
Loading