Компьютерная графика

Информация о пользователе

Привет, Гость! Войдите или зарегистрируйтесь.



Тутор по CUDA

Сообщений 1 страница 29 из 29

1

Прочитал тутор http://steps3d.narod.ru/tutorials/cuda-tutorial.html и не смог удержаться от нескольких замечаний.
Статья конечно полезная, особенно в отношении описания особенности работы памяти на cuda, но в целом слишком краткая. Её следовало бы дополнить следующим материалом:

1. Интеграция c++ и cuda программ.
2. Архитектура приложения.
3. Особенности использования графических апи, особенно использование vbo в cuda программах.
4. Работа с текстурной памятью в gpgpu программах.
5. Использование шаблонов в девайс функциях и структурах, метапрограммирование на cuda.

0

2

Это только одна часть. К ней есть продолжение, возможно будут следующие статьи, по крайней мере работу с текстурной памятью и OpenGL лучше все-таки вынести в отдельную статью, ане добавлять в эту.
А что подразумевается под п2 ?

0

3

>Это только одна часть.

Я бы сказал вводная.

>К ней есть продолжение, возможно будут следующие статьи, по крайней мере работу с текстурной памятью

Это нужно. Даже весьма.

>и OpenGL лучше все-таки вынести в отдельную статью, ане добавлять в эту.

OpenGL есть куча статей, но интересно именно взаимодействие с cuda.

>А что подразумевается под п2 ?

В основном совместная работа c++, cuda и какого-нибудь графического апи.

Я прослушал аудиолекции читаные на мгу и могу выразить общие впечатления.
Первые три лекции для меня (наверное не только для меня) оказались самыми полезными, там было детально разжевана работа с памятью.
Последующие же лекции, хоть они и затрагивали вопросы мне интересные, но к сожалению без видеокартинки следить за ними очень трудно (когда лектор говорит: вот эта строка - тут я пас, какая именно строка, могу только догадываться). Кроме того, последние лекции это не о куде, а скорее рассказ авторов, о своей работе на куде. Что совсем не одно и то же.

Кстати с замечанием, что работать со встроеными типами cuda: float2, float3 и т.д. ручками, покомпонентно, я никак согласится не могу. Cuda поддерживает перегрузку операторов, в том числе и на девайсе.
Вот например моя векторная билеотека для работы со стандартными типами Cuda.

Код:
#ifndef _VECTOR_CUH_
#define _VECTOR_CUH_

#include <vector_types.h>	//ñòàíäàðòíûå âåêòîðíûå òèïû CUDA

template <typename T>
__device__ T swap(T &a, T &b) 
{
	T temp;
	temp = a;
	a = b;
	b = temp;
}

//Øàáëîíûå ïðåîáðàçîâàíèÿ ìàññèâà float â ñòàíäàðòíûå âåêòîðíûå òèïû CUDA
template <typename T>
__device__ T makeFromAr(const float *v) 
{
	T res;
	return res;
}

template <>
__device__ float4 makeFromAr<float4>(const float *v) 
{return make_float4(*v, *(v+1), *(v+2), *(v+2));}

template <>
__device__ float3 makeFromAr<float3>(const float *v) 
{return make_float3(*v, *(v+1), *(v+2));}

template <>
__device__ float2 makeFromAr<float2>(const float *v) 
{return make_float2(*v, *(v+1));}

//Îïðåäåëåíèå àðèôìåòè÷åñêèõ îïåðàòîðîâ:

//Âåêòîðíîå ïðîèçâåäåíèå:
__device__ float3 operator ^(const float3 &a, const float3 &b) 
{return make_float3(a.y*b.z-a.z*b.y, a.z*b.x-a.x*b.z, a.x*b.y-a.y*b.x);}

//Âåêòîðíîå ïðîèçâåäåíèå (äëÿ 2D ïîíèìàåòñÿ êàê äëÿ ïëîñêîïàðàëåëüíîé ñèñòåìû):
__device__ float operator ^(const float2 &a, const float2 &b)
{return a.x*b.y-a.y*b.x;}

//âåêòîðíîå óìíîæàíèå âåêòîðà arg ëåæàùåãî íà ïëîñêîñòè XY íà îðò k (îñè Z)
__device__ float2 CrossProductOn_k(float2 arg)  
{
	return make_float2(arg.y, -arg.x);	
};

