`__Shfl_sync` не работает для 64-битного?

Question

Я реализовал сокращение ширины основы и блока, используя инструкции тасования. Все работает нормально, когда я использую 32-битные типы, но для 64-битных я всегда получаю 0 в результате. Насколько мне известно, shuffling поддерживает 64-битные аргументы. Чего мне не хватает?

#include <stdio.h>

template<typename T>
inline __device__ T warpRegSumTest(T val) {
  T result = val;
  static constexpr unsigned mask = 0xffffffff;
#pragma unroll
  for (int delta = 16; delta > 0; delta /= 2) {
    result = result + __shfl_down_sync(mask, result, delta);
  }
  return result;
}

template<int numWarpsInBlock, typename T>
inline __device__ T blockRegSumTest(T val) {
  __shared__ T part[numWarpsInBlock];
  T warppart = warpRegSumTest(val);
  if (threadIdx.x % 32 == 0) {
    part[threadIdx.x / 32] = warppart;
  }
  __syncthreads();
  if (threadIdx.x < 32) {
    int tid = threadIdx.x;
    T solution = warpRegSumTest(tid < numWarpsInBlock ? part[tid] : T(0));
    __syncwarp();
    part[0] = solution;
  }

  __syncthreads();
  T result = part[0];
  __syncthreads();
  return result;
}

__global__ void testKernel() {
  float float_result = blockRegSumTest<256 / 32>(float(threadIdx.x));
  if (threadIdx.x == 0) {
    printf("Float sum: %f\n", float_result);
  }
  double double_result = blockRegSumTest<256 / 32>(double(threadIdx.x));
  if (threadIdx.x == 0) {
    printf("Double sum: %f\n", double_result);
  }
  int int_result = blockRegSumTest<256 / 32>(int(threadIdx.x));
  if (threadIdx.x == 0) {
    printf("Int sum: %d\n", int_result);
  }
  long long longlong_result = blockRegSumTest<256 / 32>(long long(threadIdx.x));
  if (threadIdx.x == 0) {
    printf("Long long sum: %lld\n", longlong_result);
  }
}

int main()
{
  testKernel << <1, 256 >> > ();
}

Я компилирую это с compute_70,sm_70 и работаю на GTX 2070 SUPER. Он выводит:

Float sum: 32640.000000
Double sum: 0.000000
Int sum: 32640
Long long sum: 0

Я ожидал увидеть 32640 (сумма 0 + 1 + 2 + ... + 255) во всех 4 случаях.

Answer 1

У вас здесь ошибка:

part[0] = solution;

должно быть:

if (!threadIdx.x) part[0] = solution;

Вы хотите, чтобы поток 0 выполнял эту строку.