mirror of
https://gitlab.com/libeigen/eigen.git
synced 2024-12-15 07:10:37 +08:00
Resolved merge conflicts
This commit is contained in:
parent
28fcb5ca2a
commit
36369ab63c
@ -325,11 +325,7 @@ __global__ void ReductionInitFullReduxKernelHalfFloat(R, const S, I, half2*);
|
||||
template <int B, int N, typename S, typename R, typename I>
|
||||
__global__ void FullReductionKernelHalfFloat(R, const S, I, half*, half2*);
|
||||
template <int NPT, typename S, typename R, typename I>
|
||||
<<<<<<< local
|
||||
__global__ void InnerReductionKernelHalfFloat(R, const S, I, I, half*);
|
||||
=======
|
||||
__global__ void InnerReductionKernelHalfFloat(R, const S, I, I, half*, half2*);
|
||||
>>>>>>> other
|
||||
|
||||
#endif
|
||||
|
||||
@ -624,11 +620,7 @@ struct TensorEvaluator<const TensorReductionOp<Op, Dims, ArgType>, Device>
|
||||
#ifdef EIGEN_HAS_CUDA_FP16
|
||||
template <typename S, typename R, typename I> friend void internal::ReductionInitFullReduxKernelHalfFloat(R, const S, I, half2*);
|
||||
template <int B, int N, typename S, typename R, typename I> friend void internal::FullReductionKernelHalfFloat(R, const S, I, half*, half2*);
|
||||
<<<<<<< local
|
||||
template <int NPT, typename S, typename R, typename I> friend void internal::InnerReductionKernelHalfFloat(R, const S, I, I, half*);
|
||||
=======
|
||||
template <int NPT, typename S, typename R, typename I> friend void internal::InnerReductionKernelHalfFloat(R, const S, I, I, half*, half2*);
|
||||
>>>>>>> other
|
||||
#endif
|
||||
template <int NPT, typename S, typename R, typename I> friend void internal::InnerReductionKernel(R, const S, I, I, typename S::CoeffReturnType*);
|
||||
|
||||
|
@ -391,13 +391,8 @@ __global__ void InnerReductionKernelHalfFloat(Reducer reducer, const Self input,
|
||||
eigen_assert(NumPerThread % unroll_times == 0);
|
||||
eigen_assert(unroll_times % 2 == 0);
|
||||
|
||||
<<<<<<< local
|
||||
const Index input_col_blocks = divup<Index>(num_coeffs_to_reduce, blockDim.x * NumPerThread * 2);
|
||||
const Index num_input_blocks = divup<Index>(input_col_blocks * num_preserved_coeffs, 2);
|
||||
=======
|
||||
const Index input_col_blocks = divup<Index>(num_coeffs_to_reduce, blockDim.x * NumPerThread/2);
|
||||
const Index num_input_blocks = input_col_blocks * num_preserved_coeffs;
|
||||
>>>>>>> other
|
||||
|
||||
const Index num_threads = blockDim.x * gridDim.x;
|
||||
const Index thread_id = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
@ -406,12 +401,8 @@ __global__ void InnerReductionKernelHalfFloat(Reducer reducer, const Self input,
|
||||
if (gridDim.x == 1) {
|
||||
Index i = 2*thread_id;
|
||||
for (; i + 1 < num_preserved_coeffs; i += 2*num_threads) {
|
||||
<<<<<<< local
|
||||
half* loc = output + i;
|
||||
*((half2*)loc) = reducer.template initializePacket<half2>();
|
||||
=======
|
||||
((half2*)output)[i] = reducer.template initializePacket<half2>();
|
||||
>>>>>>> other
|
||||
}
|
||||
if (i < num_preserved_coeffs) {
|
||||
output[i] = reducer.initialize();
|
||||
@ -419,13 +410,8 @@ __global__ void InnerReductionKernelHalfFloat(Reducer reducer, const Self input,
|
||||
__syncthreads();
|
||||
}
|
||||
|
||||
<<<<<<< local
|
||||
for (Index i = blockIdx.x; i < num_input_blocks; i += gridDim.x) {
|
||||
const Index row = 2 * (i / input_col_blocks);
|
||||
=======
|
||||
for (Index i = 2*blockIdx.x; i < num_input_blocks; i += 2*gridDim.x) {
|
||||
const Index row = i / input_col_blocks;
|
||||
>>>>>>> other
|
||||
|
||||
if (row + 1 < num_preserved_coeffs) {
|
||||
const Index col_block = i % input_col_blocks;
|
||||
@ -446,18 +432,10 @@ __global__ void InnerReductionKernelHalfFloat(Reducer reducer, const Self input,
|
||||
}
|
||||
if (col < num_coeffs_to_reduce) {
|
||||
// Peel;
|
||||
<<<<<<< local
|
||||
const half last1 = input.m_impl.coeff(row * num_coeffs_to_reduce + col);
|
||||
=======
|
||||
const half last1 = input.m_impl.coeff(row * num_coeffs_to_reduce + col+1);
|
||||
>>>>>>> other
|
||||
const half2 val1 = __halves2half2(last1, reducer.initialize());
|
||||
reducer.reducePacket(val1, &reduced_val1);
|
||||
<<<<<<< local
|
||||
const half last2 = input.m_impl.coeff((row+1) * num_coeffs_to_reduce + col);
|
||||
=======
|
||||
const half last2 = input.m_impl.coeff((row+1) * num_coeffs_to_reduce + col+1);
|
||||
>>>>>>> other
|
||||
const half2 val2 = __halves2half2(last2, reducer.initialize());
|
||||
reducer.reducePacket(val2, &reduced_val2);
|
||||
}
|
||||
@ -466,17 +444,9 @@ __global__ void InnerReductionKernelHalfFloat(Reducer reducer, const Self input,
|
||||
// Faster version of the loop with no branches after unrolling.
|
||||
#pragma unroll
|
||||
for (int k = 0; k < unroll_times; ++k) {
|
||||
<<<<<<< local
|
||||
const Index col = col_begin + blockDim.x * (j + k) * 2;
|
||||
=======
|
||||
const Index col = col_begin + blockDim.x * (j + k);
|
||||
>>>>>>> other
|
||||
reducer.reducePacket(input.m_impl.template packet<Unaligned>(row * num_coeffs_to_reduce + col), &reduced_val1);
|
||||
<<<<<<< local
|
||||
reducer.reducePacket(input.m_impl.template packet<Unaligned>((row + 1)* num_coeffs_to_reduce + col), &reduced_val2);
|
||||
=======
|
||||
reducer.reducePacket(input.m_impl.template packet<Unaligned>((row +1)* num_coeffs_to_reduce + col), &reduced_val2);
|
||||
>>>>>>> other
|
||||
}
|
||||
}
|
||||
}
|
||||
@ -494,12 +464,8 @@ __global__ void InnerReductionKernelHalfFloat(Reducer reducer, const Self input,
|
||||
half2 val = __halves2half2(val1, val2);
|
||||
|
||||
if ((threadIdx.x & (warpSize - 1)) == 0) {
|
||||
<<<<<<< local
|
||||
half* loc = output + row;
|
||||
atomicReduce((half2*)loc, val, reducer);
|
||||
=======
|
||||
atomicReduce(&(((half2*)output)[row]), val, reducer);
|
||||
>>>>>>> other
|
||||
}
|
||||
}
|
||||
}
|
||||
@ -554,33 +520,18 @@ struct InnerReductionLauncher {
|
||||
static bool run(const Self& self, Op& reducer, const GpuDevice& device, half* output, typename Self::Index num_coeffs_to_reduce, typename Self::Index num_preserved_vals) {
|
||||
typedef typename Self::Index Index;
|
||||
|
||||
<<<<<<< local
|
||||
if (num_preserved_vals % 2 != 0) {
|
||||
// Not supported yet, revert to the slower code path
|
||||
std::cout << "BYPASSING OPTIMIZED CODE PATH" << std::endl;
|
||||
=======
|
||||
// It's faster to use the usual code.
|
||||
if (num_coeffs_to_reduce <= 32) {
|
||||
>>>>>>> other
|
||||
return true;
|
||||
}
|
||||
|
||||
const Index num_coeffs = num_coeffs_to_reduce * num_preserved_vals;
|
||||
<<<<<<< local
|
||||
const int block_size = /*256*/128;
|
||||
const int num_per_thread = /*128*/64;
|
||||
=======
|
||||
const int block_size = 256;
|
||||
const int num_per_thread = 128;
|
||||
>>>>>>> other
|
||||
const int dyn_blocks = divup<int>(num_coeffs, block_size * num_per_thread);
|
||||
const int max_blocks = device.getNumCudaMultiProcessors() *
|
||||
device.maxCudaThreadsPerMultiProcessor() / block_size;
|
||||
const int num_blocks = numext::mini<int>(max_blocks, dyn_blocks);
|
||||
<<<<<<< local
|
||||
=======
|
||||
half2* scratch = static_cast<half2*>(device.scratchpad());
|
||||
>>>>>>> other
|
||||
|
||||
if (num_blocks > 1) {
|
||||
// We initialize the outputs outside the reduction kernel when we can't be sure that there
|
||||
@ -590,19 +541,11 @@ struct InnerReductionLauncher {
|
||||
device.maxCudaThreadsPerMultiProcessor() / 1024;
|
||||
const int num_blocks = numext::mini<int>(max_blocks, dyn_blocks);
|
||||
LAUNCH_CUDA_KERNEL((ReductionInitKernelHalfFloat<Self, Op, Index>),
|
||||
<<<<<<< local
|
||||
1, 1, 0, device, reducer, self, num_preserved_vals, output);
|
||||
=======
|
||||
1, 1, 0, device, reducer, self, num_preserved_vals, scratch);
|
||||
>>>>>>> other
|
||||
}
|
||||
|
||||
LAUNCH_CUDA_KERNEL((InnerReductionKernelHalfFloat<num_per_thread, Self, Op, Index>),
|
||||
<<<<<<< local
|
||||
num_blocks, block_size, 0, device, reducer, self, num_coeffs_to_reduce, num_preserved_vals, output);
|
||||
=======
|
||||
num_blocks, block_size, 0, device, reducer, self, num_coeffs_to_reduce, num_preserved_vals, output, scratch);
|
||||
>>>>>>> other
|
||||
|
||||
return false;
|
||||
}
|
||||
@ -632,14 +575,11 @@ struct InnerReducer<Self, Op, GpuDevice> {
|
||||
if (num_coeffs == 0) {
|
||||
return true;
|
||||
}
|
||||
<<<<<<< local
|
||||
// It's faster to use the usual code.
|
||||
if (num_coeffs_to_reduce <= 128) {
|
||||
return true;
|
||||
}
|
||||
=======
|
||||
|
||||
>>>>>>> other
|
||||
return InnerReductionLauncher<Self, Op>::run(self, reducer, device, output, num_coeffs_to_reduce, num_preserved_vals);
|
||||
}
|
||||
};
|
||||
|
@ -255,13 +255,8 @@ void test_cuda_reductions(int size1, int size2, int redux) {
|
||||
|
||||
Eigen::CudaStreamDevice stream;
|
||||
Eigen::GpuDevice gpu_device(&stream);
|
||||
<<<<<<< local
|
||||
int num_elem = size1*size2;
|
||||
int result_size = (redux == 1 ? size1 : size2);
|
||||
=======
|
||||
int size = 40;
|
||||
int num_elem = size*size;
|
||||
>>>>>>> other
|
||||
|
||||
float* d_float1 = (float*)gpu_device.allocate(num_elem * sizeof(float));
|
||||
float* d_float2 = (float*)gpu_device.allocate(num_elem * sizeof(float));
|
||||
@ -291,6 +286,7 @@ void test_cuda_reductions(int size1, int size2, int redux) {
|
||||
gpu_device.synchronize();
|
||||
|
||||
for (int i = 0; i < result_size; ++i) {
|
||||
std::cout << "EXPECTED " << full_prec(i) << " GOT " << half_prec(i) << std::endl;
|
||||
VERIFY_IS_APPROX(full_prec(i), half_prec(i));
|
||||
}
|
||||
|
||||
|
Loading…
Reference in New Issue
Block a user