//Íîðìà:
__device__ float Norm(const float4 &a) 
{return (fabs(a.x)+fabs(a.y)+fabs(a.z)+fabs(a.w))*.25f;}

__device__ float Norm(const float3 &a) 
{return (fabs(a.x)+fabs(a.y)+fabs(a.z))*.33333333333f;}

__device__ float Norm(const float2 &a) 
{return (fabs(a.x)+fabs(a.y))*.5f;}

//Äëèíà:
__device__ float Length(const float4 &a) 
{return sqrtf(a.x*a.x+a.y*a.y+a.z*a.z+a.w*a.w);}

__device__ float Length(const float3 &a) 
{return sqrtf(a.x*a.x+a.y*a.y+a.z*a.z);}

__device__ float Length(const float2 &a) 
{return sqrtf(a.x*a.x+a.y*a.y);}

//Êâàäðàò äëèíû:
__device__ float Length2(const float4 &a) 
{return a.x*a.x+a.y*a.y+a.z*a.z+a.w*a.w;}

__device__ float Length2(const float3 &a) 
{return a.x*a.x+a.y*a.y+a.z*a.z;}

__device__ float Length2(const float2 &a) 
{return a.x*a.x+a.y*a.y;}

//Ñóììà:
__device__ float4 operator+(const float4 &a, const float4 &b) 
{return make_float4(a.x+b.x, a.y+b.y, a.z+b.z, a.w+b.w);}

__device__ float3 operator+(const float3 &a, const float3 &b) 
{return make_float3(a.x+b.x, a.y+b.y, a.z+b.z);}

__device__ float2 operator+(const float2 &a, const float2 &b) 
{return make_float2(a.x+b.x, a.y+b.y);}

//Ðàçíîñòü:
__device__ float4 operator-(const float4 &a, const float4 &b) 
{return make_float4(a.x-b.x, a.y-b.y, a.z-b.z, a.w-b.w);}

__device__ float3 operator-(const float3 &a, const float3 &b) 
{return make_float3(a.x-b.x, a.y-b.y, a.z-b.z);}

__device__ float2 operator-(const float2 &a, const float2 &b) 
{return make_float2(a.x-b.x, a.y-b.y);}

//Ñêàëÿðíîå ïðîèçâåäåíèå:
__device__ float operator|(const float4 &a, const float4 &b) 
{return a.x*b.x+a.y*b.y+a.z*b.z+a.w*b.w;}

__device__ float operator|(const float3 &a, const float3 &b) 
{return a.x*b.x+a.y*b.y+a.z*b.z;}

__device__ float operator|(const float2 &a, const float2 &b) 
{return a.x*b.x+a.y*b.y;}

//Ïî÷ëåííîå ïðîèçâåäåíèå:
__device__ float4 operator%(const float4 &a, const float4 &b) 
{return make_float4(a.x*b.x, a.y*b.y, a.z*b.z, a.w*b.w);}

__device__ float3 operator%(const float3 &a, const float3 &b) 
{return make_float3(a.x*b.x, a.y*b.y, a.z*b.z);}

__device__ float2 operator%(const float2 &a, const float2 &b) 
{return make_float2(a.x*b.x, a.y*b.y);}

//Ïðîèçâåäåíèå íà ñêàëÿð:
__device__ float4 operator*(const float4 &a, const float &b) 
{return make_float4(a.x*b, a.y*b, a.z*b, a.w*b);}

__device__ float3 operator*(const float3 &a, const float &b) 
{return make_float3(a.x*b, a.y*b, a.z*b);}

__device__ float2 operator*(const float2 &a, const float &b) 
{return make_float2(a.x*b, a.y*b);}

__device__ float4 operator*(const float &a, const float4 &b) 
{return make_float4(b.x*a, b.y*a, b.z*a, b.w*a);}

__device__ float3 operator*(const float &a, const float3 &b) 
{return make_float3(b.x*a, b.y*a, b.z*a);}

__device__ float2 operator*(const float &a, const float2 &b) 
{return make_float2(b.x*a, b.y*a);}

//Äåëåíèå íà ñêàëÿð:
template <typename T>
__device__ T operator/(const T &a, const float &b) 
{return (a*(1.0/b));}

