Skip to content

Fc fp16 #44558

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.

Sign up for GitHub

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

Merged
merged 5 commits into from
Jul 22, 2022
Merged

Fc fp16 #44558

Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
1eac657
(modified) fc support fp16
ming1753 Jul 22, 2022
437d665
__CUDA_ARCH__ version
ming1753 Jul 22, 2022
ffcfda3
merged
ming1753 Jul 22, 2022
a8582e0
delete half
ming1753 Jul 22, 2022
c747777
delete half
ming1753 Jul 22, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
95 changes: 1 addition & 94 deletions paddle/phi/kernels/funcs/fc_functor.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -36,14 +36,6 @@ struct FcTypeTraits<double> {
typedef double4 Type;
};

#if defined(PADDLE_WITH_CUDA)
#include <cuda_fp16.h>

template <>
struct FcTypeTraits<float16> {
typedef half2 Type;
};
#else
struct float16_4 {
float16 x, y, z, w;
};
Expand All @@ -52,7 +44,6 @@ template <>
struct FcTypeTraits<float16> {
typedef float16_4 Type;
};
#endif

template <typename T, bool DoRelu>
__global__ void bias_relu_v4(const int num, const T* bias, T* data, int K) {
Expand Down Expand Up @@ -126,95 +117,12 @@ void AddReluKernel(
}
}

#if defined(PADDLE_WITH_CUDA)
template <bool DoRelu>
__global__ void bias_relu_v2(const int num,
const half2* bias,
half2* data,
int K) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid < num) {
int bias_idx = tid % K;
const half2 bias_ptr = bias[bias_idx];
const half2 in_ptr = data[tid];
half2 packed_val = __hadd2(bias_ptr, in_ptr);
if (DoRelu) {
#if __CUDA_ARCH__ >= 800
packed_val = __hmax2(__half2(0, 0), packed_val);
#else
packed_val = __hmul2(__hgt2(__half2(0, 0), packed_val), packed_val);
#endif
}
data[tid] = packed_val;
}
}

template <bool DoRelu, int BlockDim>
__global__ void InplaceAddReluKernel(const int N,
const half* bias,
half* data) {
int offset = blockIdx.x * N;
for (int i = threadIdx.x; i < N; i += BlockDim) {
half temp;
#if defined(__HIPCC__) || __CUDA_ARCH__ >= 350
temp = __ldg(data + offset + i) + __ldg(bias + i);
#else
temp = data[offset + i] + bias[i];
#endif
if (DoRelu) {
#if __CUDA_ARCH__ >= 800
data[offset + i] = __hmax(0, temp);
#else
data[offset + i] = __hmul(__hgt(temp, 0), temp);
#endif
} else {
data[offset + i] = temp;
}
}
}

template <>
void AddReluKernel(cudaStream_t stream,
const int M,
const int N,
float16* Y,
const float16* B,
bool relu) {
if (N % 2 == 0) {
const int threads = 256;
const int num = M * N / 2;
const int blocks = (num + threads - 1) / threads;
typedef typename FcTypeTraits<float16>::Type trans_type;
auto* bias_ptr_v2 = reinterpret_cast<const trans_type*>(B);
auto* data_ptr_v2 = reinterpret_cast<trans_type*>(Y);
if (relu) {
bias_relu_v2<true><<<blocks, threads, 0, stream>>>(
num, bias_ptr_v2, data_ptr_v2, N / 2);
} else {
bias_relu_v2<false><<<blocks, threads, 0, stream>>>(
num, bias_ptr_v2, data_ptr_v2, N / 2);
}
} else {
const int threads = 256;
const int blocks = M;
auto* halfB = reinterpret_cast<const half*>(B);
auto* halfY = reinterpret_cast<half*>(Y);
if (relu) {
InplaceAddReluKernel<true, threads>
<<<blocks, threads, 0, stream>>>(N, halfB, halfY);
} else {
InplaceAddReluKernel<false, threads>
<<<blocks, threads, 0, stream>>>(N, halfB, halfY);
}
}
}

#else
template <bool DoRelu, int BlockDim>
__global__ void InplaceAddReluKernel(const int N,
const float16* bias,
float16* data) {
int offset = blockIdx.x * N;

for (int i = threadIdx.x; i < N; i += BlockDim) {
float16 temp;
temp = data[offset + i] + bias[i];
Expand Down Expand Up @@ -260,7 +168,6 @@ void AddReluKernel(gpuStream_t stream,
}
}
}
#endif

template <typename DeviceContext, typename T>
void FCFunctor<DeviceContext, T>::operator()(const DeviceContext& context,
Expand Down