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

Disable lda for swizzle_a in HipBlaslt #1699

Open
wants to merge 1 commit into
base: develop
Choose a base branch
from
Open

Conversation

mengzcai
Copy link
Contributor

Disable lda for swizzle_a in HipBlaslt

amd-jnovotny
amd-jnovotny previously approved these changes Feb 24, 2025
Copy link
Contributor

@amd-jnovotny amd-jnovotny left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Docs component OK.

Copy link
Collaborator

@jichangjichang jichangjichang left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The problem is that user may set another lda into swizzleA kernel through hipblaslt API, which may cause the result incorrect. I didn't see you fix that.

}

if(swizzleA && matA->ld != 0)
log_hints(__func__,
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

where did you clean or ignore lda?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks. The lda parameter is ignored (set to k) in both the initial problem and valid_args.

@mengzcai mengzcai force-pushed the lda branch 2 times, most recently from 619c9ce to a14872e Compare February 24, 2025 11:29
tensilelite

validateMatmulSwizzleArgs() lda
{
//support TN for swizzle
if(matmul_descr->op_A != HIPBLAS_OP_T && matmul_descr->op_B != HIPBLAS_OP_N)
{
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

swizzleA only need matmul_descr->op_A == HIPBLAS_OP_T
swizzleB only needs matmul_descr->op_B == HIPBLAS_OP_N


// matrix A
int64_t num_rows_a = matA->m;
int64_t num_cols_a = matA->n;
int num_batches_a = matA->batch_count;
a_type = matA->type;
lda = matA->ld;
lda = (swizzleA) ? k : matA->ld;
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

should we done this in Tensile runtime but not hipblaslt runtime?

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Or done it in Tensile kernel

}
}

if(swizzleA && !isValidOrderForDatatype(a_type, matA->order))
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

seems that we need to support problem type predicate for hipblasltOrder

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

5 participants