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

Question

Я не опытный программист CUDA. У меня такая проблема. Я пытаюсь загрузить плиту (32x32) большой матрицы (10K * 10K) из глобальной памяти в разделяемую память, и я синхронизирую ее, когда это происходит. Я понял, что если я загружу его в частную память (регистры), он загружается в 4-5 раз быстрее, чем загрузка общей памяти.

__global__ void speedtest(float *vel,int nx) { 

int globalx = blockDim.x * blockIdx.x + threadIdx.x+pad; 
int globalz = blockDim.y * blockIdx.y + threadIdx.y+pad; 
int localx=threadIdx.x; 
int localz=threadIdx.y; 

float ptest; 
__shared__ float stest[tile][tile]; 

//stest[localz][localx]=vel[globalz*nx+globalx]; //load to shared memory 
ptest=vel[globalz*nx+globalx]; //load to private memory 
__syncthreads(); 
} 

Я закомментировать STest и Ptest по одному и рассчитать время, прошедшее с cudaeventrecord. stest взял 3,2 мс и ptest взял 0,75 м для загрузки. Что я делаю не так? Сроки должны быть очень похожими? Что мне не хватает?

Конфигурация: Cuda 7.5, gtx 980, только 32-битные переменные и вычисления, нет конкретной цели, я просто играю с ней.

Я вывешиваю пример кода в соответствии с просьбой

#include<stdio.h> 
#include <math.h> 
#define tile 32 
#include <helper_cuda.h> 
void makeittwo(float *array,int nz,int nx) 
{ 
//this just assigns a number into the vector 
int n2; 
n2=nx*nz; 
for (int i=0;i<n2;i++) 
array[i]=2000; 
} 
__global__ void speedtest(float *vel,int nx,int nz) { 

int globalx = blockDim.x * blockIdx.x + threadIdx.x; 
int globalz = blockDim.y * blockIdx.y + threadIdx.y; 
int localx=threadIdx.x; 
int localz=threadIdx.y; 

float ptest; //declarations 
__shared__ float stest[tile][tile]; 

if (globalx<nx && globalz<nz){ 
stest[localz][localx]=vel[globalz*nx+globalx]; //shared variable 
//ptest=vel[globalz*nx+globalx];      //private variable 

//comment out ptest and stest one by one to test them 
} 
__syncthreads(); 

}  

int main(int argc,char *argv) 
{ 
int nx,nz,N; 
float *vel; 

nz=10000;nx=10000; //matrix dimensions 
N=nz*nx; //convert matrix into vector 

checkCudaErrors(cudaMallocHost(&vel,sizeof(float)*N)); //using pinned memory 
makeittwo(vel,nz,nx); 

dim3 dimBlock(tile,tile); 
dim3 dimGrid; 

int blockx=dimBlock.x; 
int blockz=dimBlock.y; 

dimGrid.x = (nx + blockx - 1)/(blockx); 
dimGrid.y = (nz + blockz - 1)/(blockz); 

float *d_vel; 
checkCudaErrors(cudaMalloc(&d_vel,sizeof(float)*(N))); //copying to device 
checkCudaErrors(cudaMemcpy(d_vel, vel, sizeof(float)*(N), cudaMemcpyHostToDevice)); 

cudaEvent_t start,stop; 
float elapsedTime; 

cudaEventCreate(&start); 
cudaEventCreate(&stop); 
cudaEventRecord(start,0); 
speedtest<<<dimGrid,dimBlock>>>(d_vel,nx,nz); //calling the function 
cudaEventRecord(stop,0); 
cudaEventSynchronize(stop); 
cudaEventElapsedTime(&elapsedTime,start,stop); 

printf("time=%3.3f ms\n",elapsedTime); 
checkCudaErrors(cudaMemcpy(vel, d_vel, sizeof(float)*N, cudaMemcpyDeviceToHost)); 
//calling the matrix back to check if all went well (this fails if out of bound calls are made) 

cudaDeviceReset(); 

} 

Answer 1

Пример кода на самом деле не измеряет то, что ОП рассчитывает измерить, потому что некоторые инструкции оптимизированы прочь компилятором.

В переменной локальная переменная пример (ptest) нагрузка не влияет на состояние вне ядра. В этом случае компилятор может полностью удалить инструкцию. Это можно увидеть в коде SASS. Код SASS тот же, когда ptest=vel[globalz*nx+globalx]; активен, или оба оператора (ptest и stest) удаляются. Чтобы проверить код SASS, вы можете запустить cuobjdump --dump-sass в объектном файле.

По-видимому, инструкции не оптимизированы в общей памяти пример, который можно проверить в коде SASS. (На самом деле, я ожидал бы инструкции будут удалены, а также. Есть ли побочные эффекты, что мисс?)

Как уже говорилось в комментариях, с помощью простого расчета (ptest*=ptest) и записи в глобальной памяти компилятор не может удалить инструкцию, поскольку он изменяет глобальное состояние.

Из комментариев ОП я предполагаю, что существует недоразумение в том, как работает операция загрузки в общую память. Фактически данные загружаются из глобальной памяти в регистры, а затем сохраняются в общей памяти. В (соответствующая) инструкция SASS (для sm_30), которые генерируются выглядеть следующим образом

LD.E R2, [R6]; // load to register R2 
STS [R0], R2; // store from register R2 to shared memory 

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

stest[localz][localx]=vel[globalz*nx+globalx]; // load to shared memory 
stest[localz][localx]*=stest[localz][localx]; // multiply 
vel[globalz*nx+globalx]=stest[localz][localx]; // save to global memory 

код SASS показывает, что переменная хранится только в общей памяти после вычисления (и никогда не читают формы общей памяти).

LD.E R2, [R6]; // load to register 
FMUL R0, R2, R2; // multiply 
STS [R3], R0; // store the result in shared memory 
ST.E [R6], R0; // store the result in global memory 

Я действительно не специалист в SASS кода, пожалуйста, поправьте меня, если я ошибаюсь, или оставить что-нибудь важное.