//Íîðìàëèçàöèÿ:
template <typename T>
__device__ T normalize(const T &a) 
{return (a/Length(a));}

//typedef float2 Vector2f;

#endif // #ifndef _VECTOR_CUH_

Эта библитека работает в одном треде, но возможно многотредная библиотка окажется выгодней. По затратам на память наверняка, но возможно и по скорости.

0

4

Про OpenGL  я имел в виду именно взаимодействие с CUDA в виде отдельной статьи, туда же работу с текстурами
Про перегшрузку операторов для встроенных типов - я писал о том, что встроенной перегрузки нет.
Да можно сделать ее самому, как сделали Вы - это я думаю очень удобный шаг, когда много именно векторных операций.

0

5

OK, но я не уверен, что сделал оптимально.
Мне собственно непонятно как лучше перегрузить данные из глобальной памяти в шеред, если предпологается их дальше использовать как встроенные типы (float2, float3 и т.д), конкретно: что там будет с конфликтом банков?
Скажем вот для такого кода

Код:
	template <unsigned int blockSize>
	__global__ void kernelTest(float3* f, const int n)
	{
    int ti = threadIdx.x;
    int i =  blockSize * blockIdx.x + threadIdx.x;
    __shared__ float3 point[blockSize];
    
    point[i] = f[ti];
    __syncthreads(); 
	};

	const int blockSize = 256;
	int nPoint = 200;
	float3 *point3;
	cudaMalloc((void**) &point3, nPoint*sizeof(float));
	
	dim3 block(blockSize);
    dim3 grid(nPoint / block.x);
    kernelTest<blockSize><<< grid, block>>>(point3, nPoint);

Здесь у меня есть подозрение, что строка

point[i] = f[ti];

эквивалентна вот такой

point[i] = make_float3(f[ti].x, f[ti].y, f[ti].z);

и даже такой

point[i].x = f[ti].x;  point[i].y = f[ti].y; point[i].z = f[ti].z;

0

6

Тут гораздо важнее не конфликт банков, а обеспечение coalescing при чтении из глобальной памяти. И в этом слуаче самое потимально - явно читать по одномму float'у за раз, т.е.

point[i].x = f[ti].x;  point[i].y = f[ti].y; point[i].z = f[ti].z;

В этом случае будет и coalescing и не будет конфликтов по банкам.
А вот

point[i] = f[ti];

Запросто может убить и то и другое.
Вообще одно из правил оптимизации - предпочтении набора массивов массиву структур. Поскольку в первом случае coalescing получается легко и просто
А выйгрышь от coalescing намного выше чем может дать разруливание конфликтов банков

0

7

Т.е получается, что cuda встроеные типы рассматривает как структуры?
Примерно, вместо float3 структуру struct point3{float x,y,z;}; ?
Но с чтением (отдельных ячеек с sizeof = 4) в тредах, это интересно. Попробую реализовать.

0

8

А как еще их рассматривать - что-то (структур)а заданного размера и при обращении читать всю сущность целиком.

0

9

Накидал решение, чтоб перезапись в шеред из float4 происходило в отдельных тредах

Код:
union wrap
{
	float4 v;
	float a[4];
};

template <unsigned int blockSize>
__global__ void kernelTestFloat4Read(wrap *data, const int n)
{
	int ti = threadIdx.x/sizeof(wrap);
	int i_of_struct = threadIdx.x % sizeof(wrap);
	int i =  blockSize * blockIdx.x + ti;
	__shared__ wrap pos[blockSize];
	pos[ti].a[i_of_struct] = data[i].a[i_of_struct];
	__syncthreads(); 
};

Думаю, что можно обойтись без wrapа, но пока ничего лучшего придумать не могу.
Это решение для предотвращения coalescing я рассматриваю как конкурирующее к решению со структурой массивов. Возможно, что оно окажется даже быстрее, особенно при чтении из кешированной памяти (тупо нужно меньше кэшев).

0

10

Сейчас делал себе рыбу из http://steps3d.narod.ru/downloads/cuda-src.zip и заметил ошибку: вместо delete[] везде написано delete.
В статье тоже самое. Пустяки конечно, но нужно исправить. Текст все-таки учебный.

0

11

