diff --git a/example/02_gemm_bilinear/gemm_bilinear_wmma_fp16.cpp b/example/02_gemm_bilinear/gemm_bilinear_wmma_fp16.cpp index 18731e810e..b50925d711 100644 --- a/example/02_gemm_bilinear/gemm_bilinear_wmma_fp16.cpp +++ b/example/02_gemm_bilinear/gemm_bilinear_wmma_fp16.cpp @@ -21,7 +21,7 @@ struct AlphaBetaAdd { - AlphaBetaAdd(float alpha, float beta) : alpha_(alpha), beta_(beta){}; + AlphaBetaAdd(float alpha, float beta) : alpha_(alpha), beta_(beta) {}; template __host__ __device__ constexpr void operator()(E& e, const C& c, const D& d) const; diff --git a/example/02_gemm_bilinear/gemm_bilinear_wmma_int8.cpp b/example/02_gemm_bilinear/gemm_bilinear_wmma_int8.cpp index 87812369bd..a9eef9c6cb 100644 --- a/example/02_gemm_bilinear/gemm_bilinear_wmma_int8.cpp +++ b/example/02_gemm_bilinear/gemm_bilinear_wmma_int8.cpp @@ -21,7 +21,7 @@ struct AlphaBetaAdd { - AlphaBetaAdd(int alpha, int beta) : alpha_(alpha), beta_(beta){}; + AlphaBetaAdd(int alpha, int beta) : alpha_(alpha), beta_(beta) {}; template __host__ __device__ constexpr void operator()(E& e, const C& c, const D& d) const; diff --git a/example/02_gemm_bilinear/gemm_bilinear_xdl_fp16.cpp b/example/02_gemm_bilinear/gemm_bilinear_xdl_fp16.cpp index c3e6ef7d5d..aa39afe277 100644 --- a/example/02_gemm_bilinear/gemm_bilinear_xdl_fp16.cpp +++ b/example/02_gemm_bilinear/gemm_bilinear_xdl_fp16.cpp @@ -20,7 +20,7 @@ struct AlphaBetaAdd { - AlphaBetaAdd(float alpha, float beta) : alpha_(alpha), beta_(beta){}; + AlphaBetaAdd(float alpha, float beta) : alpha_(alpha), beta_(beta) {}; template __host__ __device__ constexpr void operator()(E& e, const C& c, const D& d) const; diff --git a/example/60_gemm_multi_ABD/gemm_multi_ABD_xdl_fp16.cpp b/example/60_gemm_multi_ABD/gemm_multi_ABD_xdl_fp16.cpp index 93034a8b70..9b218ed583 100644 --- a/example/60_gemm_multi_ABD/gemm_multi_ABD_xdl_fp16.cpp +++ b/example/60_gemm_multi_ABD/gemm_multi_ABD_xdl_fp16.cpp @@ -83,7 +83,7 @@ struct AddScale struct AlphaBetaAdd { - AlphaBetaAdd(float alpha, float beta) : alpha_(alpha), beta_(beta){}; + AlphaBetaAdd(float alpha, float beta) : alpha_(alpha), beta_(beta) {}; template __host__ __device__ constexpr void operator()(E& e, const C& c, const D& d) const; diff --git a/example/61_contraction_multi_ABD/contraction_multi_ABD_xdl_fp16.cpp b/example/61_contraction_multi_ABD/contraction_multi_ABD_xdl_fp16.cpp index e7c1d6f0be..b50e876384 100644 --- a/example/61_contraction_multi_ABD/contraction_multi_ABD_xdl_fp16.cpp +++ b/example/61_contraction_multi_ABD/contraction_multi_ABD_xdl_fp16.cpp @@ -42,7 +42,7 @@ static constexpr ck::index_t NumDimK = 2; struct AlphaBetaAdd { - AlphaBetaAdd(float alpha, float beta) : alpha_(alpha), beta_(beta){}; + AlphaBetaAdd(float alpha, float beta) : alpha_(alpha), beta_(beta) {}; template __host__ __device__ constexpr void operator()(E& e, const C& c, const D& d) const; diff --git a/include/ck/library/utility/host_tensor_generator.hpp b/include/ck/library/utility/host_tensor_generator.hpp index ab69412c15..bc376ffcdf 100644 --- a/include/ck/library/utility/host_tensor_generator.hpp +++ b/include/ck/library/utility/host_tensor_generator.hpp @@ -483,7 +483,7 @@ struct GeneratorTensor_4 std::normal_distribution distribution; GeneratorTensor_4(float mean, float stddev, unsigned int seed = 1) - : generator(seed), distribution(mean, stddev){}; + : generator(seed), distribution(mean, stddev) {}; template T operator()(Is...) @@ -501,7 +501,7 @@ struct GeneratorTensor_4 std::normal_distribution distribution; GeneratorTensor_4(float mean, float stddev, unsigned int seed = 1) - : generator(seed), distribution(mean, stddev){}; + : generator(seed), distribution(mean, stddev) {}; template ck::f4x2_pk_t operator()(Is...) @@ -520,7 +520,7 @@ struct GeneratorTensor_4 std::normal_distribution distribution; GeneratorTensor_4(float mean, float stddev, unsigned int seed = 1) - : generator(seed), distribution(mean, stddev){}; + : generator(seed), distribution(mean, stddev) {}; template ck::f6x32_pk_t operator()(Is...) @@ -542,7 +542,7 @@ struct GeneratorTensor_4 std::normal_distribution distribution; GeneratorTensor_4(float mean, float stddev, unsigned int seed = 1) - : generator(seed), distribution(mean, stddev){}; + : generator(seed), distribution(mean, stddev) {}; template ck::bf6x32_pk_t operator()(Is...) diff --git a/include/ck/tensor_operation/gpu/element/binary_element_wise_operation.hpp b/include/ck/tensor_operation/gpu/element/binary_element_wise_operation.hpp index d86f01e255..f326c4a28d 100644 --- a/include/ck/tensor_operation/gpu/element/binary_element_wise_operation.hpp +++ b/include/ck/tensor_operation/gpu/element/binary_element_wise_operation.hpp @@ -279,7 +279,7 @@ struct Subtract struct Bilinear { - Bilinear(float alpha = 1.f, float beta = 1.f) : alpha_(alpha), beta_(beta){}; + Bilinear(float alpha = 1.f, float beta = 1.f) : alpha_(alpha), beta_(beta) {}; template __host__ __device__ constexpr void operator()(Y&, const X0&, const X1&) const; @@ -354,7 +354,7 @@ struct Bilinear struct AddClamp { AddClamp(float floor = 0.f, float ceil = NumericLimits::Max()) - : floor_(floor), ceil_(ceil){}; + : floor_(floor), ceil_(ceil) {}; template __host__ __device__ constexpr void operator()(Y& y, const X0& x0, const X1& x1) const; diff --git a/include/ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp b/include/ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp index 4a87e8a277..80b8306a51 100644 --- a/include/ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp +++ b/include/ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp @@ -756,7 +756,7 @@ struct UnarySqrt struct Clamp { Clamp(float floor = 0.f, float ceil = NumericLimits::Max()) - : floor_(floor), ceil_(ceil){}; + : floor_(floor), ceil_(ceil) {}; template __host__ __device__ constexpr void operator()(Y& y, const X& x) const; @@ -1324,7 +1324,7 @@ struct Swish struct SoftRelu { - SoftRelu(float alpha = 1.f) : alpha_(alpha){}; + SoftRelu(float alpha = 1.f) : alpha_(alpha) {}; template __host__ __device__ void operator()(T& y, const T& x) const @@ -1353,7 +1353,7 @@ struct SoftRelu struct Power { Power(float alpha = 0.f, float beta = 1.f, float gamma = 2.f) - : alpha_(alpha), beta_(beta), gamma_(gamma){}; + : alpha_(alpha), beta_(beta), gamma_(gamma) {}; template __host__ __device__ void operator()(T& y, const T& x) const @@ -1386,7 +1386,7 @@ struct Power struct ClippedRelu { - ClippedRelu(float alpha = 0.f, float beta = 1.f) : alpha_(alpha), beta_(beta){}; + ClippedRelu(float alpha = 0.f, float beta = 1.f) : alpha_(alpha), beta_(beta) {}; template __host__ __device__ void operator()(T& y, const T& x) const @@ -1415,7 +1415,7 @@ struct ClippedRelu struct LeakyRelu { - LeakyRelu(float alpha = 0.01f) : alpha_(alpha){}; + LeakyRelu(float alpha = 0.01f) : alpha_(alpha) {}; template __host__ __device__ void operator()(T& y, const T& x) const @@ -1442,7 +1442,7 @@ struct LeakyRelu struct Elu { - Elu(float alpha = 1.f) : alpha_(alpha){}; + Elu(float alpha = 1.f) : alpha_(alpha) {}; template __host__ __device__ void operator()(T& y, const T& x) const @@ -1469,7 +1469,7 @@ struct Elu struct Logistic { - Logistic(float alpha = 1.f) : alpha_(alpha){}; + Logistic(float alpha = 1.f) : alpha_(alpha) {}; template __host__ __device__ void operator()(T& y, const T& x) const diff --git a/include/ck_tile/core/config.hpp b/include/ck_tile/core/config.hpp index 97c60cf062..5f83c97968 100644 --- a/include/ck_tile/core/config.hpp +++ b/include/ck_tile/core/config.hpp @@ -152,7 +152,7 @@ // buffer atomic add: floating point #ifndef __HIP_DEVICE_COMPILE__ // for host code #define CK_TILE_USE_AMD_BUFFER_ATOMIC_ADD_FLOAT 1 -#elif defined(__gfx9__) || defined(__gfx12__)// for GPU code +#elif defined(__gfx9__) || defined(__gfx12__) // for GPU code #define CK_TILE_USE_AMD_BUFFER_ATOMIC_ADD_FLOAT 1 #else // for GPU code #define CK_TILE_USE_AMD_BUFFER_ATOMIC_ADD_FLOAT 0 diff --git a/include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp b/include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp index 0e385901ed..8a97231030 100644 --- a/include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp +++ b/include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp @@ -1308,7 +1308,7 @@ struct Swish struct SoftRelu { - SoftRelu(float alpha = 1.f) : alpha_(alpha){}; + SoftRelu(float alpha = 1.f) : alpha_(alpha) {}; template CK_TILE_HOST_DEVICE void operator()(T& y, const T& x) const @@ -1327,7 +1327,7 @@ struct SoftRelu struct Power { Power(float alpha = 0.f, float beta = 1.f, float gamma = 2.f) - : alpha_(alpha), beta_(beta), gamma_(gamma){}; + : alpha_(alpha), beta_(beta), gamma_(gamma) {}; template CK_TILE_HOST_DEVICE void operator()(T& y, const T& x) const @@ -1349,7 +1349,7 @@ struct Power struct ClippedRelu { - ClippedRelu(float alpha = 0.f, float beta = 1.f) : alpha_(alpha), beta_(beta){}; + ClippedRelu(float alpha = 0.f, float beta = 1.f) : alpha_(alpha), beta_(beta) {}; template CK_TILE_HOST_DEVICE void operator()(T& y, const T& x) const @@ -1368,7 +1368,7 @@ struct ClippedRelu struct LeakyRelu { - LeakyRelu(float alpha = 0.01f) : alpha_(alpha){}; + LeakyRelu(float alpha = 0.01f) : alpha_(alpha) {}; template CK_TILE_HOST_DEVICE void operator()(T& y, const T& x) const @@ -1385,7 +1385,7 @@ struct LeakyRelu struct Elu { - Elu(float alpha = 1.f) : alpha_(alpha){}; + Elu(float alpha = 1.f) : alpha_(alpha) {}; template CK_TILE_HOST_DEVICE void operator()(T& y, const T& x) const @@ -1402,7 +1402,7 @@ struct Elu struct Logistic { - Logistic(float alpha = 1.f) : alpha_(alpha){}; + Logistic(float alpha = 1.f) : alpha_(alpha) {}; template CK_TILE_HOST_DEVICE void operator()(T& y, const T& x) const diff --git a/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma_impl.hpp b/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma_impl.hpp index 4e5b9adc53..13727d41b1 100644 --- a/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma_impl.hpp +++ b/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma_impl.hpp @@ -77,32 +77,14 @@ struct WarpGemmAttributeWmmaImpl }; using DeviceIp = remove_cvref_t; -using WarpGemmAttributeWmmaImpl_f32_16x16x16_f16_f16 = WarpGemmAttributeWmmaImpl>; +using WarpGemmAttributeWmmaImpl_f32_16x16x16_f16_f16 = + WarpGemmAttributeWmmaImpl>; -using WarpGemmAttributeWmmaImpl_f32_16x16x16_bf16_bf16 = WarpGemmAttributeWmmaImpl>; +using WarpGemmAttributeWmmaImpl_f32_16x16x16_bf16_bf16 = + WarpGemmAttributeWmmaImpl>; -using WarpGemmAttributeWmmaImpl_i32_16x16x16_i8_i8 = WarpGemmAttributeWmmaImpl>; +using WarpGemmAttributeWmmaImpl_i32_16x16x16_i8_i8 = + WarpGemmAttributeWmmaImpl>; using WarpGemmAttributeWmmaImpl_f32_16x16x16_f8_f8 = WarpGemmAttributeWmmaImpl>;