Skip to content

Commit

Permalink
Merge remote-tracking branch 'origin/develop' into cderb/gtest_package
Browse files Browse the repository at this point in the history
  • Loading branch information
@CAHEK7
CAHEK7 committed Dec 20, 2023
2 parents a13243f + 45991db commit fed0b7c
Show file tree
Hide file tree
Showing 15 changed files with 231 additions and 145 deletions.
24 changes: 21 additions & 3 deletions Jenkinsfile
View file
Original file line number Diff line number Diff line change
Expand Up @@ -406,7 +406,7 @@ def CheckPerfDbValid(Map conf=[:]){
/// * The default compiler is usually not specified.
/// BuildType := { Release* | Debug | Install } [ BuildTypeModifier ]
/// * BuildTypeModifier := { NOCOMGR | Embedded | Static | Normal-Find | Fast-Find
/// CK | MLIR | Tensile | Tensile-Latest | Package | ... }
/// CK | NOMLIR | Tensile | Tensile-Latest | Package | ... }
/// TestSet := { All | Smoke* | Performance Dataset } [ Codecov ]
/// * "All" corresponds to "cmake -DMIOPEN_TEST_ALL=On".
/// * "Smoke" (-DMIOPEN_TEST_ALL=Off) is the default and usually not specified.
Expand Down Expand Up @@ -466,7 +466,7 @@ pipeline {
description: "")
booleanParam(
name: "TARGET_GFX908",
defaultValue: true,
defaultValue: env.BRANCH_NAME == env.NIGHTLY_BRANCH ? true : false,
description: "")
booleanParam(
name: "TARGET_GFX90A",
Expand Down Expand Up @@ -530,8 +530,9 @@ pipeline {
Bf16_flags = " -DMIOPEN_TEST_BFLOAT16=On"
Int8_flags = " -DMIOPEN_TEST_INT8=On"
Full_test = " -DMIOPEN_TEST_ALL=On"
Smoke_targets = "check MIOpenDriver"
Smoke_targets = " check MIOpenDriver"
NOCOMGR_flags = " -DMIOPEN_USE_COMGR=Off"
NOMLIR_flags = " -DMIOPEN_USE_MLIR=Off"
}
//triggers{
//
Expand Down Expand Up @@ -711,6 +712,23 @@ pipeline {
buildHipClangJobAndReboot( build_type: 'debug', setup_flags: NOCOMGR_flags, build_cmd: NOCOMGR_build_cmd, test_flags: ' --verbose ')
}
}
stage('Fp32 Hip Debug NOMLIR AnyGPU') {
when {
beforeAgent true
expression { params.TARGET_VEGA20 || params.TARGET_VEGA10 || params.TARGET_GFX908 || params.TARGET_GFX90A }
}
options {
retry(2)
}
agent{ label rocmnode("vega || gfx908 || gfx90a") }
environment{
// Can be removed altogether with when WORKAROUND_SWDEV_290754.
NOMLIR_build_cmd = "CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 MIOPEN_LOG_LEVEL=5 make -j\$(nproc) check"
}
steps{
buildHipClangJobAndReboot( build_type: 'debug', setup_flags: NOMLIR_flags, build_cmd: NOMLIR_build_cmd, test_flags: ' --verbose ')
}
}
stage('Fp32 Hip Debug Embedded Vega20') {
when {
beforeAgent true
Expand Down
36 changes: 18 additions & 18 deletions src/gemm_v2.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -142,11 +142,11 @@ template <typename T>
rocblas_status miopen_rocblas_gemm_ex3(const miopen::Handle& handle,
const miopen::GemmDescriptor& gemm_desc,
ConstData_t A,
int a_offset,
std::size_t a_offset,
ConstData_t B,
int b_offset,
std::size_t b_offset,
Data_t C,
int c_offset)
std::size_t c_offset)
{
rocblas_status rb_status =
rocblas_status::rocblas_status_internal_error; // cppcheck-suppress redundantInitialization
Expand Down Expand Up @@ -231,11 +231,11 @@ template <typename T>
rocblas_status miopen_rocblas_gemm_strided_batched_ex3(const miopen::Handle& handle,
const miopen::GemmDescriptor& gemm_desc,
ConstData_t A,
int a_offset,
std::size_t a_offset,
ConstData_t B,
int b_offset,
std::size_t b_offset,
Data_t C,
int c_offset)
std::size_t c_offset)
{
rocblas_status rb_status = rocblas_status::rocblas_status_internal_error;
// Until there is a batched counter part to the ex3 rocBlas call we need to iterate over the
Expand Down Expand Up @@ -366,11 +366,11 @@ static GemmBackend_t enforce_gemm_backend(miopenDataType_t data_type,
miopenStatus_t CallGemmTimeMeasure(const Handle& handle,
GemmDescriptor gemm_desc,
ConstData_t A,
int a_offset,
std::size_t a_offset,
ConstData_t B,
int b_offset,
std::size_t b_offset,
Data_t C,
int c_offset,
std::size_t c_offset,
bool time_precision,
CallGemmType_t call_gemm_type,
GemmBackend_t gemm_backend)
Expand Down Expand Up @@ -415,11 +415,11 @@ miopenStatus_t CallGemmTimeMeasure(const Handle& handle,
miopenStatus_t CallGemm(const Handle& handle,
GemmDescriptor gemm_desc,
ConstData_t A,
int a_offset,
std::size_t a_offset,
ConstData_t B,
int b_offset,
std::size_t b_offset,
Data_t C,
int c_offset,
std::size_t c_offset,
GemmBackend_t gemm_backend)
{
MIOPEN_LOG_I2("gemm_desc: " << gemm_desc);
Expand Down Expand Up @@ -669,11 +669,11 @@ miopenStatus_t CallGemm(const Handle& handle,
miopenStatus_t CallGemmStridedBatched(const Handle& handle,
GemmDescriptor gemm_desc,
ConstData_t A,
int a_offset,
std::size_t a_offset,
ConstData_t B,
int b_offset,
std::size_t b_offset,
Data_t C,
int c_offset,
std::size_t c_offset,
GemmBackend_t gemm_backend)
{
MIOPEN_LOG_I2("gemm_desc: " << gemm_desc);
Expand Down Expand Up @@ -944,11 +944,11 @@ miopenStatus_t CallGemmStridedBatched(const Handle& handle,
miopenStatus_t CallGemmStridedBatchedSequential(const Handle& handle,
GemmDescriptor gemm_desc,
ConstData_t A,
int a_offset,
std::size_t a_offset,
ConstData_t B,
int b_offset,
std::size_t b_offset,
Data_t C,
int c_offset,
std::size_t c_offset,
GemmBackend_t gemm_backend)
{
MIOPEN_LOG_I2("gemm_desc: " << gemm_desc);
Expand Down
4 changes: 4 additions & 0 deletions src/include/miopen/conv/asm_implicit_gemm.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -39,6 +39,10 @@
/// https://github.com/ROCm/MIOpen/issues/1979
#define WORKAROUND_ISSUE_1979 1

/// W/A for issue 2624: asm igemm fwd error when c <=4 and dilation_y > 1
/// https://github.com/ROCm/MIOpen/issues/2624
#define WORKAROUND_ISSUE_2624 1

namespace miopen {

namespace solver {
Expand Down
24 changes: 12 additions & 12 deletions src/include/miopen/gemm_v2.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -124,44 +124,44 @@ struct GemmDescriptor
miopenStatus_t CallGemmTimeMeasure(const Handle& handle,
GemmDescriptor gemm_desc,
ConstData_t A,
int a_offset,
std::size_t a_offset,
ConstData_t B,
int b_offset,
std::size_t b_offset,
Data_t C,
int c_offset,
std::size_t c_offset,
bool time_precision,
CallGemmType_t call_gemm_type,
GemmBackend_t gemm_backend = GemmBackend_t::rocblas);

miopenStatus_t CallGemm(const Handle& handle,
GemmDescriptor gemm_desc,
ConstData_t A,
int a_offset,
std::size_t a_offset,
ConstData_t B,
int b_offset,
std::size_t b_offset,
Data_t C,
int c_offset,
std::size_t c_offset,
GemmBackend_t gemm_backend = GemmBackend_t::rocblas);

miopenStatus_t CallGemmStridedBatched(const Handle& handle,
GemmDescriptor gemm_desc,
ConstData_t A,
int a_offset,
std::size_t a_offset,
ConstData_t B,
int b_offset,
std::size_t b_offset,
Data_t C,
int c_offset,
std::size_t c_offset,
GemmBackend_t gemm_backend = GemmBackend_t::rocblas);

miopenStatus_t
CallGemmStridedBatchedSequential(const Handle& handle,
GemmDescriptor gemm_desc,
ConstData_t A,
int a_offset,
std::size_t a_offset,
ConstData_t B,
int b_offset,
std::size_t b_offset,
Data_t C,
int c_offset,
std::size_t c_offset,
GemmBackend_t gemm_backend = GemmBackend_t::rocblas);

// GEMM parameters for Convolution (using Im2Col) Fwd
Expand Down
12 changes: 6 additions & 6 deletions src/include/miopen/util.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -74,8 +74,8 @@ float transpose_NCHW2CNHW(const Handle& handle,
int w_out,
ConstData_t in,
Data_t out,
int in_offset,
int out_offset,
std::size_t in_offset,
std::size_t out_offset,
int h_stride,
int w_stride,
miopenDataType_t type);
Expand All @@ -89,8 +89,8 @@ float transpose_CNHW2NCHW(const Handle& handle,
int w_in,
ConstData_t in,
Data_t out,
int in_offset,
int out_offset,
std::size_t in_offset,
std::size_t out_offset,
int h_stride,
int w_stride,
miopenDataType_t type);
Expand All @@ -108,8 +108,8 @@ float transpose_NCHW2Vec(const Handle& handle,
float transpose_packed_MN2NM(const Handle& handle,
int m,
int n,
int in_offset,
int out_offset,
std::size_t in_offset,
std::size_t out_offset,
ConstData_t in,
Data_t out,
miopenDataType_t type);
Expand Down
Loading

0 comments on commit fed0b7c

Please sign in to comment.