А разве совместно с OpenGL не лучше ли использовать OpenCL вместо CUDA? Ведь у них очень приближённо смысл один и тот же.

0

12

Насколько я знаю, по OpenCL есть пока только документация. CUDA же это отлаженая и работающая система.
Кроме того, если нвидия и будет делать OpenCL то только на базе своих процессоров G80/92/200, а это означает, что производительность выше CUDA все равно не будет.
OpenGL прекрасно интегрирована с  CUDA и там можно даже сделать нечто, напоминающее геометрический шейдер.

0

13

Наладил код для сравнения скоростей загрузки структур по элементам в тредах или в одном треде целиком:

Код:
#include <stdio.h>
#include <cutil.h>
#include <Ryba_kernel.cu>  

void determineTheValue(float *data, int &n)
{  
	for ( int i = 0; i < n; i++ ) data[i] = 1;
}

double checkTheValue(float *data, int &n)
{   	
	double sum = 0;
	for ( int i = 0; i < n; i++ )
    sum += double(data[i]);
	return sum;
}

template <typename T>
void test (bool inThread)
{
	printf("struct lenght %i \n", sizeof (T));
	int n = N_BLOCK*N_BLOCK*BLOCK_SIZE/(sizeof (T)/4);
	int	numBytes = n * sizeof (T);
	int	numFloats = numBytes / sizeof (float);

    float *data = new float[numFloats]; 

	determineTheValue(data, numFloats);

	T *data_dev;                               
    cudaMalloc((void**)&data_dev, numBytes);
	cudaMemcpy(data_dev, data, numBytes, cudaMemcpyHostToDevice);

    unsigned int timer = 0;
    CUT_SAFE_CALL(cutCreateTimer(&timer));
	cudaThreadSynchronize();
	CUT_SAFE_CALL(cutStartTimer(timer));

	dim3 blocks  (N_BLOCK, N_BLOCK);
	if (inThread)
	{
    printf("Processing in separate threads\n", numBytes);
    dim3 threads (BLOCK_SIZE);
    kernelReadInThreads<T><<<blocks, threads>>> (data_dev, n);
	}
	else
	{
    printf("Processing in one thread\n", numBytes);
    dim3 threads (BLOCK_SIZE/(sizeof(T)/4));
    kernelRead<T><<<blocks, threads>>> (data_dev, n);    
	}
	cudaThreadSynchronize();
    CUT_SAFE_CALL(cutStopTimer(timer));	
  	
    cudaMemcpy(data, data_dev, numBytes, cudaMemcpyDeviceToHost);

	printf("Array size: %i byte\n", numBytes);
    printf("Processing time: %f (ms) \n", cutGetTimerValue(timer));
	printf("check the value: %f  \n", checkTheValue(data, numFloats));

	CUT_SAFE_CALL(cutDeleteTimer(timer));
    cudaFree(data_dev);      
    delete[] data;
}

int main (int argc, char *  argv [])
{
	for ( int inThread = 0; inThread < 2; ++inThread )
	{
    test<float>(inThread);
    test<float2>(inThread);
    test<float3>(inThread);
    test<float4>(inThread); 
    test<myStruct8float>(inThread); 
    test<myStruct2float4>(inThread);    
	}
    return 0;
}

Функции ядра и девайса

Код:
//**********  file: Ryba_kernel.cu   ******************

#ifndef _RYBA_KERNEL_H_
#define _RYBA_KERNEL_H_

#define MULTKOEFF 2
#define BLOCK_SIZE 	512/(MULTKOEFF*MULTKOEFF)
#define	N_BLOCK    512*MULTKOEFF

struct myStruct8float
{
	float x,y,z,w,t,k,l,i;
};

struct myStruct2float4
{
	float x,y,z,w;
	float4 M;	
};

template <typename T>
__device__ void readInShared(const T *dataBlock, T *dataS)
{
	int ti = threadIdx.x/(sizeof(T)/4);
	int i_of_struct = threadIdx.x % (sizeof(T)/4);
	*(&dataS[ti].x +i_of_struct) = *(&(dataBlock+ti)->x +i_of_struct);
}

template <>
__device__ void readInShared<float>(const float *dataBlock, float *dataS)
{
	dataS[threadIdx.x] = *(dataBlock+threadIdx.x);
}

