时间戳

首页 > 解决方案 > 64 位的 __shfl_sync 是否损坏？

cuda - 64 位的 __shfl_sync 是否损坏？

问题描述

我已经使用 shuffle 指令实现了经纱范围和块范围的缩减。当我使用 32 位类型时，一切正常，但对于 64 位，我总是得到 0 结果。据我所知，改组支持 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

我预计在所有 4 种情况下都会看到 32640（总和 0+1+2+...+255）。

标签： cuda

解决方案

你在这里有一个错误：

part[0] = solution;

它应该是：

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

您只希望线程 0 执行该行。

推荐阅读