커널 호출 중 장치 전용 계산을위한 CUDA 메모리 (유형)

Question

현재 CUDA를 배우고 있으며 알고리즘은 일부 입력 데이터를 기반으로 과중한 계산을 수행해야합니다. 이러한 계산은 최대 1024 회 돌기 루프에서 수행됩니다. 커널 당 작은 수의 스레드 (< 100'000)가있는 한 모든 것이 잘 작동하지만, 더 많은 스레드를 활용하려면 완료하는 데 시간이 오래 걸리므로 커널이 Windows에 의해 중단됩니다.

내 솔루션은 여러 가지 커널 호출에 무거운 계산을 분할하는 것입니다 :

1. 입력 데이터를 준비하고 최초의 X 라운드를 계산합니다 주요 커널 (루프 언 롤링). 이것은 입력마다 한 번만 호출됩니다.
2. 작업 커널, 다음 x 라운드 수행 (루프 언 롤링). 이것은 필요한 모든 라운드를 계산하는 데 필요한만큼 자주 호출됩니다.
각각의 커널 호출 사이

(하나의 메인많은 일), I는 길이가 입력의 길이 (다음 호출에 사용되는 데이터의 16 + 길이 바이트를 저장해야 주 호출 당 고정됩니다. 주 커널은 처음에 해당 바이트를 쓰고 작업 커널이 커널을 읽고 다음 계산을 실행하고 원본 데이터를 새 결과로 씁니다. 장치의 데이터 만 필요하므로 호스트 액세스가 필요하지 않습니다. 어떤 종류의 기억을 사용해야합니까? 적어도 전역 메모리 여야합니다. 커널 메모리가 커널 호출 중에 지속될 수있는 쓰기 가능한 메모리 일뿐입니다. 그렇지 않습니까? 그런데, 뭐라구? 올바른 메모리 (및 최상의 성능)로 작업을 진행하는 방법에 대한 조언을 해 줄 수 있습니까?

에서 '의사'가 다음과 같을 수는 :

prepare memory to hold threads * (16 + length) bytes 

for length = 1 to x step 1 
    call mainKernel 
    rounds = 1024 - rounds_done_in_main 
    for rounds to 0 step rounds_done_in_work 
    call workKernel 
    end for 
end for 

cleanup memory 

-------- 

template <unsigned char length> __global__ mainKernel() { 
    unsigned char input[length]; 
    unsigned char output[16]; 
    const int tid = ...; 

    devPrepareInput<length>(input); 

    calc round 1: doSomething<length>(output, input) 
    calc round 2: doSomething<length>(output, output + input) // '+' == append 

    write data to memory based on tid // data == output + input 
} 

template <unsigned char length, remaining rounds> __global__ workKernel() { 
    unsigned char *input; 
    unsigned char *output; 
    const int tid = ...; 

    read data from memory based on tid 
    ouput = data 
    input = data+16 

    if rounds >= 1 
    calc round x : doSomething<length>(output, output + input) 
    if rounds >= 2 
    calc round x+1: doSomething<length>(output, output + input) // '+' == append 

    if rounds == x // x is the number of rounds in the last work call 
    do final steps on output 
    else 
    write ouput + input to memory based on tid (for next call) 
} 

Answer 1

예, 장치 메모리와 함께이 작업을 수행 할 수 있습니다. __device__으로 선언 된 변수는 cudaMemcpy 연산을 필요로하지 않고 커널이 직접 사용할 수있는 버퍼의 정적 선언을 제공하며 포인터가 커널에 명시 적으로 전달되도록 요구하지 않습니다. lifetime of the application이 있기 때문에, 그 데이터는 한 커널 호출에서 다음 커널 호출까지 지속됩니다.

#define NUM_THREADS 1024 
#define DATA_PER_THREAD 16 
__device__ int temp_data[NUM_THREADS*DATA_PER_THREAD]; 

__global__ my_kernel1(...){ 
    int my_data[DATA_PER_THREAD] = {0}; 
    int idx = threadIdx.x + blockDim.x * blockIdx.x; 
    // perform calculations 

    // write out temp data 
    for (int i = 0; i < DATA_PER_THREAD; i++) temp_data[i + (idx * DATA_PER_THREAD)] = my_data[i]; 
    } 

__global__ my_kernel2(...){ 
    int my_data[DATA_PER_THREAD]; 
    // read in temp data 
    for (int i = 0; i < DATA_PER_THREAD; i++) my_data[i] = temp_data[i + (idx * DATA_PER_THREAD)]; 
    // perform calculations 

    } 

커널의 사용 패턴에 따라이를 최적화 할 수있는 다양한 방법이 있습니다. my_data으로의 데이터 전송은 실제로 필요하지 않습니다. 당연히 커널 코드는 대신에 temp_data에 직접 액세스하여 적절한 색인을 생성 할 수 있습니다.

로드/저장하기로 결정한 경우 데이터를 끼 우고 데이터를 읽고 쓰는 동안 루프 된 액세스를 허용 할 수 있습니다.