0

使用 CUDA，我尝试使用由向量组成的数组（在我的情况下为 int2）进行编译时优化，但我无法以干净的方式实现这一点。更具体地说，我正在研究一个使用两个常量数组 c 和 w 的问题。数组 w 由浮点数组成，数组 c 由 int2 组成。现在，由于这些数组是常量，我希望编译器执行编译时优化，从而有效地优化数组访问。例如，对于以下两个设备函数，编译器展开循环并通过直接将其替换为 c 和 w 的值来优化数组访问：

__forceinline__ __device__ float someFunction1() {
  const int2  c[9] = {make_int2(0, 0), make_int2(1, 0), make_int2(0, 1), make_int2(-1, 0), make_int2(0, -1),
                      make_int2(1, 1), make_int2(-1, 1), make_int2(-1, -1), make_int2(1, -1)};
  const float w[9] = {4.0f/9.0f, 1.0f/9.0f, 1.0f/9.0f, 1.0f/9.0f, 1.0f/9.0f, 1.0f/36.0f, 1.0f/36.0f, 1.0f/36.0f, 1.0f/36.0f};
  #pragma unroll
  for (int i = 0; i < 9; ++i) {
    //Do something here, accessing c[i] and w[i]
  }
}

__forceinline__ __device__ float someFunction2() {
  const int2  c[9] = {make_int2(0, 0), make_int2(1, 0), make_int2(0, 1), make_int2(-1, 0), make_int2(0, -1),
                           make_int2(1, 1), make_int2(-1, 1), make_int2(-1, -1), make_int2(1, -1)};
  const float w[9] = {4.0f/9.0f, 1.0f/9.0f, 1.0f/9.0f, 1.0f/9.0f, 1.0f/9.0f, 1.0f/36.0f, 1.0f/36.0f, 1.0f/36.0f, 1.0f/36.0f};
  #pragma unroll
  for (int i = 0; i < 9; ++i) {
    //Do something here, accessing c[i] and w[i]
  }
}

现在，问题是我不想在每个使用 c 和 w 的设备函数中连续声明 c 和 w。我可以全局声明 w，但不允许全局声明 c，因为 CUDA 不允许我在全局变量中调用 make_int2 构造函数。也就是说，下面的程序给出了错误“无法为设备上的非空构造函数或析构函数生成代码”：

//Declaring array c like this is not allowed
__device__ const int2  c[9] = {make_int2(0, 0), make_int2(1, 0), make_int2(0, 1), make_int2(-1, 0), make_int2(0, -1),
                               make_int2(1, 1), make_int2(-1, 1), make_int2(-1, -1), make_int2(1, -1)};
__device__ const float w[9] = {4.0f/9.0f, 1.0f/9.0f, 1.0f/9.0f, 1.0f/9.0f, 1.0f/9.0f, 1.0f/36.0f, 1.0f/36.0f, 1.0f/36.0f, 1.0f/36.0f};

__forceinline__ __device__ float someFunction() {
  #pragma unroll
  for (int i = 0; i < 9; ++i) {
    //Do something here, accessing c[i] and w[i]
  }
}

我的问题是：如何防止在访问这些变量的每个函数中声明 c 和 w 并且仍然具有我想要的编译时优化？或者另有说明：是否有一种解决方法来全局声明一个向量数组？

注意：我知道我可以将 c 和 w 存储在全局或 __constant__ 内存中，但这不会给我编译时优化。__constant__ 内存在不定期访问时也可能出现问题。

4

1 回答 1

1

我不知道这是否真的可以实现您在编译器优化方面所寻找的东西，但是将指针从 int 转换为 int2 似乎对我有用：

#include <stdio.h>

  __device__ const int  ci[16] = {0, 0, 1, 0, 0, 1, -1, 0, 0, -1, 1, 1, -1, 1, -1, -1};

  __device__ const int2 *c = (const int2 *)ci;

  __global__ void mykernel(){

  int2 temp = c[1];
  int2 temp1 = c[4];
  printf("c[1].x = %d\n", temp.x);
  printf("c[4].y = %d\n", temp1.y);
  }

int main(){

  mykernel<<<1,1>>>();
  cudaDeviceSynchronize();
  printf("Done\n");
  return 0;
}

请注意，const和__ constant__不是一回事。您可以从andconst的变量定义中消除声明，但我认为将它放在那里将有助于编译器实现您想要的。cci

于 2013-03-08T05:54:03.390 回答