Hạt nhân CUDA được khởi chạy như thế nào?

Question

Tôi đã tạo một ứng dụng CUDA đơn giản để thêm hai ma trận. Nó được biên dịch tốt. Tôi muốn biết làm thế nào hạt nhân sẽ được đưa ra bởi tất cả các chủ đề và những gì sẽ lưu lượng bên trong CUDA? Ý tôi là, trong mọi thời trang, mỗi luồng sẽ thực hiện từng phần tử của ma trận.

Tôi biết đây là một khái niệm rất cơ bản, nhưng tôi không biết điều này. Tôi bối rối về dòng chảy.

Answer 1

Bạn khởi chạy một mạng lưới các khối.

Các khối được phân chia một cách không thể phân chia cho các bộ xử lý đa (nơi số lượng khối trên bộ xử lý đa xác định số lượng bộ nhớ dùng chung có sẵn).

Các khối được chia thành các phần dọc. Đối với GPU Fermi có 32 luồng hoặc thực thi cùng một lệnh hoặc không hoạt động (vì chúng phân nhánh, ví dụ như thoát khỏi vòng lặp sớm hơn so với hàng xóm trong cùng một sợi dọc hoặc không chụp if chúng đã làm). Trên một GPU Fermi nhiều nhất hai warp chạy trên một bộ xử lý tại một thời điểm.

Bất cứ khi nào có độ trễ (đó là các quầy thực thi truy cập bộ nhớ hoặc phụ thuộc dữ liệu để hoàn thành) một warp khác đang chạy (số warps phù hợp với một bộ xử lý - cùng khối hoặc khác nhau) được xác định bằng số các thanh ghi được sử dụng bởi mỗi luồng và lượng bộ nhớ dùng chung được sử dụng bởi một/(các) khối).

Quá trình lập lịch này diễn ra một cách minh bạch. Đó là, bạn không cần phải suy nghĩ về nó quá nhiều. Tuy nhiên, bạn có thể muốn sử dụng các vectơ số nguyên được xác định trước threadIdx (chủ đề của tôi nằm trong khối nào?), blockDim (khối lớn bao nhiêu?), blockIdx (khối của tôi ở đâu trong lưới?) Và gridDim là lưới?) để phân chia công việc (đọc: đầu vào và đầu ra) trong số các chủ đề. Bạn cũng có thể muốn đọc cách truy cập hiệu quả các loại bộ nhớ khác nhau (vì vậy nhiều luồng có thể được phục vụ trong một giao dịch duy nhất) - nhưng điều đó dẫn đến việc tắt chủ đề.

NSight cung cấp trình gỡ lỗi đồ họa cung cấp cho bạn ý tưởng hay về những gì đang xảy ra trên thiết bị khi bạn đã vượt qua khu rừng bí ẩn. Cũng vậy với trình thông báo của nó về những thứ bạn sẽ không thấy trong trình gỡ rối (ví dụ: các lý do gian hàng hoặc áp lực bộ nhớ).

Bạn có thể đồng bộ hóa tất cả các chuỗi trong lưới (tất cả đều có) bằng cách khởi chạy hạt nhân khác. Đối với thực thi hạt nhân tuần tự không chồng chéo, không cần đồng bộ hóa thêm nữa.

Chủ đề trong một lưới (hoặc một hạt nhân chạy - tuy nhiên bạn muốn gọi nó) có thể giao tiếp qua bộ nhớ toàn cầu bằng cách sử dụng các phép toán nguyên tử (cho số học) hoặc hàng rào bộ nhớ thích hợp (để tải hoặc truy cập cửa hàng).

Bạn có thể đồng bộ hóa tất cả các chuỗi trong một khối với hướng dẫn nội tại __syncthreads() (tất cả các chuỗi sẽ hoạt động sau đó - mặc dù, như thường lệ, tối đa hai warps có thể chạy trên GPU Fermi). Các luồng trong một khối có thể giao tiếp thông qua bộ nhớ chia sẻ hoặc toàn cục sử dụng các phép toán nguyên tử (cho số học) hoặc các hàng rào bộ nhớ thích hợp (để tải hoặc truy cập lưu trữ).

Như đã đề cập trước đó, tất cả các chuỗi trong một sợi dọc luôn được "đồng bộ", mặc dù một số luồng có thể không hoạt động. Họ có thể giao tiếp thông qua bộ nhớ chung hoặc toàn cầu (hoặc "chuyển đổi làn đường" trên phần cứng sắp tới với khả năng tính toán 3). Bạn có thể sử dụng các phép toán nguyên tử (đối với số học) và các biến được chia sẻ hoặc biến toàn cầu có đủ điều kiện bay hơi (tải hoặc truy cập lưu trữ xảy ra tuần tự trong cùng một sợi dọc).Bộ định mức dễ bay hơi cho trình biên dịch luôn luôn truy cập vào bộ nhớ và không bao giờ đăng ký có trạng thái không thể được nhìn thấy bởi các luồng khác.

Hơn nữa, có chức năng bỏ phiếu toàn cảnh có thể giúp bạn đưa ra quyết định chi nhánh hoặc tính tổng số tiền tố (tiền tố).

OK, về cơ bản là vậy. Hy vọng rằng sẽ giúp. Đã có một dòng chảy tốt bằng văn bản :-).

Answer 2

Hãy thử 'Cuda-gdb', trình gỡ lỗi CUDA.

Answer 3

Hãy lấy một ví dụ của việc bổ sung các 4 * 4 ma trận .. bạn có hai ma trận A và B, có chiều 4 * 4 ..

