Блог

Я пытаюсь отладить проблему с индексом, с которой я сталкиваюсь на моей машине CUDA

 Cuda Machine Info:
  

{1->{Имя-> Tesla C2050, тактовая частота-> 1147000, Вычислительные возможности->2., Перекрытие графического процессора-> 1, Максимальные размеры блока->{1024,1024,64}, Максимальные размеры сетки ->{65535,65535,65535}, Максимальное количество потоков на блок -> 1024, Максимальная общая память на блок -> 49152, Общая постоянная память -> 65536, Размер деформации -> 32, максимальный шаг -> 2147483647, Максимальное количество регистров на блок -> 32768, Выравнивание текстуры->512, Количество многопроцессоров-> 14, Количество ядер-> 448, Время ожидания выполнения-> 0, Встроено-> False, Может отображать память хоста-> True, Режим вычисления-> Default, Ширина Texture1D-> 65536, Ширина Texture2D-> 65536, Высота Texture2D->65535, ширина Texture3D->2048, высота Texture3D->2048, глубина Texture3D-> 2048, Ширина массива Texture2D-> 16384, Высота массива Texture2D-> 16384, Фрагменты массива Texture2D->2048, Выравнивание поверхности->512, Параллельные ядра-> True, ECC включен-> True, Общий объем памяти-> 2817982462},

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

 __global __ void cudaMatExp(
float *matrix1, float *matrixStore, int lengthx, int lengthy, int lengthz){

long UniqueBlockIndex = blockIdx.y * gridDim.x   blockIdx.x;

long index = UniqueBlockIndex * blockDim.z * blockDim.y * blockDim.x  
    threadIdx.z * blockDim.y * blockDim.x   threadIdx.y * blockDim.x  
    threadIdx.x;

if (index < lengthx*lengthy*lengthz) {

matrixStore[index] =  index;

}
}
  

По какой-то причине, как только размер моего 3D-массива становится слишком большим, индексация прекращается.

Я пробовал разные размеры блоков (blockDim.x с помощью blockDim.y с помощью blockDim.z):

8x8x8 дает правильную индексацию только до размера массива 12x12x12

9x9x9 обеспечивает правильную индексацию только до размера массива 14x14x14

10x10x10 дает правильную индексацию только до размера массива 15x15x15

Для размеров, превышающих эти, все разные размеры блоков в конечном итоге снова начинают увеличиваться, но они никогда не достигают значения dim ^ 3-1 (который является максимальным индексом, которого должен достичь поток cuda)

Вот несколько графиков, иллюстрирующих это поведение:

Например: это построение по оси x размера 3D-массива (который равен xxx), а по оси y — максимального числа индексов, которое обрабатывается во время выполнения cuda. Этот конкретный график предназначен для блоков размером 10x10x10.

Вот код (Mathematica) для генерации этого графика, но когда я запускал этот, я использовал размеры блока 1024x1x1:

 CUDAExp = CUDAFunctionLoad[codeexp, "cudaMatExp",
  {{"Float", _,"Input"}, {"Float", _,"Output"},
    _Integer, _Integer, _Integer},
  {1024, 1, 1}]; (*These last three numbers are the block dimensions*)

max = 100; (* the maximum dimension of the 3D array *)
hold = Table[1, {i, 1, max}];
compare = Table[i^3, {i, 1, max}];
Do[
   dim = ii;
   AA  = CUDAMemoryLoad[ConstantArray[1.0, {dim, dim, dim}], Real, 
                                     "TargetPrecision" -> "Single"];
   BB  = CUDAMemoryLoad[ConstantArray[1.0, {dim, dim, dim}], Real, 
                                     "TargetPrecision" -> "Single"];

   hold[[ii]] = Max[Flatten[
                  CUDAMemoryGet[CUDAExp[AA, BB, dim, dim, dim][[1]]]]];

 , {ii, 1, max}]

ListLinePlot[{compare, Flatten[hold]}, PlotRange -> All]
  

Это тот же график, но теперь выводится x ^ 3 для сравнения с тем, где он должен быть. Обратите внимание, что она расходится после того, как размерность массива равна > 32

Я проверяю размеры 3D-массива и смотрю, как далеко зашла индексация, и сравниваю ее с dim ^ 3-1. Например, для dim = 32 максимальный индекс cuda равен 32767 (что равно 32 ^ 3 -1), но для dim = 33 вывод cuda равен 33791, когда он должен быть 35936 (33 ^ 3 -1). Обратите внимание, что 33791-32767 = 1024 = blockDim.x

Вопрос:

Есть ли способ корректно проиндексировать массив с размерами, превышающими размеры блока в Mathematica?

Теперь я знаю, что некоторые люди используют __mul24(threadIdx.y,blockDim.x) в своем индексном уравнении, чтобы предотвратить ошибки при умножении битов, но, похоже, в моем случае это не помогает.

Кроме того, я видел, как кто-то упоминал, что вам следует скомпилировать свой код с помощью -arch=sm_11, потому что по умолчанию он скомпилирован для compute capability 1.0. Я не знаю, имеет ли это место в Mathematica, хотя. Я бы предположил, что CUDAFunctionLoad[] знает, что нужно компилировать с возможностью 2.0. Кто-нибудь знает?

Любые предложения были бы чрезвычайно полезны!

Итак, в Mathematica есть своего рода скрытый способ работы с размерами сетки. Чтобы исправить размер вашей сетки на что-то, что будет работать, вам нужно добавить другое число в конец вызываемой вами функции.

Аргумент обозначает количество потоков для запуска (или размер сетки, умноженный на размер блока).

Например, в моем коде выше:

 CUDAExp = 
  CUDAFunctionLoad[codeexp, 
   "cudaMatExp", {
           {"Float", _, "Input"}, {"Float", _,"Output"}, 
                        _Integer, _Integer, _Integer}, 
     {8, 8, 8}, "ShellOutputFunction" -> Print];
  

(8,8,8) обозначает размер блока.

При вызове CUDAExp[] в mathematica вы можете добавить аргумент, который обозначает количество потоков для запуска:

В этом примере я, наконец, заставил его работать со следующим:

 // AA and BB are 3D arrays of 0 with dimensions dim^3
dim = 64;
CUDAExp[AA, BB, dim, dim, dim, 4089];
  

Обратите внимание, что при компиляции с помощью CUDAFunctionLoad[] ожидается только 5 входных данных, первый — это массив, который вы ему передаете (измерений dim x dim x dim ), а второй — где хранится его память. Третье, четвертое и пятое — это измерения.

Когда вы передаете ему 6-е значение, mathematica переводит это как gridDim.x * blockDim.x , поэтому, поскольку я знаю, что мне нужен gridDim.x = 512 для обработки каждого элемента в массиве, я устанавливаю это число равным 512 * 8 = 4089.

Я надеюсь, что это понятно и полезно кому-нибудь в будущем, кто столкнется с этой проблемой.