Skip to content

Commit

Permalink
resolve PREfast code scanning
Browse files Browse the repository at this point in the history
  • Loading branch information
fs-eire committed Oct 26, 2024
1 parent 93be3f4 commit fa67aa2
Show file tree
Hide file tree
Showing 5 changed files with 8 additions and 8 deletions.
2 changes: 1 addition & 1 deletion onnxruntime/contrib_ops/webgpu/bert/multihead_attention.cc
Original file line number Diff line number Diff line change
Expand Up @@ -350,7 +350,7 @@ Status ComputeVxAttentionScore(onnxruntime::webgpu::ComputeContext& context, int
int total_sequence_length) {
const bool feed_past_value = present_value != nullptr && past_value != nullptr && past_value->SizeInBytes() > 0;
const bool has_present_value = output_count > 1 && past_value != nullptr;
const int tile_size = 12;
constexpr int tile_size = 12;

VxAttentionScoreProgram program{"VxAttentionScore", feed_past_value, has_present_value, tile_size};
program.AddInputs({{probs, ProgramTensorMetadataDependency::TypeAndRank},
Expand Down
8 changes: 4 additions & 4 deletions onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits.cc
Original file line number Diff line number Diff line change
Expand Up @@ -61,7 +61,7 @@ Status MatMulNBitsProgram::GenerateShaderCode(ShaderHelper& shader) const {
const uint32_t workgroup_size = WorkgroupSizeX() * WorkgroupSizeY();
const uint32_t tile_size = WorkgroupSizeX() * components_b_ * 8; // each uint32 has 8 data.
const uint32_t a_length_per_tile = tile_size / a.NumComponents();
const uint32_t block_size = 32;
constexpr uint32_t block_size = 32;
const uint32_t blocks_per_tile = tile_size / block_size;
shader.AdditionalImplementation() << "var<workgroup> sub_a: array<input_a_value_t, " << a_length_per_tile << ">;\n"
<< "var<workgroup> inter_results: array<array<output_value_t, " << WorkgroupSizeX() << ">, " << WorkgroupSizeY() << ">;\n";
Expand Down Expand Up @@ -344,7 +344,7 @@ Status MatMulNBits::ComputeInternal(onnxruntime::webgpu::ComputeContext& context
const uint32_t N = SafeInt<uint32_t>(helper.N());
const uint32_t K = SafeInt<uint32_t>(helper.K());
const uint32_t block_size = SafeInt<uint32_t>(block_size_);
const uint32_t nbits = 4;
constexpr uint32_t nbits = 4;

const uint32_t n_blocks_per_col = (K + block_size - 1) / block_size;
const uint32_t blob_size = (block_size / 8) * nbits;
Expand All @@ -357,12 +357,12 @@ Status MatMulNBits::ComputeInternal(onnxruntime::webgpu::ComputeContext& context
const bool has_zero_points = zero_points != nullptr;
// TODO: Support output_number > 1. Some cases are failed when output_number > 1.
// const uint32_t output_number = M > 1 && (N / components) % 2 == 0 ? 2 : 1;
const uint32_t output_number = 1;
constexpr uint32_t output_number = 1;
MatMulNBitsProgram program{output_number, SafeInt<int>(components_b), has_zero_points, use_block32};

if (use_block32) {
components = 1;
const uint32_t workgroup_size = 128;
constexpr uint32_t workgroup_size = 128;
const uint32_t workgroup_y = N % 8 == 0 ? 8 : N % 4 == 0 ? 4
: 1;
const uint32_t workgroup_x = workgroup_size / workgroup_y;
Expand Down
2 changes: 1 addition & 1 deletion onnxruntime/core/providers/webgpu/buffer_manager.cc
Original file line number Diff line number Diff line change
Expand Up @@ -7,7 +7,7 @@
namespace onnxruntime {
namespace webgpu {

size_t NormalizeBufferSize(size_t size) {
constexpr size_t NormalizeBufferSize(size_t size) {
return (size + 15) / 16 * 16;
}

Expand Down
2 changes: 1 addition & 1 deletion onnxruntime/core/providers/webgpu/tensor/where.cc
Original file line number Diff line number Diff line change
Expand Up @@ -126,7 +126,7 @@ Status Where::ComputeInternal(ComputeContext& context) const {
TensorShape output_shape;
ORT_RETURN_IF_ERROR(ComputeOutputShape(cond_shape, x_shape, y_shape, output_shape));
auto* output_tensor = context.Output(0, output_shape);
const auto component = 4;
constexpr int component = 4;
uint32_t vec_size = gsl::narrow_cast<uint32_t>((output_shape.Size() + 3) / component);
const auto is_broadcast = !(x_shape == y_shape &&
y_shape == cond_shape);
Expand Down
2 changes: 1 addition & 1 deletion onnxruntime/core/providers/webgpu/webgpu_context.cc
Original file line number Diff line number Diff line change
Expand Up @@ -346,7 +346,7 @@ Status WebGpuContext::Run(ComputeContext& context, const ProgramBase& program) {

// Meet alignment of struct here: https://www.w3.org/TR/WGSL/#alignment-and-size. For simplicity, set
// max_alignment_of_field to 16 since the underlying buffer has been rounded up to 16.
const size_t max_alignment_of_field = 16;
constexpr size_t max_alignment_of_field = 16;
const size_t uniform_buffer_total_size = (current_offset + max_alignment_of_field - 1) / max_alignment_of_field * max_alignment_of_field;

WGPUBuffer uniform_buffer = nullptr;
Expand Down

0 comments on commit fa67aa2

Please sign in to comment.