CUDA Где объявить константу для распределения общей памяти

Question

Я запускаю функцию пригодности для 1024 матриц, каждая матрица получает свой собственный блок и имеет тот же размер. Каждый блок имеет потоки n*n (размерность матрицы) и должен иметь общую память n*n, чтобы я мог легко сократить сумму. Однако размер n для всех матриц является переменным до времени выполнения (т. Е. Его можно вручную изменить, хотя всегда имеет значение 2, поэтому суммирование простое). Проблема здесь в том, что разделяемая память должна выделяться с использованием константы, но мне также нужно передать значение ядру из хоста. Где объявить размер n так, чтобы он был видимым для CPU (для передачи в ядро) и может использоваться для объявления размера общей памяти (в ядре)?

Мой код структурирована следующим образом:

из main.cu я называю ядро:

const int num_states = 1024 
const int dimension = 4 

fitness <<< num_states, dimension * dimension >>> (device_array_of_states, dimension, num_states, device_fitness_return); 

, а затем в kernel.cu у меня есть:

__global__ void fitness(
    int *numbers, 
    int dimension, 
    int num_states, 
    int *fitness_return) { 
    __shared__ int fitness[16]; <<-- needs to be dimension * dimension 
    //code 
} 

numbers представляет собой массив, представляющий 1024 матрицы , dimension - длина строки и столбца, num_states - 1024, fitness_return i s массив с длиной 1024, который содержит значение пригодности для каждой матрицы. В ядре общая память жестко закодирована с квадратом dimension (поэтому dimension - это 4 в этом примере).

Где и как я могу объявить dimension, чтобы его можно было использовать для распределения общей памяти, а также для вызова ядра, поэтому мне нужно только обновить dimension в одном месте? Спасибо за вашу помощь.

Answer 1

Количество выделенной разделяемой памяти равномерно по всем блокам. Вы можете _actually_use_ различного количества разделяемой памяти в каждом блоке, но он все еще доступен. Кроме того, количество разделяемой памяти весьма ограничено независимо, поэтому n * n элементов не может превышать maximum amount of space (обычно 48KiB); для float элементов типа (по 4 байта), что будет означать n < 340 или около того.

Теперь существует два способа распределения общей памяти: статический и динамический.

Статический распределение является то, что вы дали в качестве примера, который не будет работать:

__shared__ int fitness[16]; 

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

С Dynamic общего распределения памяти, вы не указать размер в коде ядра - вы leave it empty и предварять extern:

extern __shared__ int fitness[]; 

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

Но в вашем случае нитки do должны знать, что такое n. Ну, просто передайте его как аргумент ядра.Так, Parallel-for-all блог

__global__ void fitness(
    int *numbers, 
    int dimension, 
    int num_states, 
    int *fitness_return, 
    unsigned short fitness_matrix_order /* that's your n*/) 
{ 
    extern __shared__ int fitness[]; 
    /* ... etc ... */ 
} 

NVidia имеет a nice post с внедрением более углубленным с использованием разделяемой памяти, которая специально охватывает статическое и динамическое распределение разделяемой памяти.