Skip to content

Commit 62aae36

Browse files
illsilinillsilin
authored
Merge branch 'develop' into import/develop/ROCm_composable_kernel/pr-3652
2 parents 1d1c3a9 + 52def72 commit 62aae36

File tree

222 files changed

+19152
-25577
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

222 files changed

+19152
-25577
lines changed

‎dnn-providers/hipblaslt-provider/README.md‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10,3 +10,7 @@ This plugin is built as a standalone plugin. To build the plugin, first install
1010
1. Make a build directory using `mkdir build && cd build`.
1111
1. Configure the build using `cmake -DCMAKE_CXX_COMPILER=<path to amdclang>/clang++ ..`.
1212
1. Finally, run `ninja` to build the plugin.
13+
14+
## Operation support
15+
16+
The list of supported operations is described in [Operation Support](docs/OperationSupport.md) documentation.

‎dnn-providers/hipblaslt-provider/docs/OperationSupport.md‎

Lines changed: 25 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,25 @@
1+
# hipBLASLt Provider Plugin - Operation Support
2+
3+
This document provides detailed information about the operations supported by the hipBLASLt Provider Plugin for hipDNN.
4+
5+
For general information about hipDNN's operation support, please see the [hipDNN Operation Support](../../../projects/hipdnn/docs/OperationSupport.md) documentation.
6+
7+
## Current Operation Support
8+
9+
hipBLASLt Provider Plugin currently supports only stand-alone Matmul (GEMM, general matrix multiplication) operations with the following features and constraints:
10+
- Input and output data types: FP32, FP16, BF16
11+
- Compute data type: FP32
12+
- Transposed inputs: supported
13+
- Batched matmuls: only equal batch sizes are supported, or broadcasting when one input has a single batch (batch=1)
14+
15+
## Notes
16+
17+
> [!NOTE]
18+
> **Fused Operations:** Currently the kernel provider doesn't support any fusions.
19+
20+
## Legend
21+
22+
### Datatypes
23+
- **FP16**: Half-precision floating point (16-bit)
24+
- **BFP16**: Brain floating point (16-bit)
25+
- **FP32**: Single-precision floating point (32-bit)

‎dnn-providers/hipblaslt-provider/integration_tests/IntegrationGpuMatmul.cpp‎

Lines changed: 6 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -5,6 +5,7 @@
55
#include <random>
66

77
#include <hip/hip_runtime.h>
8+
#include <hipdnn_data_sdk/types.hpp>
89
#include <hipdnn_data_sdk/utilities/PlatformUtils.hpp>
910
#include <hipdnn_test_sdk/utilities/CpuFpReferenceValidation.hpp>
1011
#include <hipdnn_test_sdk/utilities/TestTolerances.hpp>
@@ -26,7 +27,7 @@ template <typename DataType>
2627
class IntegrationGpuMatmul : public IntegrationGraphVerificationHarness<DataType, MatmulTestCase>
2728
{
2829
protected:
29-
void runGraphTest(DataType tolerance) override
30+
void runGraphTest(float tolerance) override
3031
{
3132
const MatmulTestCase& testCase = this->GetParam();
3233

@@ -73,8 +74,8 @@ class IntegrationGpuMatmul : public IntegrationGraphVerificationHarness<DataType
7374
};
7475

7576
using IntegrationGpuMatmulFp32 = IntegrationGpuMatmul<float>;
76-
using IntegrationGpuMatmulFp16 = IntegrationGpuMatmul<half>;
77-
using IntegrationGpuMatmulBf16 = IntegrationGpuMatmul<hip_bfloat16>;
77+
using IntegrationGpuMatmulFp16 = IntegrationGpuMatmul<hipdnn_data_sdk::types::half>;
78+
using IntegrationGpuMatmulBf16 = IntegrationGpuMatmul<hipdnn_data_sdk::types::bfloat16>;
7879

7980
} // namespace
8081

@@ -85,12 +86,12 @@ TEST_P(IntegrationGpuMatmulFp32, Correctness)
8586

8687
TEST_P(IntegrationGpuMatmulFp16, Correctness)
8788
{
88-
runGraphTest(matmul::getTolerance<half>());
89+
runGraphTest(matmul::getTolerance<hipdnn_data_sdk::types::half>());
8990
}
9091

