[SYCLomatic] Enable migration of CUBLASLT: EPILOGUE_BIAS, EPILOGUE_GELU, EPILOGUE_GELU_AUX, EPILOGUE_GELU_AUX_BIAS #2460

abhilash1910 · 2024-11-06T14:18:32Z

zhiweij1 · 2024-11-07T00:05:00Z

clang/runtime/dpct-rt/include/dpct/blas_gemm_utils.hpp

+    } else if (compute_desc->_epilogue == epilogue_t::bias) {
+      matmul_ops.append_binary(::dnnl::algorithm::binary_add, bias_md);
+    } else if (compute_desc-> epilogue == epilogue::gelu_aux_bias) {
+      matmul_ops.append_eltwise(::dnnl::algorithm::eltwise_gelu_erf, 0.f, 0.f);


CUBLASLT_EPILOGUE_GELU_AUX_BIAS needs to output the bias result to CUBLASLT_MATMUL_DESC_EPILOGUE_AUX_POINTER. But here, your implementation does not.

Please write an E2E test to validate your implementation. I tried something similar before but encountered a runtime error or an incorrect result.

Ok. Please help to review.

zhiweij1

Please refine

zhiweij1 · 2024-11-19T00:55:48Z

clang/runtime/dpct-rt/include/dpct/blas_gemm_utils.hpp

+    } else if (compute_desc-> epilogue == epilogue::gelu_aux_bias) {
+      matmul_ops.append_eltwise(::dnnl::algorithm::eltwise_gelu_erf, 0.f, 0.f);
+      matmul_ops.append_binary(::dnnl::algorithm::binary_add, bias_md);
+      compute_desc->set_attribute(matmul_desc_t::attribute::epilogue_aux_pointer, bias_mem);


We need to copy the data in bias_mem to _epilogue_aux_pointer. Meanwhile, please also consider the data layout (row_major/col_major) and the leading dimension (CUBLASLT_MATMUL_DESC_EPILOGUE_AUX_LD).

And according to the cublaslt doc, CUBLASLT_MATMUL_DESC_BIAS_POINTER is also related to this epilogue. Please check the usage.

I directly added matrix memcpy to copy epilogue_aux_pointer contents from bias_mem pointer, using the epilogue_aux_ld is and new_ldc (bias ld). This usage is for matrix memcpy, and will try to see if tests are ok .
Please suggest in meantime.

zhiweij1

Please refine

zhiweij1 · 2024-11-20T03:20:13Z

Also could you please prepare an E2E test first? like https://github.com/oneapi-src/SYCLomatic-test/blob/SYCLomatic/features/feature_case/cublasLt/matmul.cu .
Make sure your test case (cuda code) can run pass on NV HW. Then migrate it to sycl. Compile the sycl code within your PR, make sure the result is same as the cuda code output.

zhiweij1 · 2024-11-28T07:50:40Z

clang/runtime/dpct-rt/include/dpct/blas_gemm_utils.hpp

+    } else if (compute_desc->_epilogue == epilogue_t::gelu_aux_bias) {
+      matmul_ops.append_eltwise(::dnnl::algorithm::eltwise_gelu_erf, 0.f, 0.f);
+      matmul_ops.append_binary(::dnnl::algorithm::binary_add, bias_md);
+      dpct::blas::matrix_mem_copy(matmul_desc_t::attribute::epilogue_aux_pointer, bias_mem,


matmul_desc_t::attribute::epilogue_aux_pointer is a enum, not a pointer. bias_mem is a pointer of dnnl::memory, not a pointer to data.

zhiweij1 · 2024-12-02T08:48:16Z

clang/runtime/dpct-rt/include/dpct/blas_gemm_utils.hpp

+      matmul_ops.append_eltwise(::dnnl::algorithm::eltwise_gelu_erf, 0.f, 0.f);
+    }  else if (compute_desc->_epilogue == epilogue_t::gelu_aux) {
+      matmul_ops.append_eltwise(::dnnl::algorithm::eltwise_gelu_erf, 0.f, 0.f);
+      dpct::blas::matrix_mem_copy(compute_desc->_epilogue_aux_pointer, new_c,


I forget one issue. This dpct::blas::matrix_mem_copy does not return sycl::event. So we cannot control the dependency chain when using non-USM.

Another issue is new_c is the bias, not the result of (A*B+C). I did not find there is a method to get that immediate result when using oneDNN post-op.
Maybe we need to use a separate oneDNN elewise primitive.

Yes correct, for the second point, I see it is raised in JIRA. Currently I could not find a way to extract immediate result of A*B+C for the aux memcpy.

For the first case, do I have to force use USM always then ? What do you suggest ?

zhiweij1 · 2024-12-02T08:48:32Z

clang/runtime/dpct-rt/include/dpct/blas_gemm_utils.hpp

+     if (compute_desc->_epilogue == epilogue_t::relu) {
+      matmul_ops.append_eltwise(::dnnl::algorithm::eltwise_relu, 0.f, 0.f);
+    }  else if (compute_desc->_epilogue == epilogue_t::gelu) {
+      matmul_ops.append_eltwise(::dnnl::algorithm::eltwise_gelu_erf, 0.f, 0.f);


Suggested change

matmul_ops.append_eltwise(::dnnl::algorithm::eltwise_gelu_erf, 0.f, 0.f);

matmul_ops.append_eltwise(::dnnl::algorithm::eltwise_gelu_tanh, 0.f, 0.f);

zhiweij1 · 2024-12-02T08:51:51Z

clang/runtime/dpct-rt/include/dpct/blas_gemm_utils.hpp

+                                  compute_desc->_epilogue_aux_ld, new_ldc, m, n,
+                                  sizeof(size_t) , dpct::device_to_device, queue);
+    } else if (compute_desc->_epilogue == epilogue_t::bias) {
+      matmul_ops.append_binary(::dnnl::algorithm::binary_add, bias_md);


bias_md represents the C in the D=A*B+C, not the extra bias in epilogue.

Need to create new bias descriptor. A*B+C+bias.?

abhilash1910 requested a review from a team as a code owner November 6, 2024 14:18

abhilash1910 requested review from tomflinda and AndyCHHuang November 6, 2024 14:18

zhiweij1 reviewed Nov 7, 2024

View reviewed changes

zhiweij1 requested changes Nov 7, 2024

View reviewed changes

abhilash1910 closed this Nov 18, 2024

abhilash1910 force-pushed the bias_au_gelu_epilogues branch from 0a657ab to c5b19af Compare November 18, 2024 03:32

abhilash1910 reopened this Nov 18, 2024

zhiweij1 reviewed Nov 19, 2024

View reviewed changes

zhiweij1 requested changes Nov 19, 2024

View reviewed changes

abhilash1910 closed this Nov 26, 2024

abhilash1910 force-pushed the bias_au_gelu_epilogues branch from 6b14fa2 to a2dc3f5 Compare November 26, 2024 07:35

rebase

eb5a2f8

abhilash1910 reopened this Nov 26, 2024

abhilash1910 and others added 5 commits November 26, 2024 03:56

fix

9e48680

Merge branch 'oneapi-src:SYCLomatic' into bias_au_gelu_epilogues

b830f22

fix

bba1b9a

fix test

33c0f33

add pointer to pointer copy

266ccec

zhiweij1 reviewed Nov 28, 2024

View reviewed changes

abhilash1910 added 6 commits November 28, 2024 06:56

add header

03deb90

fix the memcpy

ceca900

fix

9f92c61

add gelu

1eb5880

fix

d52bc6b

add gelu aux

c837886

abhilash1910 changed the title ~~[SYCLomatic] Enable migration of CUBLASLT_EPILOGUE_BIAS & CUBLASLT_EPILOGUE_GELU_AUX_BIAS~~ [SYCLomatic] Enable migration of CUBLASLT: EPILOGUE_BIAS, EPILOGUE_GELU, EPILOGUE_GELU_AUX, EPILOGUE_GELU_AUX_BIAS Nov 29, 2024

abhilash1910 mentioned this pull request Nov 29, 2024

[SYCLomatic] Enable migration of CUBLASLT_EPILOGUE_GELU & GELU_AUX #2436

Closed

fix

29c1411

abhilash1910 and others added 3 commits December 2, 2024 08:38

Merge branch 'oneapi-src:SYCLomatic' into bias_au_gelu_epilogues

2cbc0cf

fix syntax

9bfb234

Merge branch 'oneapi-src:SYCLomatic' into bias_au_gelu_epilogues

fc018b0

zhiweij1 reviewed Dec 2, 2024

View reviewed changes

add epilogue bias mem descriptor and refactor

77356a0

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[SYCLomatic] Enable migration of CUBLASLT: EPILOGUE_BIAS, EPILOGUE_GELU, EPILOGUE_GELU_AUX, EPILOGUE_GELU_AUX_BIAS #2460

[SYCLomatic] Enable migration of CUBLASLT: EPILOGUE_BIAS, EPILOGUE_GELU, EPILOGUE_GELU_AUX, EPILOGUE_GELU_AUX_BIAS #2460

abhilash1910 commented Nov 6, 2024

zhiweij1 Nov 7, 2024

abhilash1910 Nov 18, 2024

zhiweij1 left a comment

zhiweij1 Nov 19, 2024 •

edited

Loading

abhilash1910 Nov 20, 2024

zhiweij1 left a comment

zhiweij1 commented Nov 20, 2024

zhiweij1 Nov 28, 2024

zhiweij1 Dec 2, 2024

zhiweij1 Dec 2, 2024

abhilash1910 Dec 2, 2024

zhiweij1 Dec 2, 2024

zhiweij1 Dec 2, 2024

abhilash1910 Dec 2, 2024 •

edited

Loading

	matmul_ops.append_eltwise(::dnnl::algorithm::eltwise_gelu_erf, 0.f, 0.f);
	matmul_ops.append_eltwise(::dnnl::algorithm::eltwise_gelu_tanh, 0.f, 0.f);

[SYCLomatic] Enable migration of CUBLASLT: EPILOGUE_BIAS, EPILOGUE_GELU, EPILOGUE_GELU_AUX, EPILOGUE_GELU_AUX_BIAS #2460

Are you sure you want to change the base?

[SYCLomatic] Enable migration of CUBLASLT: EPILOGUE_BIAS, EPILOGUE_GELU, EPILOGUE_GELU_AUX, EPILOGUE_GELU_AUX_BIAS #2460

Conversation

abhilash1910 commented Nov 6, 2024

Choose a reason for hiding this comment

Choose a reason for hiding this comment

zhiweij1 left a comment

Choose a reason for hiding this comment

zhiweij1 Nov 19, 2024 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

zhiweij1 left a comment

Choose a reason for hiding this comment

zhiweij1 commented Nov 20, 2024

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

abhilash1910 Dec 2, 2024 • edited Loading

Choose a reason for hiding this comment

zhiweij1 Nov 19, 2024 •

edited

Loading

abhilash1910 Dec 2, 2024 •

edited

Loading