sum += texfetch(tex,rowN,colN);
for(colN=tid; colN < SIZE; colN+=threads){
for(rowN=bx; rowN < SIZE; rowN+=blocks){
float sum = 0.0;
unsigned int rowN,colN;
const unsigned int bx = blockIdx.x;
const unsigned int tid = threadIdx.x;
const unsigned int threads = blockDim.x;
const unsigned int blocks = gridDim.x;
Sum_h(int run,FTYPE *g_idata, float *g_odata)
texture<float, 2, cudaReadModeElementType> tex;
Тестовая задача практически неотличима от , но вместо доступа к массиву по индексу используется texfetch:
Все это, как обычно, нуждается в изучении.
В то же время, скорость доступа при неоптимальном паттерне очень маленькая. Для решения этой проблемы (помимо оптимизации паттерна) NVidia CUDA предлагает доступ к памяти как к текстуре. При этом работает двумерное кэширование (оптимизированное под локальный доступ), пиковые скорости должны получаться меньше, а наихудшие варианты - наоборот лучше.
В мы рассматривали чтение из глобальной памяти Geforce 8800 напрямую ("как из массива C"). При этом отсутствует кэширование, но при оптимальной схеме доступа получается (согласно указаниям NVidia) наибольшая производительность.
ЧТ, 03/08/2007 - 20:21 lexa
NVidia 8800GTX: скорость чтения текстур
» NVidia 8800GTX: скорость чтения текстур
(не)структурированные мысли
NVidia 8800GTX: скорость чтения текстур | blog.lexa.ru
Комментариев нет:
Отправить комментарий