9192
TEST_P(IntegrationGpuMatmulBf16, Correctness)
9293
{
93-
runGraphTest(matmul::getTolerance<hip_bfloat16>());
94+
runGraphTest(matmul::getTolerance<hipdnn_data_sdk::types::bfloat16>());
9495
}
9596

9697
INSTANTIATE_TEST_SUITE_P(IntegrationGpuMatmul,

‎dnn-providers/hipblaslt-provider/integration_tests/IntegrationGraphVerificationHarness.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -66,7 +66,7 @@ class IntegrationGraphVerificationHarness : public ::testing::TestWithParam<Test
6666
}
6767
}
6868

69-
virtual void runGraphTest(DataType tolerance) = 0;
69+
virtual void runGraphTest(float tolerance) = 0;
7070

7171
protected:
7272
void verifyGraph(hipdnn_frontend::graph::Graph& graph, unsigned int seed)

‎dnn-providers/miopen-provider/integration_tests/IntegrationGpuBatchnormBackward.cpp‎

Lines changed: 17 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -58,7 +58,7 @@ class BatchnormBackward : public IntegrationGraphVerificationHarness<DataType, B
5858
}
5959
}
6060

61-
void runGraphTest(DataType tolerance, const TensorLayout& layout = TensorLayout::NCHW)
61+
void runGraphTest(float tolerance, const TensorLayout& layout = TensorLayout::NCHW)
6262
{
6363
const BatchnormTestCase& testCase = this->GetParam();
6464

@@ -134,53 +134,49 @@ class BatchnormBackward : public IntegrationGraphVerificationHarness<DataType, B
134134

135135
using IntegrationGpuBatchnormBackwardNchwFp32 = BatchnormBackward<float, float>;
136136

137-
using IntegrationGpuBatchnormBackwardNchwBfp16 = BatchnormBackward<hip_bfloat16, float>;
137+
using IntegrationGpuBatchnormBackwardNchwBfp16 = BatchnormBackward<bfloat16, float>;
138138

139139
using IntegrationGpuBatchnormBackwardNchwFp16 = BatchnormBackward<half, float>;
140140

141141
using IntegrationGpuBatchnormBackwardNhwcFp32 = BatchnormBackward<float, float>;
142142

143-
using IntegrationGpuBatchnormBackwardNhwcBfp16 = BatchnormBackward<hip_bfloat16, float>;
143+
using IntegrationGpuBatchnormBackwardNhwcBfp16 = BatchnormBackward<bfloat16, float>;
144144

145145
using IntegrationGpuBatchnormBackwardNhwcFp16 = BatchnormBackward<half, float>;
146146

147147
using IntegrationGpuBatchnormBackwardNcdhwFp32 = BatchnormBackward<float, float>;
148148

149-
using IntegrationGpuBatchnormBackwardNcdhwBfp16 = BatchnormBackward<hip_bfloat16, float>;
149+
using IntegrationGpuBatchnormBackwardNcdhwBfp16 = BatchnormBackward<bfloat16, float>;
150150

151151
using IntegrationGpuBatchnormBackwardNcdhwFp16 = BatchnormBackward<half, float>;
152152

153153
using IntegrationGpuBatchnormBackwardNdhwcFp32 = BatchnormBackward<float, float>;
154154

155-
using IntegrationGpuBatchnormBackwardNdhwcBfp16 = BatchnormBackward<hip_bfloat16, float>;
155+
using IntegrationGpuBatchnormBackwardNdhwcBfp16 = BatchnormBackward<bfloat16, float>;
156156

157157
using IntegrationGpuBatchnormBackwardNdhwcFp16 = BatchnormBackward<half, float>;
158158

159159
using IntegrationGpuBatchnormBackwardCalcStatsNchwFp32 = BatchnormBackward<float, float, true>;
160160

161-
using IntegrationGpuBatchnormBackwardCalcStatsNchwBfp16
162-
= BatchnormBackward<hip_bfloat16, float, true>;
161+
using IntegrationGpuBatchnormBackwardCalcStatsNchwBfp16 = BatchnormBackward<bfloat16, float, true>;
163162

164163
using IntegrationGpuBatchnormBackwardCalcStatsNchwFp16 = BatchnormBackward<half, float, true>;
165164

166165
using IntegrationGpuBatchnormBackwardCalcStatsNhwcFp32 = BatchnormBackward<float, float, true>;
167166

168-
using IntegrationGpuBatchnormBackwardCalcStatsNhwcBfp16
169-
= BatchnormBackward<hip_bfloat16, float, true>;
167+
using IntegrationGpuBatchnormBackwardCalcStatsNhwcBfp16 = BatchnormBackward<bfloat16, float, true>;
170168

171169
using IntegrationGpuBatchnormBackwardCalcStatsNhwcFp16 = BatchnormBackward<half, float, true>;
172170

173171
using IntegrationGpuBatchnormBackwardCalcStatsNcdhwFp32 = BatchnormBackward<float, float, true>;
174172

175-
using IntegrationGpuBatchnormBackwardCalcStatsNcdhwBfp16
176-
= BatchnormBackward<hip_bfloat16, float, true>;
173+
using IntegrationGpuBatchnormBackwardCalcStatsNcdhwBfp16 = BatchnormBackward<bfloat16, float, true>;
177174

178175
using IntegrationGpuBatchnormBackwardCalcStatsNcdhwFp16 = BatchnormBackward<half, float, true>;
179176

180177
using IntegrationGpuBatchnormBackwardCalcStatsNdhwcFp32 = BatchnormBackward<float, float, true>;
181178

182-
using IntegrationGpuBatchnormBackwardCalcStatsNdhwcBfp16
183-
= BatchnormBackward<hip_bfloat16, float, true>;
179+
using IntegrationGpuBatchnormBackwardCalcStatsNdhwcBfp16 = BatchnormBackward<bfloat16, float, true>;
184180

185181
using IntegrationGpuBatchnormBackwardCalcStatsNdhwcFp16 = BatchnormBackward<half, float, true>;
186182

@@ -201,7 +197,7 @@ INSTANTIATE_TEST_SUITE_P(Full,
201197

202198
TEST_P(IntegrationGpuBatchnormBackwardNchwBfp16, Correctness)
203199
{
204-
runGraphTest(batchnorm::getToleranceBackward<hip_bfloat16>(), TensorLayout::NCHW);
200+
runGraphTest(batchnorm::getToleranceBackward<bfloat16>(), TensorLayout::NCHW);
205201
}
206202

207203
INSTANTIATE_TEST_SUITE_P(Smoke,
@@ -240,7 +236,7 @@ INSTANTIATE_TEST_SUITE_P(Full,
240236

241237
TEST_P(IntegrationGpuBatchnormBackwardNhwcBfp16, Correctness)
242238
{
243-
runGraphTest(batchnorm::getToleranceBackward<hip_bfloat16>(), TensorLayout::NHWC);
239+
runGraphTest(batchnorm::getToleranceBackward<bfloat16>(), TensorLayout::NHWC);
244240
}
245241

246242
INSTANTIATE_TEST_SUITE_P(Smoke,
@@ -275,7 +271,7 @@ INSTANTIATE_TEST_SUITE_P(Smoke,
275271

276272
TEST_P(IntegrationGpuBatchnormBackwardNcdhwBfp16, Correctness)
277273
{
278-
runGraphTest(batchnorm::getToleranceBackward<hip_bfloat16>(), TensorLayout::NCDHW);
274+
runGraphTest(batchnorm::getToleranceBackward<bfloat16>(), TensorLayout::NCDHW);
279275
}
280276

281277
INSTANTIATE_TEST_SUITE_P(Smoke,
@@ -302,7 +298,7 @@ INSTANTIATE_TEST_SUITE_P(Smoke,
302298

303299
TEST_P(IntegrationGpuBatchnormBackwardNdhwcBfp16, Correctness)
304300
{
305-
runGraphTest(batchnorm::getToleranceBackward<hip_bfloat16>(), TensorLayout::NDHWC);
301+
runGraphTest(batchnorm::getToleranceBackward<bfloat16>(), TensorLayout::NDHWC);
306302
}
307303

308304
INSTANTIATE_TEST_SUITE_P(Smoke,
@@ -329,7 +325,7 @@ INSTANTIATE_TEST_SUITE_P(Smoke,
329325

330326
TEST_P(IntegrationGpuBatchnormBackwardCalcStatsNchwBfp16, Correctness)
331327
{
332-
runGraphTest(batchnorm::getToleranceBackward<hip_bfloat16>(), TensorLayout::NCHW);
328+
runGraphTest(batchnorm::getToleranceBackward<bfloat16>(), TensorLayout::NCHW);
333329
}
334330

335331
INSTANTIATE_TEST_SUITE_P(Smoke,
@@ -356,7 +352,7 @@ INSTANTIATE_TEST_SUITE_P(Smoke,
356352

357353
TEST_P(IntegrationGpuBatchnormBackwardCalcStatsNhwcBfp16, Correctness)
358354
{
359-
runGraphTest(batchnorm::getToleranceBackward<hip_bfloat16>(), TensorLayout::NHWC);
355+
runGraphTest(batchnorm::getToleranceBackward<bfloat16>(), TensorLayout::NHWC);
360356
}
361357

362358
INSTANTIATE_TEST_SUITE_P(Smoke,
@@ -383,7 +379,7 @@ INSTANTIATE_TEST_SUITE_P(Smoke,
383379

384380
TEST_P(IntegrationGpuBatchnormBackwardCalcStatsNcdhwBfp16, Correctness)
385381
{
386-
runGraphTest(batchnorm::getToleranceBackward<hip_bfloat16>(), TensorLayout::NCDHW);
382+
runGraphTest(batchnorm::getToleranceBackward<bfloat16>(), TensorLayout::NCDHW);
387383
}
388384

389385
INSTANTIATE_TEST_SUITE_P(Smoke,
@@ -410,7 +406,7 @@ INSTANTIATE_TEST_SUITE_P(Smoke,
410406

411407
TEST_P(IntegrationGpuBatchnormBackwardCalcStatsNdhwcBfp16, Correctness)
412408
{
413-
runGraphTest(batchnorm::getToleranceBackward<hip_bfloat16>(), TensorLayout::NDHWC);
409+
runGraphTest(batchnorm::getToleranceBackward<bfloat16>(), TensorLayout::NDHWC);
414410
}
415411

416412
INSTANTIATE_TEST_SUITE_P(Smoke,

‎dnn-providers/miopen-provider/integration_tests/IntegrationGpuBatchnormBackwardActivation.cpp‎

Lines changed: 10 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -60,7 +60,7 @@ class BatchnormBackwardActivation
6060
->fillTensorWithRandomValues(1.9f, 2.0f, seed);
6161
}
6262

63-
void runGraphTest([[maybe_unused]] DataType tolerance, const TensorLayout& layout)
63+
void runGraphTest([[maybe_unused]] float tolerance, const TensorLayout& layout)
6464
{
6565
namespace fe = hipdnn_frontend;
6666

@@ -175,7 +175,7 @@ class BatchnormBackwardActivation
175175

176176
auto intermediateTolerance = batchnorm::getToleranceBackward<float>();
177177

178-
this->registerValidator(dxOut, static_cast<float>(tolerance));
178+
this->registerValidator(dxOut, tolerance);
179179
this->registerValidator(dscaleOut, intermediateTolerance);
180180
this->registerValidator(dbiasOut, intermediateTolerance);
181181

@@ -185,29 +185,25 @@ class BatchnormBackwardActivation
185185

186186
using IntegrationGpuBatchnormBackwardActivationNchwFp32 = BatchnormBackwardActivation<float>;
187187

188-
using IntegrationGpuBatchnormBackwardActivationNchwBfp16
189-
= BatchnormBackwardActivation<hip_bfloat16>;
188+
using IntegrationGpuBatchnormBackwardActivationNchwBfp16 = BatchnormBackwardActivation<bfloat16>;
190189

191190
using IntegrationGpuBatchnormBackwardActivationNchwFp16 = BatchnormBackwardActivation<half>;
192191

193192
using IntegrationGpuBatchnormBackwardActivationNhwcFp32 = BatchnormBackwardActivation<float>;
194193

195-
using IntegrationGpuBatchnormBackwardActivationNhwcBfp16
196-
= BatchnormBackwardActivation<hip_bfloat16>;
194+
using IntegrationGpuBatchnormBackwardActivationNhwcBfp16 = BatchnormBackwardActivation<bfloat16>;
197195

198196
using IntegrationGpuBatchnormBackwardActivationNhwcFp16 = BatchnormBackwardActivation<half>;
199197

200198
using IntegrationGpuBatchnormBackwardActivationNcdhwFp32 = BatchnormBackwardActivation<float>;
201199

202-
using IntegrationGpuBatchnormBackwardActivationNcdhwBfp16
203-
= BatchnormBackwardActivation<hip_bfloat16>;
200+
using IntegrationGpuBatchnormBackwardActivationNcdhwBfp16 = BatchnormBackwardActivation<bfloat16>;
204201

205202
using IntegrationGpuBatchnormBackwardActivationNcdhwFp16 = BatchnormBackwardActivation<half>;
206203

207204
using IntegrationGpuBatchnormBackwardActivationNdhwcFp32 = BatchnormBackwardActivation<float>;
208205

209-
using IntegrationGpuBatchnormBackwardActivationNdhwcBfp16
210-
= BatchnormBackwardActivation<hip_bfloat16>;
206+
using IntegrationGpuBatchnormBackwardActivationNdhwcBfp16 = BatchnormBackwardActivation<bfloat16>;
211207

212208
using IntegrationGpuBatchnormBackwardActivationNdhwcFp16 = BatchnormBackwardActivation<half>;
213209

@@ -234,7 +230,7 @@ INSTANTIATE_TEST_SUITE_P(
234230

235231
TEST_P(IntegrationGpuBatchnormBackwardActivationNchwBfp16, Correctness)
236232
{
237-
runGraphTest(batchnorm::getToleranceBackward<hip_bfloat16>(), TensorLayout::NCHW);
233+
runGraphTest(batchnorm::getToleranceBackward<bfloat16>(), TensorLayout::NCHW);
238234
}
239235

240236
INSTANTIATE_TEST_SUITE_P(
@@ -291,7 +287,7 @@ INSTANTIATE_TEST_SUITE_P(
291287

292288
TEST_P(IntegrationGpuBatchnormBackwardActivationNhwcBfp16, Correctness)
293289
{
294-
runGraphTest(batchnorm::getToleranceBackward<hip_bfloat16>(), TensorLayout::NHWC);
290+
runGraphTest(batchnorm::getToleranceBackward<bfloat16>(), TensorLayout::NHWC);
295291
}
296292

297293
INSTANTIATE_TEST_SUITE_P(
@@ -341,7 +337,7 @@ INSTANTIATE_TEST_SUITE_P(
341337

342338
TEST_P(IntegrationGpuBatchnormBackwardActivationNcdhwBfp16, Correctness)
343339
{
344-
runGraphTest(batchnorm::getToleranceBackward<hip_bfloat16>(), TensorLayout::NCDHW);
340+
runGraphTest(batchnorm::getToleranceBackward<bfloat16>(), TensorLayout::NCDHW);
345341
}
346342

347343
INSTANTIATE_TEST_SUITE_P(
@@ -377,7 +373,7 @@ INSTANTIATE_TEST_SUITE_P(
377373

378374
TEST_P(IntegrationGpuBatchnormBackwardActivationNdhwcBfp16, Correctness)
379375
{
380-
runGraphTest(batchnorm::getToleranceBackward<hip_bfloat16>(), TensorLayout::NDHWC);
376+
runGraphTest(batchnorm::getToleranceBackward<bfloat16>(), TensorLayout::NDHWC);
381377
}
382378

383379
INSTANTIATE_TEST_SUITE_P(

‎dnn-providers/miopen-provider/integration_tests/IntegrationGpuBatchnormForwardInference.cpp‎

Lines changed: 9 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -29,7 +29,7 @@ class BatchnormForwardInference
2929
: public IntegrationGraphVerificationHarness<DataType, BatchnormTestCase>
3030
{
3131
protected:
32-
void runGraphTest(DataType tolerance, const TensorLayout& layout = TensorLayout::NCHW)
32+
void runGraphTest(float tolerance, const TensorLayout& layout = TensorLayout::NCHW)
3333
{
3434
const BatchnormTestCase& testCase = this->GetParam();
3535

@@ -86,29 +86,27 @@ class BatchnormForwardInference
8686

8787
using IntegrationGpuBatchnormForwardInferenceNchwFp32 = BatchnormForwardInference<float, float>;
8888

89-
using IntegrationGpuBatchnormForwardInferenceNchwBfp16
90-
= BatchnormForwardInference<hip_bfloat16, float>;
89+
using IntegrationGpuBatchnormForwardInferenceNchwBfp16 = BatchnormForwardInference<bfloat16, float>;
9190

9291
using IntegrationGpuBatchnormForwardInferenceNchwFp16 = BatchnormForwardInference<half, float>;
9392

9493
using IntegrationGpuBatchnormForwardInferenceNhwcFp32 = BatchnormForwardInference<float, float>;
9594

96-
using IntegrationGpuBatchnormForwardInferenceNhwcBfp16
97-
= BatchnormForwardInference<hip_bfloat16, float>;
95+
using IntegrationGpuBatchnormForwardInferenceNhwcBfp16 = BatchnormForwardInference<bfloat16, float>;
9896

9997
using IntegrationGpuBatchnormForwardInferenceNhwcFp16 = BatchnormForwardInference<half, float>;
10098

10199
using IntegrationGpuBatchnormForwardInferenceNcdhwFp32 = BatchnormForwardInference<float, float>;
102100

103101
using IntegrationGpuBatchnormForwardInferenceNcdhwBfp16
104-
= BatchnormForwardInference<hip_bfloat16, float>;
102+
= BatchnormForwardInference<bfloat16, float>;
105103

106104
using IntegrationGpuBatchnormForwardInferenceNcdhwFp16 = BatchnormForwardInference<half, float>;
107105

108106
using IntegrationGpuBatchnormForwardInferenceNdhwcFp32 = BatchnormForwardInference<float, float>;
109107

110108
using IntegrationGpuBatchnormForwardInferenceNdhwcBfp16
111-
= BatchnormForwardInference<hip_bfloat16, float>;
109+
= BatchnormForwardInference<bfloat16, float>;
112110

113111
using IntegrationGpuBatchnormForwardInferenceNdhwcFp16 = BatchnormForwardInference<half, float>;
114112

@@ -129,7 +127,7 @@ INSTANTIATE_TEST_SUITE_P(Full,
129127

130128
TEST_P(IntegrationGpuBatchnormForwardInferenceNchwBfp16, Correctness)
131129
{
132-
runGraphTest(batchnorm::getToleranceInference<hip_bfloat16>(), TensorLayout::NCHW);
130+
runGraphTest(batchnorm::getToleranceInference<bfloat16>(), TensorLayout::NCHW);
133131
}
134132

135133
INSTANTIATE_TEST_SUITE_P(Smoke,
@@ -168,7 +166,7 @@ INSTANTIATE_TEST_SUITE_P(Full,
168166

169167
TEST_P(IntegrationGpuBatchnormForwardInferenceNhwcBfp16, Correctness)
170168
{
171-
runGraphTest(batchnorm::getToleranceInference<hip_bfloat16>(), TensorLayout::NHWC);
169+
runGraphTest(batchnorm::getToleranceInference<bfloat16>(), TensorLayout::NHWC);
172170
}
173171

174172
INSTANTIATE_TEST_SUITE_P(Smoke,
@@ -203,7 +201,7 @@ INSTANTIATE_TEST_SUITE_P(Smoke,
203201

204202
TEST_P(IntegrationGpuBatchnormForwardInferenceNcdhwBfp16, Correctness)
205203
{
206-
runGraphTest(batchnorm::getToleranceInference<hip_bfloat16>(), TensorLayout::NCDHW);
204+
runGraphTest(batchnorm::getToleranceInference<bfloat16>(), TensorLayout::NCDHW);
207205
}
208206

209207
INSTANTIATE_TEST_SUITE_P(Smoke,
@@ -230,7 +228,7 @@ INSTANTIATE_TEST_SUITE_P(Smoke,
230228

231229
TEST_P(IntegrationGpuBatchnormForwardInferenceNdhwcBfp16, Correctness)
232230
{
233-
runGraphTest(batchnorm::getToleranceInference<hip_bfloat16>(), TensorLayout::NDHWC);
231+
runGraphTest(batchnorm::getToleranceInference<bfloat16>(), TensorLayout::NDHWC);
234232
}
235233

236234
INSTANTIATE_TEST_SUITE_P(Smoke,

0 commit comments

Comments
 (0)