int main() 
{ 
int *a, *b, *c;   //To store your matrix A & B in RAM. Result will be stored in matrix C 
int *ad, *bd, *cd;   // To store matrices into GPU's RAM. 
int N =4;     //No of rows and columns. 

size_t size=sizeof(float)* N * N; 

a=(float*)malloc(size);  //Allocate space of RAM for matrix A 
b=(float*)malloc(size);  //Allocate space of RAM for matrix B 

//allocate memory on device 
    cudaMalloc(&ad,size); 
    cudaMalloc(&bd,size); 
    cudaMalloc(&cd,size); 

//initialize host memory with its own indices 
    for(i=0;i<N;i++) 
     { 
    for(j=0;j<N;j++) 
     { 
      a[i * N + j]=(float)(i * N + j); 
      b[i * N + j]= -(float)(i * N + j); 
     } 
     } 

//copy data from host memory to device memory 
    cudaMemcpy(ad, a, size, cudaMemcpyHostToDevice); 
    cudaMemcpy(bd, b, size, cudaMemcpyHostToDevice); 

//calculate execution configuration 
    dim3 grid (1, 1, 1); 
    dim3 block (16, 1, 1); 

//each block contains N * N threads, each thread calculates 1 data element 

    add_matrices<<<grid, block>>>(ad, bd, cd, N); 

    cudaMemcpy(c,cd,size,cudaMemcpyDeviceToHost); 
    printf("Matrix A was---\n"); 
    for(i=0;i<N;i++) 
    { 
     for(j=0;j<N;j++) 
      printf("%f ",a[i*N+j]); 
     printf("\n"); 
    } 

    printf("\nMatrix B was---\n"); 
    for(i=0;i<N;i++) 
    { 
     for(j=0;j<N;j++) 
      printf("%f ",b[i*N+j]); 
     printf("\n"); 
    } 

    printf("\nAddition of A and B gives C----\n"); 
    for(i=0;i<N;i++) 
    { 
     for(j=0;j<N;j++) 
      printf("%f ",c[i*N+j]); //if correctly evaluated, all values will be 0 
     printf("\n"); 
    } 



    //deallocate host and device memories 
    cudaFree(ad); 
    cudaFree(bd); 
    cudaFree (cd); 

    free(a); 
    free(b); 
    free(c); 

    getch(); 
    return 1; 
} 

/////Kernel Part 

__global__ void add_matrices(float *ad,float *bd,float *cd,int N) 
{ 
    int index; 
    index = blockIDx.x * blockDim.x + threadIDx.x    

    cd[index] = ad[index] + bd[index]; 
} 

Hãy lấy một ví dụ bổ sung 16 * 16 ma trận .. bạn có hai ma trận A và B, có kích thước 16 * 16 ..

Trước hết, bạn phải quyết định cấu hình chuỗi của mình. Giả sử bạn khởi động một hàm hạt nhân, nó sẽ thực hiện phép tính song song của phép thêm ma trận, sẽ được thực hiện trên thiết bị GPU của bạn.

Bây giờ, một lưới được khởi chạy với một hàm hạt nhân. Lưới có thể có tối đa 65,535 không có khối nào có thể được sắp xếp theo các chiều 3 chiều. (65535 * 65535 * 65535).

Mỗi khối trong lưới có thể có tối đa 1024 không có các chủ đề threads.Those cũng có thể được sắp xếp theo 3 cách chiều (1024 * 1024 * 64)

Bây giờ vấn đề của chúng ta là bổ sung 16 * 16 ma trận ..

A | 1 2 3 4 |  B | 1 2 3 4 |  C| 1 2 3 4 | 
    | 5 6 7 8 | +  | 5 6 7 8 | = | 5 6 7 8 | 
    | 9 10 11 12 |   | 9 10 11 12 |  | 9 10 11 12 | 
    | 13 14 15 16|   | 13 14 15 16|  | 13 14 15 16| 

Chúng tôi cần 16 chuỗi để thực hiện tính toán.

i.e. A(1,1) + B (1,1) = C(1,1) 
    A(1,2) + B (1,2) = C(1,2) 
    .  .   . 
    .  .   . 
    A(4,4) + B (4,4) = C(4,4) 

Tất cả các chủ đề này sẽ được thực hiện đồng thời. Vì vậy, chúng tôi cần một khối với 16 chủ đề. Để thuận tiện cho chúng tôi, chúng tôi sẽ sắp xếp chủ đề theo cách (16 * 1 * 1) theo khối Vì không có chủ đề nào là 16 nên chúng tôi chỉ cần một khối để lưu trữ 16 luồng đó.

vì vậy, cấu hình lưới sẽ là dim3 Grid(1,1,1) tức là lưới sẽ chỉ có một khối và cấu hình chặn sẽ là dim3 block(16,1,1) tức là khối sẽ có 16 chuỗi được sắp xếp một cách khôn ngoan.

Chương trình sau sẽ cung cấp cho bạn ý tưởng rõ ràng về việc thực thi của nó .. Hiểu phần lập chỉ mục (ví dụ: threadID, blockDim, blockID) là phần quan trọng. Bạn cần phải trải qua văn học CUDA. Một khi bạn có ý tưởng rõ ràng về lập chỉ mục, bạn sẽ giành chiến thắng trong một nửa trận chiến! Vì vậy, dành một chút thời gian với sách cuda, các thuật toán khác nhau và bút chì giấy tất nhiên!