Я зарисовка возможное решение здесь.Вам понадобятся некоторые эксперименты, чтобы максимизировать тепловую нагрузку вашего графического процессора.Вообще говоря, перемещение данных энергетически дорого, гораздо больше, чем вычисления в современных процессорах.Так что перетасовка большого количества данных приведет к увеличению энергопотребления.В то же время мы хотим получить дополнительный вклад в энергопотребление от вычислительных единиц.Множители, как правило, являются самыми мощными свиньями;в современных процессорах нам может потребоваться использовать модули FMA (плавное умножение и сложение).
Различные графические процессоры имеют низкую пропускную способность математических операций двойной точности, другие имеют низкую пропускную способность математических операций половинной точности.Поэтому мы бы хотели сосредоточиться на математике с одинарной точностью для вычислительной части нашей нагрузки.Мы хотим иметь возможность легко изменять соотношение вычислений и активности памяти.Один из подходов состоит в том, чтобы использовать развернутую оценку полинома со схемой Хорнера в качестве основного строительного блока, используя POLY_DEPTH шагов.Это мы повторяем REPS раз в цикле.До цикла мы извлекаем исходные данные из глобальной памяти, а после завершения цикла сохраняем результат в глобальной памяти.Изменяя REPS, мы можем экспериментировать с различными настройками баланса вычислений / памяти.
Можно дополнительно поэкспериментировать с параллелизмом на уровне команд, шаблонами данных (поскольку энергопотребление умножителей часто различается в зависимости от битовых шаблонов),и добавление активности PCIe с использованием потоков CUDA для достижения перекрытия выполнения ядра и передачи данных PCIe.Ниже я просто использовал несколько случайных констант в качестве данных множителя.
Очевидно, что мы хотели бы заполнить графический процессор большим количеством потоков.Для этого мы можем использовать довольно маленькое значение THREADS_PER_BLK, дающее нам тонкую гранулярность для заполнения каждого SM.Возможно, мы захотим выбрать количество блоков, кратное количеству SM, чтобы распределить нагрузку как можно более равномерно, или использовать значение MAX_BLOCKS, которое равномерно делит общее число SM.Какой объем памяти источника и получателя мы должны затронуть, будет зависеть от эксперимента: мы можем определить массивы LEN элементов как кратное число блоков.Наконец, мы хотим запустить определенное и настроенное таким образом ядро ITER количество раз, чтобы создать непрерывную загрузку в течение некоторого времени.
Обратите внимание, что при применении нагрузки графический процессор будет нагреваться, и это, в свою очередь, будетдальнейшее увеличение его мощности.Для достижения максимальной тепловой нагрузки необходимо запустить приложение, генерирующее нагрузку, на 5 и более минут.Также обратите внимание, что управление питанием графического процессора может динамически снижать тактовые частоты и напряжения для снижения энергопотребления, и ограничение мощности может сработать, прежде чем вы достигнете температурного предела.В зависимости от графического процессора вы можете установить ограничение мощности выше, чем значение, используемое по умолчанию с утилитой nvidia-smi.
Программа, приведенная ниже, поддерживает Quadro P2000 в режиме ожидания.cap, с загрузкой графического процессора 98% и загрузкой контроллера памяти 83% -86%, как сообщает утилита TechPowerUp GPU-Z.Это, безусловно, потребует корректировки для других графических процессоров.
#include <stdlib.h>
#include <stdio.h>
#define THREADS_PER_BLK (128)
#define MAX_BLOCKS (65520)
#define LEN (MAX_BLOCKS * 1024)
#define POLY_DEPTH (30)
#define REPS (2)
#define ITER (100000)
// Macro to catch CUDA errors in CUDA runtime calls
#define CUDA_SAFE_CALL(call) \
do { \
cudaError_t err = call; \
if (cudaSuccess != err) { \
fprintf (stderr, "Cuda error in file '%s' in line %i : %s.\n",\
__FILE__, __LINE__, cudaGetErrorString(err) ); \
exit(EXIT_FAILURE); \
} \
} while (0)
// Macro to catch CUDA errors in kernel launches
#define CHECK_LAUNCH_ERROR() \
do { \
/* Check synchronous errors, i.e. pre-launch */ \
cudaError_t err = cudaGetLastError(); \
if (cudaSuccess != err) { \
fprintf (stderr, "Cuda error in file '%s' in line %i : %s.\n",\
__FILE__, __LINE__, cudaGetErrorString(err) ); \
exit(EXIT_FAILURE); \
} \
/* Check asynchronous errors, i.e. kernel failed (ULF) */ \
err = cudaDeviceSynchronize(); \
if (cudaSuccess != err) { \
fprintf (stderr, "Cuda error in file '%s' in line %i : %s.\n",\
__FILE__, __LINE__, cudaGetErrorString( err) ); \
exit(EXIT_FAILURE); \
} \
} while (0)
__global__ void burn (const float * __restrict__ src,
float * __restrict__ dst, int len)
{
int stride = gridDim.x * blockDim.x;
int tid = blockDim.x * blockIdx.x + threadIdx.x;
for (int i = tid; i < len; i += stride) {
float p = src[i] + 1.0;
float q = src[i] + 3.0f;
for (int k = 0; k < REPS; k++) {
#pragma unroll POLY_DEPTH
for (int j = 0; j < POLY_DEPTH; j++) {
p = fmaf (p, 0.68073987f, 0.8947237f);
q = fmaf (q, 0.54639739f, 0.9587058f);
}
}
dst[i] = p + q;
}
}
int main (int argc, char *argv[])
{
float *d_a, *d_b;
/* Allocate memory on device */
CUDA_SAFE_CALL (cudaMalloc((void**)&d_a, sizeof(d_a[0]) * LEN));
CUDA_SAFE_CALL (cudaMalloc((void**)&d_b, sizeof(d_b[0]) * LEN));
/* Initialize device memory */
CUDA_SAFE_CALL (cudaMemset(d_a, 0x00, sizeof(d_a[0]) * LEN)); // zero
CUDA_SAFE_CALL (cudaMemset(d_b, 0xff, sizeof(d_b[0]) * LEN)); // NaN
/* Compute execution configuration */
dim3 dimBlock(THREADS_PER_BLK);
int threadBlocks = (LEN + (dimBlock.x - 1)) / dimBlock.x;
if (threadBlocks > MAX_BLOCKS) threadBlocks = MAX_BLOCKS;
dim3 dimGrid(threadBlocks);
printf ("burn: using %d threads per block, %d blocks, %f GB\n",
dimBlock.x, dimGrid.x, 2e-9*LEN*sizeof(d_a[0]));
for (int k = 0; k < ITER; k++) {
burn<<<dimGrid,dimBlock>>>(d_a, d_b, LEN);
CHECK_LAUNCH_ERROR();
}
CUDA_SAFE_CALL (cudaFree(d_a));
CUDA_SAFE_CALL (cudaFree(d_b));
return EXIT_SUCCESS;
}