Skip to content

Commit bfde520

Browse files
tsenwangtsenwang
committed
add vector scale a and b
add a and b vector scale change change 2
1 parent d001ff3 commit bfde520

4 files changed

Lines changed: 457 additions & 1 deletion

File tree

‎clients/samples/26_gemm_scale_a_b_vector/sample_hipblaslt_gemm_scale_a_b_vector.cpp‎

Lines changed: 210 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,210 @@
1+
/*******************************************************************************
2+
*
3+
* MIT License
4+
*
5+
* Copyright (C) 2022-2024 Advanced Micro Devices, Inc.
6+
*
7+
* Permission is hereby granted, free of charge, to any person obtaining a copy
8+
* of this software and associated documentation files (the "Software"), to deal
9+
* in the Software without restriction, including without limitation the rights
10+
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
11+
* copies of the Software, and to permit persons to whom the Software is
12+
* furnished to do so, subject to the following conditions:
13+
*
14+
* The above copyright notice and this permission notice shall be included in
15+
* all copies or substantial portions of the Software.
16+
*
17+
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
18+
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
19+
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
20+
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
21+
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
22+
* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
23+
* SOFTWARE.
24+
*
25+
*******************************************************************************/
26+
#include "helper.h"
27+
#include <hip/hip_fp16.h>
28+
#include <hip/hip_runtime.h>
29+
#include <hip/library_types.h>
30+
#include <hipblaslt/hipblaslt.h>
31+
#include <iostream>
32+
#include <vector>
33+
34+
void simpleGemmScaleABVec(hipblasLtHandle_t handle,
35+
hipblasOperation_t trans_a,
36+
hipblasOperation_t trans_b,
37+
int64_t m,
38+
int64_t n,
39+
int64_t k,
40+
int64_t batch_count,
41+
float& alpha,
42+
float& beta,
43+
void* d_a,
44+
void* d_b,
45+
void* d_c,
46+
void* d_d,
47+
void* d_workspace,
48+
int64_t max_workspace_size,
49+
hipStream_t stream);
50+
51+
int main()
52+
{
53+
Runner<hipblaslt_f8_fnuz, hipblaslt_f8_fnuz, hipblasLtHalf, float, float> runner(
54+
128, 128, 128, 1, 1.f, 1.f, 32 * 1024 * 1024);
55+
56+
std::cout << "Running with Scale A and Scale B alternating between 0.5 and 2.0" << std::endl;
57+
58+
runner.run([&runner] {
59+
simpleGemmScaleABVec(runner.handle,
60+
HIPBLAS_OP_T,
61+
HIPBLAS_OP_N,
62+
runner.m,
63+
runner.n,
64+
runner.k,
65+
runner.batch_count,
66+
runner.alpha,
67+
runner.beta,
68+
runner.d_a,
69+
runner.d_b,
70+
runner.d_c,
71+
runner.d_d,
72+
runner.d_workspace,
73+
runner.max_workspace_size,
74+
runner.stream);
75+
});
76+
77+
return 0;
78+
}
79+
80+
void simpleGemmScaleABVec(hipblasLtHandle_t handle,
81+
hipblasOperation_t trans_a,
82+
hipblasOperation_t trans_b,
83+
int64_t m,
84+
int64_t n,
85+
int64_t k,
86+
int64_t batch_count,
87+
float& alpha,
88+
float& beta,
89+
void* d_a,
90+
void* d_b,
91+
void* d_c,
92+
void* d_d,
93+
void* d_workspace,
94+
int64_t max_workspace_size,
95+
hipStream_t stream)
96+
{
97+
std::vector<float> h_scale_a(128);
98+
std::vector<float> h_scale_b(128);
99+
100+
// Initialize scale_a and scale_b to alternate between 0.5 and 2.0
101+
for (int i = 0; i < 128; ++i) {
102+
if (i % 2 == 0) {
103+
h_scale_a[i] = 0.5f;
104+
h_scale_b[i] = 0.5f;
105+
} else {
106+
h_scale_a[i] = 2.0f;
107+
h_scale_b[i] = 2.0f;
108+
}
109+
}
110+
111+
float* d_scale_a;
112+
float* d_scale_b;
113+
114+
// Allocate device memory for scale vectors
115+
CHECK_HIP_ERROR(hipMalloc(&d_scale_a, h_scale_a.size() * sizeof(float)));
116+
CHECK_HIP_ERROR(hipMalloc(&d_scale_b, h_scale_b.size() * sizeof(float)));
117+
118+
// Copy scale vectors from host to device
119+
CHECK_HIP_ERROR(hipMemcpyAsync(d_scale_a, h_scale_a.data(), h_scale_a.size() * sizeof(float), hipMemcpyHostToDevice, stream));
120+
CHECK_HIP_ERROR(hipMemcpyAsync(d_scale_b, h_scale_b.data(), h_scale_b.size() * sizeof(float), hipMemcpyHostToDevice, stream));
121+
122+
hipblasLtMatrixLayout_t matA, matB, matC, matD;
123+
CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutCreate(&matA, HIP_R_8F_E4M3_FNUZ, m, k, m));
124+
CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutCreate(&matB, HIP_R_8F_E4M3_FNUZ, k, n, k));
125+
CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutCreate(&matC, HIP_R_16BF, m, n, m));
126+
CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutCreate(&matD, HIP_R_16BF, m, n, m));
127+
128+
hipblasLtMatmulDesc_t matmul;
129+
CHECK_HIPBLASLT_ERROR(hipblasLtMatmulDescCreate(&matmul, HIPBLAS_COMPUTE_32F, HIP_R_32F));
130+
CHECK_HIPBLASLT_ERROR(hipblasLtMatmulDescSetAttribute(
131+
matmul, HIPBLASLT_MATMUL_DESC_TRANSA, &trans_a, sizeof(int32_t)));
132+
CHECK_HIPBLASLT_ERROR(hipblasLtMatmulDescSetAttribute(
133+
matmul, HIPBLASLT_MATMUL_DESC_TRANSB, &trans_b, sizeof(int32_t)));
134+
135+
// Set A and B matrix scale vectors
136+
CHECK_HIPBLASLT_ERROR(hipblasLtMatmulDescSetAttribute(
137+
matmul, HIPBLASLT_MATMUL_DESC_A_SCALE_POINTER_VEC_EXT, &d_scale_a, sizeof(float*)));
138+
CHECK_HIPBLASLT_ERROR(hipblasLtMatmulDescSetAttribute(
139+
matmul, HIPBLASLT_MATMUL_DESC_B_SCALE_POINTER_VEC_EXT, &d_scale_b, sizeof(float*)));
140+
141+
hipblasLtMatmulPreference_t pref;
142+
CHECK_HIPBLASLT_ERROR(hipblasLtMatmulPreferenceCreate(&pref));
143+
CHECK_HIPBLASLT_ERROR(
144+
hipblasLtMatmulPreferenceSetAttribute(pref,
145+
HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES,
146+
&max_workspace_size,
147+
sizeof(max_workspace_size)));
148+
149+
const int request_solutions = 5;
150+
hipblasLtMatmulHeuristicResult_t heuristicResult[request_solutions];
151+
int returnedAlgoCount = 0;
152+
CHECK_HIPBLASLT_ERROR(hipblasLtMatmulAlgoGetHeuristic(handle,
153+
matmul,
154+
matA,
155+
matB,
156+
matC,
157+
matD,
158+
pref,
159+
request_solutions,
160+
heuristicResult,
161+
&returnedAlgoCount));
162+
163+
if(returnedAlgoCount == 0)
164+
{
165+
std::cerr << "No valid solution found!" << std::endl;
166+
CHECK_HIP_ERROR(hipFree(d_scale_a));
167+
CHECK_HIP_ERROR(hipFree(d_scale_b));
168+
CHECK_HIPBLASLT_ERROR(hipblasLtMatmulPreferenceDestroy(pref));
169+
CHECK_HIPBLASLT_ERROR(hipblasLtMatmulDescDestroy(matmul));
170+
CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutDestroy(matA));
171+
CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutDestroy(matB));
172+
CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutDestroy(matC));
173+
CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutDestroy(matD));
174+
return;
175+
}
176+
177+
uint64_t workspace_size = max_workspace_size;
178+
for(int i = 0; i < returnedAlgoCount; i++)
179+
workspace_size = std::max(workspace_size, heuristicResult[i].workspaceSize);
180+
181+
// Perform matrix multiplication
182+
CHECK_HIPBLASLT_ERROR(hipblasLtMatmul(handle,
183+
matmul,
184+
&alpha,
185+
d_a,
186+
matA,
187+
d_b,
188+
matB,
189+
&beta,
190+
d_c,
191+
matC,
192+
d_d,
193+
matD,
194+
&heuristicResult[0].algo,
195+
d_workspace,
196+
workspace_size,
197+
stream));
198+
199+
// Clean up resources
200+
CHECK_HIP_ERROR(hipFree(d_scale_a));
201+
CHECK_HIP_ERROR(hipFree(d_scale_b));
202+
CHECK_HIPBLASLT_ERROR(hipblasLtMatmulPreferenceDestroy(pref));
203+
CHECK_HIPBLASLT_ERROR(hipblasLtMatmulDescDestroy(matmul));
204+
CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutDestroy(matA));
205+
CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutDestroy(matB));
206+
CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutDestroy(matC));
207+
CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutDestroy(matD));
208+
209+
std::cout << "Matrix multiplication completed successfully." << std::endl;
210+
}

