Skip to content
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

improve dropout #29465

Merged
merged 4 commits into from
Dec 11, 2020
Merged

improve dropout #29465

Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
91f4709
improve drop out
Dec 7, 2020
3a29f8d
add VectorizedRandomGeneratorWithGenerator
Dec 8, 2020
4964a47
fix bug
Dec 8, 2020
581430e
modify according to comments
Dec 9, 2020
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
191 changes: 99 additions & 92 deletions paddle/fluid/operators/dropout_op.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,6 +17,7 @@ limitations under the License. */
#include <thrust/iterator/counting_iterator.h>
#include <thrust/random.h>
#include <thrust/transform.h>
#include <algorithm>
#include <string>
#include "paddle/fluid/memory/memcpy.h"
#include "paddle/fluid/operators/dropout_op.h"
Expand All @@ -26,60 +27,35 @@ limitations under the License. */
namespace paddle {
namespace operators {

template <typename T, typename MaskType>
__global__ void RandomGenerator(const size_t n, const int seed,
const float dropout_prob, const T* src,
MaskType* mask_data, T* dst,
bool is_upscale_in_train) {
curandStatePhilox4_32_10_t state;
int idx = blockDim.x * blockIdx.x + threadIdx.x;
int step_size = 0;
// aligned vector generates vectorized load/store on CUDA
template <typename T, int Size>
struct alignas(sizeof(T) * Size) AlignedVector {
T val[Size];
};

MaskType mask;
T dest;
for (; idx < n; idx += blockDim.x * gridDim.x) {
T s = src[idx];
if (step_size == 0) {
curand_init(seed, idx, idx, &state);
step_size = blockDim.x * gridDim.x;
} else {
curand_init(seed, idx, step_size, &state);
}
if (curand_uniform(&state) < dropout_prob) {
mask = 0;
dest = 0;
} else {
mask = 1;
if (is_upscale_in_train) {
dest = s / static_cast<T>(1.0f - dropout_prob);
} else {
dest = s;
}
}
mask_data[idx] = mask;
dst[idx] = dest;
template <typename T>
inline int VectorizedSize(const T* pointer) {
uint64_t address = reinterpret_cast<uint64_t>(pointer);
constexpr int vec4 = std::alignment_of<AlignedVector<T, 4>>::value; // NOLINT
if (address % vec4 == 0) {
return 4;
}
return 1;
}

template <typename T, typename MaskType>
__global__ void RandomGeneratorWithSeed(const size_t n, const int* seed,
const float dropout_prob, const T* src,
MaskType* mask_data, T* dst,
bool is_upscale_in_train) {
__global__ void RandomGenerator(const size_t n, uint64_t seed,
const float dropout_prob, const T* src,
MaskType* mask_data, T* dst,
bool is_upscale_in_train, uint64_t increment) {
curandStatePhilox4_32_10_t state;
int idx = blockDim.x * blockIdx.x + threadIdx.x;
int step_size = 0;
curand_init(seed, idx, increment, &state);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

这个修改是等价的吗？原来curand_init是放在for循环里面，每次迭代都会调用一次。increment参数是指什么？

Copy link
Contributor Author

@zhangting2020 zhangting2020 Dec 9, 2020 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

我认为这样改是没有问题的，并不需要每次迭代里面去初始化，并且动态修改offset，以下详细解释：

curand_init (
    unsigned long long seed, unsigned long long subsequence,
    unsigned long long offset, curandState_t *state)

curand_init根据给定的seed、subsequence和offset设置初始的state。在生成随机数时，根据subsequence和offset来决定开始取随机数的位置。

|seq0  ...   2^67-1|seq1   ...   2^68-1|seq2   ... 
                           ^
                   |offset |                     (determined by offset parameter)
                           |
                           RNG begins here for given seed, sequence(=seq1), offset

根据原来的逻辑：它每次迭代重新init，只是重新设置了取随机数的位置，也就是（sequence的索引和offset）。

for (; idx < n; idx += blockDim.x * gridDim.x) {
    T s = src[idx];
    if (step_size == 0) {
      curand_init(seed, idx, idx, &state);
      step_size = blockDim.x * gridDim.x;
    } else {
      curand_init(seed, idx, step_size, &state);
    }
...

首先用一段程序解释curand_init 原来写在for循环里，和现在写在for循环外的区别：

  • curand_init写在for循环里：
_global__ void testrand1(unsigned long seed, float *a, int N){
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    curandState state;
    int step_size = 0;
    for (int i = idx; i < N; i += blockDim.x * gridDim.x) {
        if (step_size == 0) {
          curand_init(seed, idx, i, &state);
          step_size = blockDim.x * gridDim.x;
        } else {
          curand_init(seed, idx, step_size, &state);
        }
        a[i] = curand_uniform(&state);
    }
}

0 0.145468
1 0.926417
2 0.782640
3 0.535606
4 0.650189
5 0.629326
6 0.713179
7 0.448197
8 0.300772
9 0.136307
  • 下面这段是目前PR里的写法：
__global__ void testrand1(unsigned long seed, float *a, int N){
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    curandState state;
    curand_init(seed, idx, idx, &state);
    int step_size = 0;
    for (int i = idx; i < N; i += blockDim.x * gridDim.x) {
        a[i] = curand_uniform(&state);
    }
}

int main() {

    const int N = 10;

    float *h_a  = (float*)malloc(N*sizeof(float));
    float *d_a;
    cudaMalloc((void**)&d_a, N*sizeof(float));
    int thread = 2;
    // int grid = (N + thread - 1) / thread;
    int grid = 4;
    testrand1<<<grid, thread>>>(1234, d_a, N);
    cudaPeekAtLastError();
    cudaDeviceSynchronize();

    cudaMemcpy(h_a, d_a, N*sizeof(float), cudaMemcpyDeviceToHost);

    for (int i=0; i<N; i++) printf("%i %f\n", i, h_a[i]);

    getchar();
}

0 0.145468
1 0.926417
2 0.782640
3 0.535606
4 0.650189
5 0.629326
6 0.713179
7 0.448197
8 0.434899
9 0.511765

idx取值为0~7。
第一种写法是：

  • idx=0时
    • 第一次迭代：产生子序列sequence0，然后设置随机数a[0]=sequence0[0]，同时offset=8
    • 第二次迭代：产生子序列sequence8，然后设置随机数a[8]=sequence8[8]；
  • idx=1时
    • 第一次迭代：产生子序列sequence1，然后设置随机数a[1]=sequence1[1]，同时offset=8
    • 第二次迭代：产生子序列sequence8，然后设置随机数a[9]=sequence9[8]；

第二种写法是：

  • idx=0时，产生子序列sequence0，然后设置随机数a[0]=sequence0[0]和a[8]=sequence0[1]；
  • idx=1时，产生子序列sequence1，然后设置随机数a[1]=sequence1[1]和a[9]=sequence1[2]；

curand_uniform每调用一次，就往下取一个数。因此前8个数，2种写法都是相同的。从上面的分析可以看到，a[8]和a[9]两种写法取得数不同。

只是取的随机数不同而已，但是我理解并不影响这个op的功能。而且我认为原始的写法没有必要。

最后再说PR里的写法：

  • offset是在调用cuda kernel前算好的，计算一个increment用来设置offset，increment不能小于每个线程产生的随机数个数。每一次调用drop out后，会设置offset = offset + increament。 那么下一次调用drop out时，每个线程就会跳过前一次drop out时取过的那些随机数。例如：

  • 第一次调用drop out产生10个数，当grid=4， thread=2时，会将curand_init的当前的offset设置为0，同时会将线程的offset置为 offset= 0 + 2，因为本次每个线程最多会产生2个数。

  • 第二次调用drop out时，curand_init的offset就会使用2，然后再设置offset = offset + increament（increament根据这次drop out每个线程产生的随机数确定）。从上一次取过的随机数之后去取数，以免和上一次调用drop out产生的随机数出现重叠。

    • idx=0时，产生子序列sequence0，然后设置随机数a[0]=sequence0[2]和a[8]=sequence0[2+1]；
    • idx=1时，产生子序列sequence1，然后设置随机数a[1]=sequence1[2]和a[9]=sequence1[2+1]；


MaskType mask;
T dest;
for (; idx < n; idx += blockDim.x * gridDim.x) {
T s = src[idx];
if (step_size == 0) {
curand_init(seed[0], idx, idx, &state);
step_size = blockDim.x * gridDim.x;
} else {
curand_init(seed[0], idx, step_size, &state);
}
if (curand_uniform(&state) < dropout_prob) {
mask = 0;
dest = 0;
Expand All @@ -96,39 +72,49 @@ __global__ void RandomGeneratorWithSeed(const size_t n, const int* seed,
}
}

template <typename T, typename MaskType>
__global__ void RandomGeneratorWithGenerator(const size_t n, uint64_t seed,
const float dropout_prob,
const T* src, MaskType* mask_data,
T* dst, bool is_upscale_in_train,
uint64_t increment) {
template <typename T, typename MaskType, int VecSize>
__global__ void VectorizedRandomGenerator(const size_t n, uint64_t seed,
const float dropout_prob,
const T* src, MaskType* mask_data,
T* dst, bool is_upscale_in_train,
uint64_t increment) {
int64_t idx = blockDim.x * blockIdx.x + threadIdx.x;
curandStatePhilox4_32_10_t state;
int idx = blockDim.x * blockIdx.x + threadIdx.x;
int step_size = 0;
curand_init(seed, idx, increment, &state);

MaskType mask;
T dest;
for (; idx < n; idx += blockDim.x * gridDim.x) {
T s = src[idx];
if (step_size == 0) {
curand_init(seed, idx, increment, &state);
step_size = blockDim.x * gridDim.x;
} else {
curand_init(seed, idx, increment, &state);
}
if (curand_uniform(&state) < dropout_prob) {
mask = 0;
dest = 0;
} else {
mask = 1;
if (is_upscale_in_train) {
dest = s / static_cast<T>(1.0f - dropout_prob);
using LoadT = AlignedVector<T, VecSize>;
using MaskLoadT = AlignedVector<MaskType, VecSize>;
T factor = static_cast<T>(1.0f / (1.0f - dropout_prob));
for (int i = idx * VecSize; i < n; i += blockDim.x * gridDim.x * VecSize) {
T src_vec[VecSize];
LoadT* value = reinterpret_cast<LoadT*>(&src_vec);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

有没有试过直接用float4 __ldg(const float4 *ptr)来加载？

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

我试了，这里不能这么使用，问了下wangchao，看上去是不支持，

*value = *reinterpret_cast<const LoadT*>(&src[i]);
float4 rand = curand_uniform4(&state);

T dest_vec[VecSize];
MaskType mask_vec[VecSize];

#pragma unroll
for (int ii = 0; ii < VecSize; ii++) {
if ((&rand.x)[ii] < dropout_prob) {
dest_vec[ii] = 0;
mask_vec[ii] = 0;
} else {
dest = s;
if (is_upscale_in_train) {
dest_vec[ii] = src_vec[ii] * factor;
} else {
dest_vec[ii] = src_vec[ii];
}
mask_vec[ii] = 1;
}
}
mask_data[idx] = mask;
dst[idx] = dest;

*(reinterpret_cast<LoadT*>(&dst[i])) =
*reinterpret_cast<LoadT*>(&dest_vec[0]);
*(reinterpret_cast<MaskLoadT*>(&mask_data[i])) =
*reinterpret_cast<MaskLoadT*>(&mask_vec[0]);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

这个写回好长。。。

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

这个我好像也不确定还有什么更好的改法

}
}

Expand Down Expand Up @@ -170,36 +156,57 @@ class GPUDropoutKernel : public framework::OpKernel<T> {

int threads = 512;
int grid = (x_numel + threads - 1) / threads;
const auto& dev_ctx = context.cuda_device_context();
int blocks_per_sm =
dev_ctx.GetMaxPhysicalThreadCount() / dev_ctx.GetSMCount() / threads;
grid = std::min(dev_ctx.GetSMCount() * blocks_per_sm, grid);

// increment is used to set the args(offset) of curand_init, which defines
// offset in subsequence.
// The detail:
// https://docs.nvidia.com/cuda/curand/device-api-overview.html
// Increment should be at least the number of curand() random numbers used
// in each thread to avoid the random number generated this time being the
// same as the previous calls.
uint64_t seed_data;
uint64_t increment;
int vec_size = VectorizedSize<T>(x_data);
auto offset =
((x_numel - 1) / (threads * grid * vec_size) + 1) * vec_size;
int device_id = BOOST_GET_CONST(platform::CUDAPlace, context.GetPlace())
.GetDeviceId();
auto gen_cuda = framework::GetDefaultCUDAGenerator(device_id);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

这个随机数种子生成器，是说CPUGenerator生成的种子Tensor是在CPU上，CUDAGenerator生成的种子Tensor在GPU上吗？我们的种子看来是希望在CPU上访问的，那是不是可以直接用CPUGenerator？

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

这里应该不是区分种子是否在CPU还是GPU上。 原来的RandomGeneratorWithGenerator kernel需要用到increment，来设置curand_init中的offset参数。op里会用到下面的IncrementOffset去获得当前的offset。因为dropout在训练中会多次调用，没调用一次，就应该改变offset，跳过这一次产生的这些随机数。这样下一次调用drop out产生的随机数和之前就不会发生重叠。

std::pair<uint64_t, uint64_t> Generator::IncrementOffset(
    uint64_t increament_offset) {
  uint64_t cur_offset = this->state_.thread_offset;
#ifdef PADDLE_WITH_CUDA
  std::lock_guard<std::mutex> lock(this->mu_);

  this->state_.thread_offset += increament_offset;

#else
  PADDLE_THROW(platform::errors::PermissionDenied(
      "Increment Offset only support in CUDA place"));
#endif
  return std::make_pair(static_cast<int>(this->state_.current_seed),
                        cur_offset);
}


if (seed && platform::is_gpu_place(seed->place())) {
auto seed_gpu_data = seed->data<int>();
RandomGeneratorWithSeed<T, uint8_t><<<grid, threads, 0, stream>>>(
size, seed_gpu_data, dropout_prob, x_data, mask_data, y_data,
upscale_in_train);
return;
}
int seed_data;
std::random_device rnd;
if (seed) {
seed_data = *(seed->data<int>());
framework::Tensor seed_cpu_tensor;
TensorCopySync(*seed, platform::CPUPlace(), &seed_cpu_tensor);
seed_data = static_cast<uint64_t>(seed_cpu_tensor.data<int>()[0]);
increment = offset;
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

感觉这个increment跟原来的设置不完全一样，这个参数的影响是什么？代码里面能不能加个注释说明下。

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

其他的2个kernel并不会用到这个参数，为了将3个kernel进行统一，所以其他的情况也设置了。increment作用就是用来设置curand_init的offset，上面已经解释过了原理。

} else if (gen_cuda->GetIsInitPy() && (!context.Attr<bool>("fix_seed"))) {
auto seed_offset = gen_cuda->IncrementOffset(offset);
seed_data = seed_offset.first;
increment = seed_offset.second;
} else {
seed_data =
context.Attr<bool>("fix_seed") ? context.Attr<int>("seed") : rnd();
if (seed) {
seed_data = *(seed->data<int>());
} else {
std::random_device rnd;
seed_data = context.Attr<bool>("fix_seed") ? context.Attr<int>("seed")
: rnd();
}
increment = offset;
}

int device_id = BOOST_GET_CONST(platform::CUDAPlace, context.GetPlace())
.GetDeviceId();
auto gen_cuda = framework::GetDefaultCUDAGenerator(device_id);
if (gen_cuda->GetIsInitPy() && (!context.Attr<bool>("fix_seed"))) {
auto seed_offset = gen_cuda->IncrementOffset(1);
RandomGeneratorWithGenerator<T, uint8_t><<<grid, threads, 0, stream>>>(
size, seed_offset.first, dropout_prob, x_data, mask_data, y_data,
upscale_in_train, seed_offset.second);
return;
if (vec_size == 4) {
VectorizedRandomGenerator<T, uint8_t, 4><<<grid, threads, 0, stream>>>(
size, seed_data, dropout_prob, x_data, mask_data, y_data,
upscale_in_train, increment);
} else {
RandomGenerator<T, uint8_t><<<grid, threads, 0, stream>>>(
size, seed_data, dropout_prob, x_data, mask_data, y_data,
upscale_in_train, increment);
}

RandomGenerator<T, uint8_t><<<grid, threads, 0, stream>>>(
size, seed_data, dropout_prob, x_data, mask_data, y_data,
upscale_in_train);
} else {
auto X = EigenMatrix<T>::Reshape(*x, 1);
auto Y = EigenMatrix<T>::Reshape(*y, 1);
Expand Down