Есть ли способ явно сопоставить поток с определенным перекосом в CUDA?

Question

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

Например, давайте представим, что у нас есть 1024 потока cuda и размер основы 32. После динамического анализа мы обнаруживаем, что потоки 989, 243, 819, ..., 42 (всего 32 потока в списке) должны находиться в одной и той же основе. Мы определили, что они должны быть в одном и том же перекосе, потому что они практически не расходятся при выполнении кода (возможно, они не обязательно были в одном перекосе при выполнении динамического анализа программы CUDA).

Есть ли способ управления планированием потоков для деформации в CUDA? Если нет, есть ли другой язык программирования GPU, который предложил бы это явное планирование деформации. Если нет, что можно сделать (возможно, даже очень низкоуровневый подход для решения этой проблемы)? Я надеюсь, что есть хотя бы ответ на этот последний вопрос, так как, вероятно, именно так была реализована CUDA - если планирование деформации не выполняется на аппаратном уровне, что было бы неудачно. Спасибо!

Answer 1

Нет, вы не можете выбрать назначение потоков деформациям. Поддержка этого заявления освещена здесь.

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

Одним из основных способов, которым поток CUDA приобретает свой характер, является генерация уникального идентификатора потока, который является типичным образцом для любого кода CUDA, например:

int idx=threadIdx.x+blockDim.x*blockIdx.x;

создает канонический, глобально уникальный 1D-индекс потока для каждого потока.

Но нет особой причины, по которой так должно быть. Я мог бы так же легко сделать:

int private_idx = threadIdx.x+blockDim.x*blockIdx.x;
int idx = desired_idx[private_idx];

и тогда потоки могут быть пронумерованы в любом порядке. Если твой desired_idx массив имел группу чисел, как вы предлагаете:

989, 243, 819, ..., 42

Тогда эти смежные потоки будут вести себя в соответствии с этим порядком.