template <typename T>
__device__ void addWithShered(T *dataBlock, T *dataS)
{
	int ti = threadIdx.x/(sizeof(T)/4);
	int i_of_struct = threadIdx.x % (sizeof(T)/4);
	*(&(dataBlock+ti)->x +i_of_struct) += (*(&dataS[ti].x +i_of_struct))*3.0f;
}

template <>
__device__ void addWithShered<float>(float *dataBlock, float *dataS)
{
	*(dataBlock+threadIdx.x) += dataS[threadIdx.x]*3.0f;
} 

template <typename T>
__device__ void addWithShered(T *data, T *dataS, const int ib){}

template <>
__device__ void addWithShered<myStruct8float>
        	(myStruct8float *data, myStruct8float *dataS, const int ib)
{
	myStruct8float temp = data[ib]; 
	temp.x += dataS[threadIdx.x].x*3.0f;
	temp.y += dataS[threadIdx.x].y*3.0f;
	temp.z += dataS[threadIdx.x].z*3.0f;
	temp.w += dataS[threadIdx.x].w*3.0f;
	temp.t += dataS[threadIdx.x].t*3.0f;
	temp.k += dataS[threadIdx.x].k*3.0f;
	temp.l += dataS[threadIdx.x].l*3.0f;
	temp.i += dataS[threadIdx.x].i*3.0f;
	data[ib] = temp; 
}

template <>
__device__ void addWithShered<myStruct2float4>
        	(myStruct2float4 *data, myStruct2float4 *dataS, const int ib)
{
	float4 temp = make_float4(data[ib].x, data[ib].y, data[ib].z, data[ib].w); 
	float4 tempM = data[ib].M;

	temp.x += dataS[threadIdx.x].x*3.0f;
	temp.y += dataS[threadIdx.x].y*3.0f;
	temp.z += dataS[threadIdx.x].z*3.0f;
	temp.w += dataS[threadIdx.x].w*3.0f;

	tempM.x += dataS[threadIdx.x].M.x*3.0f;
	tempM.y += dataS[threadIdx.x].M.y*3.0f;
	tempM.z += dataS[threadIdx.x].M.z*3.0f;
	tempM.w += dataS[threadIdx.x].M.w*3.0f;

	data[ib].x = temp.x; data[ib].y = temp.y; data[ib].z = temp.z; data[ib].w = temp.w; 
	data[ib].M = tempM; 
}

template <>
__device__ void addWithShered<float4>(float4 *data, float4 *dataS, const int ib)
{
	float4 temp = data[ib];
	temp.x += dataS[threadIdx.x].x*3.0f;
	temp.y += dataS[threadIdx.x].y*3.0f;
	temp.z += dataS[threadIdx.x].z*3.0f;
	temp.w += dataS[threadIdx.x].w*3.0f;
	data[ib] = temp; 
}

template <>
__device__ void addWithShered<float3>(float3 *data, float3 *dataS, const int ib)
{
	float3 temp = data[ib];
	temp.x += dataS[threadIdx.x].x*3.0f;
	temp.y += dataS[threadIdx.x].y*3.0f;
	temp.z += dataS[threadIdx.x].z*3.0f;
	data[ib] = temp; 
}

template <>
__device__ void addWithShered<float2>(float2 *data, float2 *dataS, const int ib)
{
	float2 temp = data[ib];
	temp.x += dataS[threadIdx.x].x*3.0f;
	temp.y += dataS[threadIdx.x].y*3.0f;
	data[ib] = temp; 
}

template <>
__device__ void addWithShered<float>(float *data, float *dataS, const int ib)
{   	
	data[ib] += dataS[threadIdx.x]*3.0f;
}

template <typename T>
__global__ void kernelReadInThreads(T *data, const int n)
{
	int ib =  (N_BLOCK * blockIdx.y + blockIdx.x)*BLOCK_SIZE/(sizeof(T)/4);
	__shared__ T pos[BLOCK_SIZE/(sizeof(T)/4)];
	readInShared(data+ib, pos);
	__syncthreads();
	addWithShered(data+ib, pos);
	__syncthreads();
};

