Глобальный массив работы с векторным типом

Question

С CUDA Я пытаюсь оптимизировать время компиляции с массивом, состоящим из векторов (int2 в моем случае), но я не могу добиться этого в чистом виде. Позвольте быть более конкретным, я работаю над проблемой, которая использует два постоянных массива c и w. Array w состоит из float и array c состоит из int2. Теперь, поскольку эти массивы являются постоянными, я хочу, чтобы компилятор выполнял оптимизацию времени компиляции, тем самым эффективно оптимизируя доступ к массиву. Например, для следующих двух функций устройства компилятор разворачивает цикл и оптимизирует прочь массив доступ, заменив его непосредственно со значениями с и ш:

__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] 
    } 
} 

Мой вопрос: как я могу предотвратить декларацию с и ш в каждой функции, доступ к этим переменным и все еще есть оптимизация времени компиляции, которые я хочу? Или указано иначе: существует ли работа для объявления массива векторов по всему миру?

N.B .: Я знаю, что я могу хранить c и w в глобальной или __constant__ памяти, но это не даст мне оптимизацию времени компиляции. Память __constant__ также может стать проблематичной при нерегламентированном доступе.

Answer 1

Я не знаю, если это действительно дает то, что вы ищете с точки зрения оптимизации компиляторов, но литье указатель от ИНТ на 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__ - это не одно и то же. Вы можете исключить объявление const из определений переменных для c и ci, но я предполагаю, что наличие там поможет компилятору достичь желаемого.