 
     if (compute_virial)
         {
-        virialxx = aniso_warp_reduce(tpp, threadIdx.x % tpp, virialxx, &sh[cta_offs]);
-        virialxy = aniso_warp_reduce(tpp, threadIdx.x % tpp, virialxy, &sh[cta_offs]);
-        virialxz = aniso_warp_reduce(tpp, threadIdx.x % tpp, virialxz, &sh[cta_offs]);
-        virialyy = aniso_warp_reduce(tpp, threadIdx.x % tpp, virialyy, &sh[cta_offs]);
-        virialyz = aniso_warp_reduce(tpp, threadIdx.x % tpp, virialyz, &sh[cta_offs]);
-        virialzz = aniso_warp_reduce(tpp, threadIdx.x % tpp, virialzz, &sh[cta_offs]);
+        virialxx = reducer.Sum(virialxx);
+        virialxy = reducer.Sum(virialxy);
+        virialxz = reducer.Sum(virialxz);
+        virialyy = reducer.Sum(virialyy);
+        virialyz = reducer.Sum(virialyz);
+        virialzz = reducer.Sum(virialzz);
 
         // if we are the first thread in the cta, write out virial to global mem
-        if (threadIdx.x %tpp == 0)
+        if (active && threadIdx.x %tpp == 0)
             {
             d_virial[0*virial_pitch+idx] = virialxx;
             d_virial[1*virial_pitch+idx] = virialxy;

    assert(pair_args.d_rcutsq);
    assert(pair_args.ntypes > 0);

    Index2D typpair_idx(pair_args.ntypes);
    unsigned int shared_bytes = (2*sizeof(Scalar) + sizeof(typename evaluator::param_type))
                                * typpair_idx.getNumElements();

    unsigned int block_size = pair_args.block_size;
    unsigned int tpp = pair_args.threads_per_particle;

    // iterate over active GPUs in reverse, to end up on first GPU when returning from this function
    for (int idev = pair_args.gpu_partition.getNumActiveGPUs() - 1; idev >= 0; --idev)
        {
        auto range = pair_args.gpu_partition.getRangeAndSetGPU(idev);
        unsigned int nwork = range.second - range.first;

        // run the kernel
        if (pair_args.compute_virial)

#include "hoomd/TextureTools.h"
#include "hoomd/GPUPartition.cuh"

#ifdef NVCC
#include "hoomd/WarpTools.cuh"
#endif // NVCC

/*! \file AnisoPotentialPairGPU.cuh
    \brief Defines templated GPU kernel code for calculating the anisotropic ptl pair forces and torques