template <typename T>
__global__ void kernelRead(T *data, const int n)
{
	int ti = threadIdx.x;	
	int ib =  (N_BLOCK * blockIdx.y + blockIdx.x)*BLOCK_SIZE/(sizeof(T)/4) + ti;
	__shared__ T pos[BLOCK_SIZE];
	pos[ti] = data[ib];    
	__syncthreads();
	addWithShered(data, pos, ib);
	__syncthreads();
}; 

#endif // #ifndef _RYBA_KERNEL_H_

0

14

Эта программа
1. Определяет для флоат элементов data = 1
2. Перезаписыввет data в шеред.
3. Вычисляет для всех data += шеред*3

Результат такой:

Код:
struct lenght 4
Processing in one thread
flrray size: 536870912 byte
Processing tine: 35.177002 <ns>
check the value: 536870912.000000
struct lenght 8
Processing in one thread
flrray size: 536870912 byte
Processing tine: 33.363297 <ns>
check the value: 536870912.000000
struct lenght 12
Processing in one thread
flrray size: 53687B904 byte
Processing tine: 476.622314 (ns)
check the value: 530579454.000000
struct lenght 16
Processing in one thread
flrray size: 536870912 byte
Processing tine: 36.444752 <ns>
check the value: 536870912.000000
struct lenght 32
Processing in one thread
flrray size: 536870912 byte
Processing tine: 506.948730 (ns)
check the value: 536870912.000000
struct lenght 32
Processing in one thread
flrray size: 53687B912 byte
Processing tine: 88.071953 <ns>
check the value: 536870912.000000
struct lenght 4
Processing in separate threads
flrray size: 536870912 byte
Processing tine: 35.399235 <ns>
check the value: 536870912.000000
struct lenght 8
Processing in separate threads
flrray size: 536870912 byte
Processing tine: 35.645550 <ns>
check the value: 536870912.000000
struct lenght 12
Processing in separate threads
flrray size: 536870904 byte
Processing tine: 323.152191 (ns)
check the value: 536517240.000000
struct lenght 16
Processing in separate threads
flrray size: 536870912 byte
Processing tine: 35.643211 <ns>
check the value: 536870912.000090 
struct lenght 32 
Processing in separate threads 
flrray size: 53687B912 byte 
Processing tine: 35.644894 <ns> 
check the value: 536870912.000000 
struct lenght 32 
Processing in separate threads 
flrray size: 536870912 byte 
Processing tine: 35.784363 <ns) 
check the value: 536870912.000000

Т.е. получается, что
1. Обмен данных сепаратно по тредам дает такую же скорость, как и обмен с массивом float (это даже несмотря на оверхед на вычисление индексов).
2. Обмен данных со стандартными типами float2, float4 как сепаратно элементов по тредам, так и целиком в одном треде дает такую же скорость, как и обмен с массивом float.
3. Невыровненые структуры, включая и стандартный float3 дают провал производительности более чем на порядок.

Отсюда выводы:
Хотя в целом предпочтение структуры массивов массиву структур сохраняет свою силу, но при необходимости можно и нужно использовать массивы структур, особенно стандартных.
В целом массивы нужно организовать таким образом, чтоб они хранили только информацию необходимую и достаточную (ни в коем случае не больше!!) для работы вызываемой функции ядра.

0

15

Уважаемый О. Федор, ты случайно не знаешь как задействовать CUDA совместно с Delphi. Очень надеюсь, что это возможно.

0

16

В документации есть описание как использовать библиотеки CUDA из фортрана. Думаю что можно сделать аналогично.
Но я бы этого не делал. Лучше работать в c++. В крайнем случае, если есть большой законченый проект не дельфях, можно вычислительные модули сделанные на CUDA выделить как DLL.
Если же проекту далеко до завершения, лучше сразу переходить на c++.

0

17

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

0

18

Жаль нельзя Delphi использовать... Ну да ладно.

Я тем временем нашёл опечатку в статье:"Сами программы пишутся на "расширенном" С, при этом их параллельная часть (ядра) выполняются на GPU, а обычная часть - на CPU."
нужно было написать так: ...их параллельная часть (ядра) выполняЕтся... И непонятно ещё что означает слово "(ядра)" и зачем оно взято в скобки.

А вот ещё ляпус (сразу за предыдущим предложением): "Это разделение, также как и управление запуском частей на GPU занимается CUDA.".
можно так писать: Разделение частей и управление их запуском CUDA осуществляет автоматически.

