Два главных приоритета оптимизации для любого программиста CUDA - это эффективное использование памяти и предоставление достаточного параллелизма, чтобы скрыть задержки.Мы будем использовать их, чтобы руководствоваться нашим алгоритмическим выбором.
Очень простая стратегия потока (стратегия потока отвечает на вопрос «что каждый поток будет делать или за что будет отвечать?») В любом преобразовании *Проблема типа 1004 * (в отличие от сокращение ) состоит в том, чтобы каждый поток отвечал за 1 выходное значение.Ваша проблема соответствует описанию преобразование - размер выходного набора данных соответствует порядку размера (ов) входных данных.
Я предполагаю, что вы намеревались иметь два равныхвекторы длины, содержащие ваши трехмерные векторы, а также то, что вы хотите взять перекрестное произведение первых трехмерных векторов в каждом и 2-х трехмерных векторов в каждом и т. д.
Если мы выберем стратегию потока с 1 выходомточка на поток (т. е. result[x] или result[y] или result[z], все вместе будет 3 выходными точками), тогда нам понадобится 3 потока для вычисления выходных данных каждого векторного перекрестного произведения.Если у нас будет достаточно векторов для умножения, у нас будет достаточно потоков, чтобы наша машина была «занята» и хорошо справлялась с сокрытием задержки.Как правило, ваша проблема начнет становиться интересной на графических процессорах, если число потоков составляет 10000 или более, поэтому это означает, что мы бы хотели, чтобы ваши одномерные векторы состояли примерно из 3000 трехмерных векторов или более.Давайте предположим, что это так.
Чтобы решить задачу повышения эффективности использования памяти, нашей первой задачей является загрузка ваших векторных данных из глобальной памяти.В идеале мы хотим, чтобы это было слитым , что примерно означает, что смежные потоки обращаются к соседним элементам в памяти.Мы также хотим, чтобы хранилища выходных данных были объединены, и наша стратегия потоков, заключающаяся в выборе одной выходной точки / одного векторного компонента для каждого потока, будет хорошо работать для поддержки этого.
Для эффективного использования памяти мы хотели быидеально загружать каждый элемент из глобальной памяти только один раз.Ваш алгоритм, естественно, включает в себя небольшое количество повторного использования данных .Повторное использование данных очевидно, поскольку вычисление result[y] зависит от vec2[z], а вычисление result[x] также зависит от vec2[z], чтобы выбрать только один пример.Поэтому типичная стратегия при повторном использовании данных состоит в том, чтобы сначала загрузить данные в общую память CUDA, а затем позволить потокам выполнять свои вычисления на основе данных в общей памяти.Как мы увидим, это позволяет нам довольно легко / удобно организовывать объединенные нагрузки из глобальной памяти, поскольку глобальная схема загрузки данных больше не тесно связана с потоками или использованием данных для вычислений.
Последняя задача - определить шаблон индексации, чтобы каждый поток выбирал нужные элементы из общей памяти для умножения вместе.Если мы посмотрим на ваш шаблон расчета, который вы изобразили в своем вопросе, мы увидим, что первая загрузка из vec1 следует шаблону смещения +1 (по модулю 3) из индекса, для которого вычисляется результат.Итак x -> y, y -> z и z -> x.Аналогичным образом мы видим +2 (по модулю 3) для следующей загрузки из vec2, другой шаблон +2 (по модулю 3) для следующей загрузки из vec1 и еще один шаблон +1 (по модулю 3) для окончательной загрузки из vec2.
Если мы объединим все эти идеи, мы сможем написать ядро, которое должно иметь в целом эффективные характеристики:
$ cat t1003.cu
#include <stdio.h>
#define TV1 1
#define TV2 2
const size_t N = 4096; // number of 3D vectors
const int blksize = 192; // choose as multiple of 3 and 32, and less than 1024
typedef float mytype;
//pairwise vector cross product
template <typename T>
__global__ void vcp(const T * __restrict__ vec1, const T * __restrict__ vec2, T * __restrict__ res, const size_t n){
__shared__ T sv1[blksize];
__shared__ T sv2[blksize];
size_t idx = threadIdx.x+blockDim.x*blockIdx.x;
while (idx < 3*n){ // grid-stride loop
// load shared memory using coalesced pattern to global memory
sv1[threadIdx.x] = vec1[idx];
sv2[threadIdx.x] = vec2[idx];
// compute modulo/offset indexing for thread loads of shared data from vec1, vec2
int my_mod = threadIdx.x%3; // costly, but possibly hidden by global load latency
int off1 = my_mod+1;
if (off1 > 2) off1 -= 3;
int off2 = my_mod+2;
if (off2 > 2) off2 -= 3;
__syncthreads();
// each thread loads its computation elements from shared memory
T t1 = sv1[threadIdx.x-my_mod+off1];
T t2 = sv2[threadIdx.x-my_mod+off2];
T t3 = sv1[threadIdx.x-my_mod+off2];
T t4 = sv2[threadIdx.x-my_mod+off1];
// compute result, and store using coalesced pattern, to global memory
res[idx] = t1*t2-t3*t4;
idx += gridDim.x*blockDim.x;} // for grid-stride loop
}
int main(){
mytype *h_v1, *h_v2, *d_v1, *d_v2, *h_res, *d_res;
h_v1 = (mytype *)malloc(N*3*sizeof(mytype));
h_v2 = (mytype *)malloc(N*3*sizeof(mytype));
h_res = (mytype *)malloc(N*3*sizeof(mytype));
cudaMalloc(&d_v1, N*3*sizeof(mytype));
cudaMalloc(&d_v2, N*3*sizeof(mytype));
cudaMalloc(&d_res, N*3*sizeof(mytype));
for (int i = 0; i<N; i++){
h_v1[3*i] = TV1;
h_v1[3*i+1] = 0;
h_v1[3*i+2] = 0;
h_v2[3*i] = 0;
h_v2[3*i+1] = TV2;
h_v2[3*i+2] = 0;
h_res[3*i] = 0;
h_res[3*i+1] = 0;
h_res[3*i+2] = 0;}
cudaMemcpy(d_v1, h_v1, N*3*sizeof(mytype), cudaMemcpyHostToDevice);
cudaMemcpy(d_v2, h_v2, N*3*sizeof(mytype), cudaMemcpyHostToDevice);
vcp<<<(N*3+blksize-1)/blksize, blksize>>>(d_v1, d_v2, d_res, N);
cudaMemcpy(h_res, d_res, N*3*sizeof(mytype), cudaMemcpyDeviceToHost);
// verification
for (int i = 0; i < N; i++) if ((h_res[3*i] != 0) || (h_res[3*i+1] != 0) || (h_res[3*i+2] != TV1*TV2)) { printf("mismatch at %d, was: %f, %f, %f, should be: %f, %f, %f\n", i, h_res[3*i], h_res[3*i+1], h_res[3*i+2], (float)0, (float)0, (float)(TV1*TV2)); return -1;}
printf("%s\n", cudaGetErrorString(cudaGetLastError()));
return 0;
}
$ nvcc t1003.cu -o t1003
$ cuda-memcheck ./t1003
========= CUDA-MEMCHECK
no error
========= ERROR SUMMARY: 0 errors
$
Обратите внимание, что я решил написать ядро с использованием петля шага сетки .Это не очень важно для этого обсуждения и не так важно для этой проблемы, потому что я выбрал размер сетки, равный размеру проблемы (4096 * 3).Однако для гораздо больших размеров задач вы можете выбрать меньший размер сетки, чем общий размер проблемы, для некоторого возможного небольшого повышения эффективности.
Для такой простой задачи, как эта, довольно легко определить «оптимальность».Тем не менее, оптимальный сценарий будет долгим, чтобы загрузить входные данные (только один раз) и записать выходные данные.Если мы рассмотрим более крупную версию тестового кода выше, изменив N на 40960 (и не внося никаких других изменений), то общее количество прочитанных и записанных данных составит 40960 * 3 * 4 * 3 байта.Если мы профилируем этот код, а затем сравниваем с bandwidthTest в качестве прокси-сервера для пиковой достижимой пропускной способности памяти, мы наблюдаем:
$ CUDA_VISIBLE_DEVICES="1" nvprof ./t1003
==27861== NVPROF is profiling process 27861, command: ./t1003
no error
==27861== Profiling application: ./t1003
==27861== Profiling result:
Type Time(%) Time Calls Avg Min Max Name
GPU activities: 65.97% 162.22us 2 81.109us 77.733us 84.485us [CUDA memcpy HtoD]
30.04% 73.860us 1 73.860us 73.860us 73.860us [CUDA memcpy DtoH]
4.00% 9.8240us 1 9.8240us 9.8240us 9.8240us void vcp<float>(float const *, float const *, float*, unsigned long)
API calls: 99.10% 249.79ms 3 83.263ms 6.8890us 249.52ms cudaMalloc
0.46% 1.1518ms 96 11.998us 374ns 454.09us cuDeviceGetAttribute
0.25% 640.18us 3 213.39us 186.99us 229.86us cudaMemcpy
0.10% 255.00us 1 255.00us 255.00us 255.00us cuDeviceTotalMem
0.05% 133.16us 1 133.16us 133.16us 133.16us cuDeviceGetName
0.03% 71.903us 1 71.903us 71.903us 71.903us cudaLaunchKernel
0.01% 15.156us 1 15.156us 15.156us 15.156us cuDeviceGetPCIBusId
0.00% 7.0920us 3 2.3640us 711ns 4.6520us cuDeviceGetCount
0.00% 2.7780us 2 1.3890us 612ns 2.1660us cuDeviceGet
0.00% 1.9670us 1 1.9670us 1.9670us 1.9670us cudaGetLastError
0.00% 361ns 1 361ns 361ns 361ns cudaGetErrorString
$ CUDA_VISIBLE_DEVICES="1" /usr/local/cuda/samples/bin/x86_64/linux/release/bandwidthTest
[CUDA Bandwidth Test] - Starting...
Running on...
Device 0: Tesla K20Xm
Quick Mode
Host to Device Bandwidth, 1 Device(s)
PINNED Memory Transfers
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 6375.8
Device to Host Bandwidth, 1 Device(s)
PINNED Memory Transfers
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 6554.3
Device to Device Bandwidth, 1 Device(s)
PINNED Memory Transfers
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 171220.3
Result = PASS
NOTE: The CUDA Samples are not meant for performance measurements. Results may vary when GPU Boost is enabled.
$
Ядру требуется 9,8240us для выполнения, и в это время загружается или сохраняется общее количество40960 * 3 * 4 * 3 байта данных.Поэтому достигнутая ядром пропускная способность памяти составляет 40960 * 3 * 4 * 3 / 0,000009824 или 150 ГБ / с.Измерение прокси для пика, достижимого на этом графическом процессоре, составляет 171 ГБ / с, поэтому это ядро достигает 88% от оптимальной пропускной способности.Благодаря более тщательному бенчмаркингу для запуска ядра два раза подряд, для 2-го выполнения требуется всего 8,99us.Это обеспечивает достижимую полосу пропускания в этом случае до 96% максимальной достижимой пропускной способности.