Почему переменные-члены не могут быть общими?

Question

Почему переменные-члены не могут быть общими?

Я хотел бы создать экземпляр класса в коде CUDA, который разделяет некоторые его члены с другими потоками в том же блоке.

Однако при попытке скомпилировать следующий код я получаю сообщение об ошибке: »атрибут« общий доступ »здесь не применяется« (nvcc версия 4.2).

class SharedSomething {

public:
__shared__ int i; // this is not allowed
};

__global__ void run() {

SharedSomething something;
}

В чем причина этого? Есть ли обходной путь для достижения желаемого поведения (общие члены класса в одном блоке)?

6

c++cuda

Решение

Другие решения

Объекты отмечены как __shared__ находиться в разделяемой памяти, которая выделена на блок потока. Он имеет ограниченный размер и тот же срок службы, что и блок потока.

Так вот почему вы не можете объявить членов класса как общие — их время жизни определяется не экземпляром класса, а блоком потока. возможно static ученики могли делиться, но не проверяли.

Увидеть Руководство по программированию CUDA для деталей, раздел B.2.3.

7

Источник

Accepted Answer

Рост объяснил причину ограничения. Чтобы ответить на вторую часть вопроса, простой обходной путь — заставить ядро объявить разделяемую память и инициализировать указатель на нее, принадлежащий классу, например, в конструкторе класса. Пример.

class Foo
{
public:
__device__
Foo(int *sPtr) : sharedPointer(sPtr, gPtr) {
sharedPointer[threadIdx.x] = gPtr[blockIdx.x * blockDim.x + threadIdx.x];
__syncthreads();
}

__device__
void useSharedData() { printf("my data: %f\n", sharedPointer[threadIdx.x]); }

private:
int *sharedPointer;
};

__global__ void example(int *gData)
{
__shared__ int sData[BLOCKDIM];

Foo f(sData, gData);

f.useSharedData();
}

Предостережение: код, написанный в браузере, непроверенный, непроверенный (и это тривиальный пример, но концепция распространяется на реальный код — я сам использовал эту технику).

6