CUDA By Example (1)

9148 단어 example
cudaMalloc는 장치에 메모리를 분배합니다. 첫 번째 파라미터는 새로 분배된 메모리의 주소이고, 두 번째 파라미터는 분배 메모리 크기입니다.호스트에서 이 메모리에 대해 어떠한 수정도 할 수 없습니다.호스트 포인터는 호스트 코드의 메모리만 액세스할 수 있고 장치 포인터는 장치 코드의 메모리만 액세스할 수 있습니다.장치 메모리를 사용한 후 cudaFree로 분배된 메모리를 방출해야 합니다.호스트 코드에서 cudaMemcpy를 호출하여 장치의 메모리에 접근할 수 있습니다.
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include<conio.h>
#include <stdio.h>

__global__ void add(int a,int b,int *c){    //        
	*c=a+b;
}

int main(){
	int c;
	int *dev_c;
	cudaMalloc((void**)&dev_c,sizeof(int));    //       dev_c         
	add<<<1,1>>>(2,7,dev_c);    //  grid     block，  block     thread     kernel
	cudaMemcpy(&c,dev_c,sizeof(int),cudaMemcpyDeviceToHost);  //   dev_c          c
	printf("2+7=%d",c);
	cudaFree(dev_c);    //    dev_c     
	getche();
	return 0;
}

 
뾰족괄호에는 두 개의 파라미터가 있는데, 첫 번째 파라미터는 장치가 핵 함수를 실행할 때 사용하는 병렬 블록 블록의 수량을 가리킨다.kernel <<<256,1>>>은 이 kernel을 실행하기 위해 Grid에 256개의 스레드 블록이 있다는 것을 의미합니다.BlockIdx를 사용할 수 있습니다.x로 라인 블록의 인덱스를 가져옵니다.CUDA는 2차원 스레드 블록 배열을 지원하여 행렬을 처리할 때나 도형을 처리할 때 매우 편리하다.스레드 블록 그룹을 시작할 때, 그룹의 매 비트는 65335를 초과할 수 없습니다.두 번째 파라미터는 모든 스레드 블록에서 실행되는 스레드의 수량을 표시합니다.스레드 블록의 병렬 스레드는 병렬 스레드 블록(block)이 할 수 없는 작업을 완성할 수 있습니다.스레드 블록의 스레드는threadIdx를 사용할 수 있습니다.색인
  
#ifndef _use_h
#define _use_h
    #include "cuda_runtime.h"
    #include "device_launch_parameters.h"
    #include <conio.h>
    #include <stdio.h>
    #define N 100
    __global__ void add(int *a,int *b,int *c);
#endif
#include "use.h"
__global__ void add(int *a,int *b,int *c){
	//int tid=blockIdx.x;    //tid  i block
	int tid=threadIdx.x;     //tid  i thread
	if(tid<N)
		c[tid]=a[tid]+b[tid];
}

  
#include "use.h"