Отредактировано DungeonLords (2009-05-16 11:13:18)

0

19

Спасибо, исправлю

0

20

в статье написал(а):

при этом их параллельная часть (ядра) выполняются на GPU, а обычная часть - на CPU.

Может объяснишь, что ещё за ядра? И где ставить ударение, иначе слово можно прочесть двояко?

0

21

Здесь используется потоковая вычислительная модель - в ней есть потоки однородных данных, которые маожно обрабатывать независимо друг от друга.
Например отдельные вершины или фрагменты в графике.
В потоковой модели обработчик потока называется ядром.
Т.е. ядро - это нечто, на вход посутпаю входные потоки, на выход выходные. В принципе может быть много идущих друг за другом ядер.
Соответственно CUDA -код, выполняемый на GPU это фактически такое ядро

0

22

А можно подробнее изложить как же использовать C++ и DUDA? Вот у меня есть дармовой Visual C++ 2008 и CUDA Toolkit; всё установлено. Что мне запускать? С C++ не работал, пожалуйста попроще.

Отредактировано DungeonLords (2009-07-13 10:56:54)

0

23

Поскольку CUDA основан на С, то для начала хорошо выучить С :))
Запускать можно примеры из CUDA SDK - там готовые проекты для Visual C++

0

24

DungeonLords написал(а):

А можно подробнее изложить как же использовать C++ и DUDA? Вот у меня есть дармовой Visual C++ 2008 и CUDA Toolkit; всё установлено. Что мне запускать? С C++ не работал, пожалуйста попроще.

С дармовым Visual C++ 2008 запустить CUDA не удасться. Нужен atl сервер.
Выход. Переходить на линюкс и работать с gcc.

0

25

Зачем нужен "atl сервер" ?
И какое отношение atl имеет к CUDA ?
Вообще то нужен просто компилятор с С/С++
Попробуйте поставить CUDA SDKи собрать примеры к нему

0

26

Steps3D написал(а):

Зачем нужен "atl сервер" ?
И какое отношение atl имеет к CUDA ?
Вообще то нужен просто компилятор с С/С++
Попробуйте поставить CUDA SDKи собрать примеры к нему

Присоединяюсь к заявлению. Только что успешно перестроил демку OceanFFT в Visual C++ 2008 (кстати, она создавалась не в 2008 C++, а в старше, но 2008 версия сама перестроила решение(за что ей спасибо)).

Отредактировано DungeonLords (2009-08-04 21:55:29)

0

27

Steps3D написал(а):

Зачем нужен "atl сервер" ?
И какое отношение atl имеет к CUDA ?
Вообще то нужен просто компилятор с С/С++
Попробуйте поставить CUDA SDKи собрать примеры к нему

Давно уже CUDA SDK стоит, естественно примеры собирал, причем без проблеми на vc7.1, сейчас гружу последнию версию SDK.
Пытался пару раз (давно уже) собирать под vs2008 express, тогда не получилось, причем со странным сообщением, которое при разборке манулов указывало, что есть какие-то проблемы с atl.
Я тогда не стал глубже вникать, поскольку vc7.1 меня устраивал.
Но сейчас понимаю, что нужно двигаться дальше, поэтому сегодня ревизовал все настройки vs2008, как оказалось вполне удачно: проект из OpenCL SDK (он тоже работает через CUDA) собрался на ура.

0

28

DungeonLords написал(а):

Присоединяюсь к заявлению. Только что успешно перестроил демку OceanFFT в Visual C++ 2008 (кстати, она создавалась не в 2008 C++, а в старше, но 2008 версия сама перестроила решение(за что ей спасибо)).

У меня пока сообщает

1>nvcc fatal   : nvcc cannot find a supported cl version. Only MSVC 7.1 and MSVC 8.0 are supported

Сейчас поставлю новую версию, тогда посмотрим.

Но вообще, признаться, я очень рад, что ошибался. Собирать cuda под vs2008 express - это то, что мне давно было нужно.

0

29

Точно, пришлось поставить новую версию тулкит с новым компилятором nvcc, сейчас  cuda под vs2008 express компилируется.
Неприятный допэффект: теперь не компилируются старые проекты под vc7.1 :(
Пришлось дополнительно сохранить старую версию nvcc.

0



создать форум