‎clients/samples/27_gemm_scale_a_b_vector_ext/sample_hipblaslt_gemm_scale_a_b_vector_ext.cpp‎

Lines changed: 177 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,177 @@
1+
/*******************************************************************************
2+
*
3+
* MIT License
4+
*
5+
* Copyright (C) 2022-2024 Advanced Micro Devices, Inc.
6+
*
7+
* Permission is hereby granted, free of charge, to any person obtaining a copy
8+
* of this software and associated documentation files (the "Software"), to deal
9+
* in the Software without restriction, including without limitation the rights
10+
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
11+
* copies of the Software, and to permit persons to whom the Software is
12+
* furnished to do so, subject to the following conditions:
13+
*
14+
* The above copyright notice and this permission notice shall be included in
15+
* all copies or substantial portions of the Software.
16+
*
17+
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
18+
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
19+
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
20+
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
21+
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
22+
* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
23+
* SOFTWARE.
24+
*
25+
*******************************************************************************/
26+
27+
#include <hip/hip_runtime.h>
28+
#include <hipblaslt/hipblaslt-ext.hpp>
29+
#include <iostream>
30+
31+
#include "helper.h"
32+
33+
void simpleGemmScaleABVecExt(hipblasLtHandle_t handle,
34+
hipblasOperation_t trans_a,
35+
hipblasOperation_t trans_b,
36+
int64_t m,
37+
int64_t n,
38+
int64_t k,
39+
int64_t batch_count,
40+
float& alpha,
41+
float& beta,
42+
void* d_a,
43+
void* d_b,
44+
void* d_c,
45+
void* d_d,
46+
void* d_workspace,
47+
int64_t max_workspace_size,
48+
hipStream_t stream);
49+
50+
int main()
51+
{
52+
// This is an example using hipblaslt extension API: ScaleA & ScaleB
53+
Runner<hipblaslt_f8_fnuz, hipblaslt_f8_fnuz, hipblasLtHalf, float, float> runner(
54+
128, 128, 128, 1, 1.f, 1.f, 32 * 128 * 128);
55+
56+
std::cout << "Running with Scale A and Scale B alternating between 0.5 and 2.0" << std::endl;
57+
58+
runner.run([&runner] {
59+
simpleGemmScaleABVecExt(runner.handle,
60+
HIPBLAS_OP_T,
61+
HIPBLAS_OP_N,
62+
runner.m,
63+
runner.n,
64+
runner.k,
65+
runner.batch_count,
66+
runner.alpha,
67+
runner.beta,
68+
runner.d_a,
69+
runner.d_b,
70+
runner.d_c,
71+
runner.d_d,
72+
runner.d_workspace,
73+
runner.max_workspace_size,
74+
runner.stream);
75+
});
76+
77+
return 0;
78+
}
79+
80+
void simpleGemmScaleABVecExt(hipblasLtHandle_t handle,
81+
hipblasOperation_t trans_a,
82+
hipblasOperation_t trans_b,
83+
int64_t m,
84+
int64_t n,
85+
int64_t k,
86+
int64_t batch_count,
87+
float& alpha,
88+
float& beta,
89+
void* d_a,
90+
void* d_b,
91+
void* d_c,
92+
void* d_d,
93+
void* d_workspace,
94+
int64_t max_workspace_size,
95+
hipStream_t stream)
96+
{
97+
hipblaslt_ext::GemmPreferenceV2 gemmPref;
98+
gemmPref.setMaxWorkspaceBytes(max_workspace_size);
99+
hipblaslt_ext::Gemm gemm(handle,
100+
trans_a,
101+
trans_b,
102+
HIP_R_8F_E4M3_FNUZ,
103+
HIP_R_8F_E4M3_FNUZ,
104+
HIP_R_16BF,
105+
HIP_R_16BF,
106+
HIPBLAS_COMPUTE_32F);
107+
108+
hipblaslt_ext::GemmEpilogueV2
109+
epilogue; // No action needed, default is HIPBLASLT_EPILOGUE_DEFAULT. (Gemm only)
110+
hipblaslt_ext::GemmInputsV2 inputs;
111+
112+
std::vector<float> h_scale_a(128);
113+
std::vector<float> h_scale_b(128);
114+
115+
// Initialize scale_a and scale_b to alternate between 0.5 and 2.0
116+
for (int i = 0; i < 128; ++i) {
117+
if (i % 2 == 0) {
118+
h_scale_a[i] = 0.5f;
119+
h_scale_b[i] = 0.5f;
120+
} else {
121+
h_scale_a[i] = 2.0f;
122+
h_scale_b[i] = 2.0f;
123+
}
124+
}
125+
126+
127+
float* d_scale_a;
128+
float* d_scale_b;
129+
130+
// Allocate device memory for scale vectors
131+
CHECK_HIP_ERROR(hipMalloc(&d_scale_a, h_scale_a.size() * sizeof(float)));
132+
CHECK_HIP_ERROR(hipMalloc(&d_scale_b, h_scale_b.size() * sizeof(float)));
133+
134+
// Copy scale vectors from host to device
135+
CHECK_HIP_ERROR(hipMemcpyAsync(d_scale_a, h_scale_a.data(), h_scale_a.size() * sizeof(float), hipMemcpyHostToDevice, stream));
136+
CHECK_HIP_ERROR(hipMemcpyAsync(d_scale_b, h_scale_b.data(), h_scale_b.size() * sizeof(float), hipMemcpyHostToDevice, stream));
137+
138+
inputs.setA(d_a);
139+
inputs.setB(d_b);
140+
inputs.setC(d_c);
141+
inputs.setD(d_d);
142+
inputs.setAlpha(&alpha);
143+
inputs.setBeta(&beta);
144+
epilogue.setScalingAType(1);
145+
epilogue.setScalingBType(1);
146+
inputs.setScaleA(d_scale_a);
147+
inputs.setScaleB(d_scale_b);
148+
gemm.setProblem(m, n, k, batch_count, epilogue, inputs);
149+
150+
const int request_solutions = 1;
151+
std::vector<hipblasLtMatmulHeuristicResult_t> heuristicResult;
152+
CHECK_HIPBLASLT_ERROR(gemm.algoGetHeuristic(request_solutions, gemmPref, heuristicResult));
153+
154+
if(heuristicResult.empty())
155+
{
156+
CHECK_HIP_ERROR(hipFree(d_scale_a));
157+
CHECK_HIP_ERROR(hipFree(d_scale_b));
158+
std::cerr << "No valid solution found!" << std::endl;
159+
return;
160+
}
161+
162+
// In this sample, the workspace is already allocated with max_workspace_size
163+
// If not, calculate the needed workspace_size and allocate d_workspace here
164+
// uint64_t workspace_size = 0;
165+
// for(int i = 0; i < returnedAlgoCount; i++)
166+
// workspace_size = max(workspace_size, heuristicResult[i].workspaceSize);
167+
// CHECK_HIP_ERRORhipMalloc(&d_workspace, workspace_size));
168+
169+
// Make sure to initialize every time when algo changes
170+
CHECK_HIPBLASLT_ERROR(gemm.initialize(heuristicResult[0].algo, d_workspace));
171+
CHECK_HIPBLASLT_ERROR(gemm.run(stream));
172+
173+
CHECK_HIP_ERROR(hipFree(d_scale_a));
174+
CHECK_HIP_ERROR(hipFree(d_scale_b));
175+
std::cout << "Matrix multiplication completed successfully." << std::endl;
176+
return;
177+
}

0 commit comments

Comments
 (0)