int main(){
	int a[N],b[N],c[N];
	int *dev_a,*dev_b,*dev_c;

	cudaMalloc((void**)&dev_a,N*sizeof(int));
	cudaMalloc((void**)&dev_b,N*sizeof(int));
	cudaMalloc((void**)&dev_c,N*sizeof(int));

	//CPU a,b  
	for(int i=0;i<N;i++){
		a[i]=-i;
		b[i]=i*i;
	}

	// a，b   GPU 
	cudaMemcpy(dev_a,a,N*sizeof(int),cudaMemcpyHostToDevice);
	cudaMemcpy(dev_b,b,N*sizeof(int),cudaMemcpyHostToDevice);

	//add<<<N,1>>>(dev_a,dev_b,dev_c);  //N block，  block          add
	add<<<1,N>>>(dev_a,dev_b,dev_c);  //1 block，  block  N        add

	// c GPU   CPU 
	cudaMemcpy(c,dev_c,N*sizeof(int),cudaMemcpyDeviceToHost);

	for(int i=0;i<N;i++){
		printf("%d+%d=%d
",a[i],b[i],c[i]);
	}

	//  GPU      
	cudaFree(dev_a);
	cudaFree(dev_b);
	cudaFree(dev_c);
	getche();
}

  
gridDim을 사용할 수 있습니다.x는 grid당 몇 개의 Block을 가져옵니다. BlockDim을 사용하십시오.Block마다 thread가 얼마나 있는지 가져오십시오.이 두 수는 라인 포지셔닝에 매우 유용하다.상기 벡터를 더하면 스레드 수량과 하드웨어의 일부 제한을 고려하지 않았다. 예를 들어 스레드 칸은 한 명당 65535를 초과하지 않는다.다음과 같이 향상되었습니다.
#include "use.h"  
__global__ void add(int *a,int *b,int *c){   
    int tid=blockIdx.x*blockDim.x+threadIdx.x;     // i block   i*blockDim   ，tid        
    while(tid<N){
        c[tid]=a[tid]+b[tid];
	    tid+=blockDim.x*gridDim.x;    //  N         ，      i      （i+k*   ）   
	}
}

  
#include "use.h"

int main(){
	int a[N],b[N],c[N];
	int *dev_a,*dev_b,*dev_c;

	//GPU     
	cudaMalloc((void**)&dev_a,N*sizeof(int));
	cudaMalloc((void**)&dev_b,N*sizeof(int));
	cudaMalloc((void**)&dev_c,N*sizeof(int));

	//CPU a,b  
	for(int i=0;i<N;i++){
		a[i]=i;
		b[i]=i*i;
	}

	// a，b   GPU 
	cudaMemcpy(dev_a,a,N*sizeof(int),cudaMemcpyHostToDevice);
	cudaMemcpy(dev_b,b,N*sizeof(int),cudaMemcpyHostToDevice);

	add<<<128,128>>>(dev_a,dev_b,dev_c);  //128 block，  block  128        add， 16384，  N 1/3

	// c GPU   CPU 
	cudaMemcpy(c,dev_c,N*sizeof(int),cudaMemcpyDeviceToHost);
	int success=1;
	for(int i=0;i<N;i++){
		//printf("%d+%d=%d
",a[i],b[i],c[i]);
		if(a[i]+b[i]!=c[i])
			{success=0;break;}
	}
	if(success==1)
		printf("sucess
");
	else
		printf("failure
");

	//  GPU      
	cudaFree(dev_a);
	cudaFree(dev_b);
	cudaFree(dev_c);
	getche();
}

 
키워드share__변수 성명에 사용되며 공유 메모리에 변수를 주재할 수 있습니다.CUDA C는 공유 변수를 일반 변수와 다르게 처리합니다.공유 메모리의 변수에 대해 컴파일러는 그 부본을 만들 것입니다. 이 Block의 모든thread는 이 부본을 공유합니다. (다른 Block의 루트는 접근할 수 없습니다.) 공유 메모리 버퍼가 물리적 GPU에 주재하도록 합니다.공유 메모리에 접근하는 지연은 일반 버퍼에 접근하는 지연보다 훨씬 낮다.그러나 라인 간의 통신은 동기화 메커니즘을 실현해야 한다.
동기화 메커니즘용syncthreads () 로 실행합니다.코드에서 만약 이 함수를 사용한다면, 모든 라인이 이 함수를 실행한 후에야 나머지 부분을 더 실행할 수 있다.
다음 예는 두 벡터의 내적을 계산하는 것입니다.
#ifndef _use_h
#define _use_h
    #include "cuda_runtime.h"
    #include "device_launch_parameters.h"
    #include "device_functions.h"
    #include <conio.h>
    #include <stdio.h>
    #include <stdlib.h>
    #define imin(a,b) (a<b?a:b)
    const int N=33*1024;
    const int threadPerBlock=256;
    const int blockPerGrid=imin(32,(N+threadPerBlock-1) / threadPerBlock);
    __global__ void dot(float *a,float *b,float *c);
#endif

  
#include "use.h"

__global__ void dot(float *a,float *b,float *c){
	__shared__ float cache[threadPerBlock];    //    ，  a[i]*b[i]   ，  block       
	int tid=threadIdx.x+blockIdx.x*blockDim.x;    //  grid      
	int cacheIndex=threadIdx.x;    //block      

	float temp=0;
	while(tid<N){
		temp+=a[tid]*b[tid];    //             
		tid+=blockDim.x*gridDim.x;    //tid       
	}

	cache[cacheIndex]=temp;    //  block   chace  

	__syncthreads();    //            ，             

	int i=blockDim.x/2;
	while(i!=0){
		if(cacheIndex<i)
			cache[cacheIndex]+=cache[cacheIndex+i];
		__syncthreads();  //cache[0] cache[i]     
		i=i/2;
	}
	if(cacheIndex==0)
		c[blockIdx.x]=cache[0];
}

  
#include "use.h"

int main(){
	float *a,*b,c,*partial_c;  //a,b        ，c  ，partial_c  grid  block ，
	                          //partial_c[i]   i block           
	float *dev_a,*dev_b,*dev_partial_c;

	// CPU     
	a=(float*)malloc(N*sizeof(float));
	b=(float*)malloc(N*sizeof(float));
	partial_c=(float*)malloc(blockPerGrid*sizeof(float));

	// GPU     
	cudaMalloc((void**)&dev_a,N*sizeof(float));
	cudaMalloc((void**)&dev_b,N*sizeof(float));
	cudaMalloc((void**)&dev_partial_c,blockPerGrid*sizeof(float));

	//      
	for(int i=0;i<N;i++){
		a[i]=i;
		b[i]=i*2;
	}

	// a，b   GPU
	cudaMemcpy(dev_a,a,N*sizeof(float),cudaMemcpyHostToDevice);
	cudaMemcpy(dev_b,b,N*sizeof(float),cudaMemcpyHostToDevice);

	dot<<<blockPerGrid,threadPerBlock>>>(dev_a,dev_b,dev_partial_c);

	//  partial_c HOST ， HOST   
	cudaMemcpy(partial_c,dev_partial_c,blockPerGrid*sizeof(float),cudaMemcpyDeviceToHost);

	
	c=0;
	for(int i=0;i<blockPerGrid;i++){
		c+=partial_c[i];
	}

		
    #define sum_squares(x) (x*(x+1)*(2*x+1)/6)
    printf("Does GPU value %.6g=%.6g?
",c,2*sum_squares((float)(N-1)));

	//  GPU   
	cudaFree(dev_a);
	cudaFree(dev_b);
	cudaFree(dev_partial_c);

	free(a);
	free(b);
	free(partial_c);
	getche();
}

이 구내적 과정 중 추상적으로는 하나의 입력 수조에 대해 어떤 계산을 한 다음에 더욱 작은 결과 수조를 생성하고 이 과정도 귀약(reduction)이 된다.귀약을 실현하는 가장 간단한 방법은 공유 메모리에서 하나의 라인을 교체하여 총화 값을 계산하는 것이다.
코드 중;
int i=blockDim.x/2;
while(i!=0){
    if(cacheIndex<i)
	cache[cacheIndex]+=cache[cacheIndex+i];
    __syncthreads();
    i=i/2;
}
여기서는 일부 스레드만 작동하고 나머지 스레드는 항상syncthreads 기다림은 다음과 같은 최적화가 있을 수 있습니다. 즉, 일을 하지 않는 라인을 기다림이 아니라 끝까지 계속 실행하도록 하는 것입니다. 그러면 효율이 좀 높을 수도 있습니다.
int i=blockDim.x/2;
while(i!=0){
    if(cacheIndex<i){
	cache[cacheIndex]+=cache[cacheIndex+i];
        __syncthreads();
    }
    i=i/2;
}

하지만, 이것은 일할 수 없습니다. 왜냐하면syncthreads는 모든 스레드가 여기까지 실행되어야만 스레드가 다음 작업을 진행할 수 있지만, 일부 스레드는 영원히 실행되지 않습니다syncthreads (), 이렇게 하면 프로그램이 끝없이 기다릴 수 있습니다.
 
전: Osgi 실전 중 질문
Heroku와 Cloudinary를 사용하여 Strapiv4 무료 이동

좋은 웹페이지 즐겨찾기

개발자 우수 사이트 수집

개발자가 알아야 할 필수 사이트 100선 추천 우리는 당신을 위해 100개의 자주 사용하는 개발자 학습 사이트를 정리했습니다
best100-homepage