Thực hiện một phần quan trọng trong CUDA

Question

Tôi đang cố gắng triển khai một phần quan trọng trong CUDA bằng cách sử dụng hướng dẫn nguyên tử, nhưng tôi gặp phải một số sự cố. Tôi đã tạo ra các chương trình thử nghiệm để hiển thị các vấn đề:

#include <cuda_runtime.h> 
#include <cutil_inline.h> 
#include <stdio.h> 

__global__ void k_testLocking(unsigned int* locks, int n) { 
    int id = threadIdx.x % n; 
    while (atomicExch(&(locks[id]), 1u) != 0u) {} //lock 
    //critical section would go here 
    atomicExch(&(locks[id]),0u); //unlock 
} 

int main(int argc, char** argv) { 
    //initialize the locks array on the GPU to (0...0) 
    unsigned int* locks; 
    unsigned int zeros[10]; for (int i = 0; i < 10; i++) {zeros[i] = 0u;} 
    cutilSafeCall(cudaMalloc((void**)&locks, sizeof(unsigned int)*10)); 
    cutilSafeCall(cudaMemcpy(locks, zeros, sizeof(unsigned int)*10, cudaMemcpyHostToDevice)); 

    //Run the kernel: 
    k_testLocking<<<dim3(1), dim3(256)>>>(locks, 10); 

    //Check the error messages: 
    cudaError_t error = cudaGetLastError(); 
    cutilSafeCall(cudaFree(locks)); 
    if (cudaSuccess != error) { 
     printf("error 1: CUDA ERROR (%d) {%s}\n", error, cudaGetErrorString(error)); 
     exit(-1); 
    } 
    return 0; 
} 

Mã này, thật không may, cứng đóng băng máy tính của tôi trong vài giây và cuối cùng thoát ra, in ra thông điệp:

fcudaSafeCall() Runtime API error in file <XXX.cu>, line XXX : the launch timed out and was terminated. 

có nghĩa là một trong số những vòng lặp đó không quay trở lại, nhưng có vẻ như nó sẽ hoạt động.

Làm lời nhắc atomicExch(unsigned int* address, unsigned int val) đặt nguyên tử giá trị của vị trí bộ nhớ được lưu trữ trong địa chỉ thành val và trả về giá trị old. Vì vậy, ý tưởng đằng sau cơ chế khóa của tôi là ban đầu nó là 0u, do đó, một chủ đề sẽ vượt qua vòng lặp while và tất cả các chuỗi khác sẽ đợi trên vòng lặp while vì chúng sẽ đọc locks[id] là 1u. Sau đó, khi chuỗi được thực hiện với phần quan trọng , nó sẽ đặt lại khóa thành 0u để một chuỗi khác có thể nhập.

Tôi đang thiếu gì?

Bằng cách này, tôi đang biên soạn với:

nvcc -arch sm_11 -Ipath/to/cuda/C/common/inc XXX.cu 

Answer 1

Được rồi, tôi figured it out, và điều này vẫn chưa-một-một-of-the-CUDA-mô-đau.

Như bất kỳ lập trình viên cuda giỏi nào cũng biết (thông báo rằng tôi không nhớ điều này khiến tôi trở thành một lập trình viên cuda xấu), tất cả các chủ đề trong một warp phải thực thi cùng một mã. Mã tôi đã viết sẽ làm việc hoàn hảo nếu không cho thực tế này. Vì nó là, tuy nhiên, có khả năng là hai chủ đề trong cùng một warp truy cập vào cùng một khóa. Nếu một trong số họ mua lại khóa, nó chỉ quên về việc thực thi vòng lặp, nhưng nó không thể tiếp tục qua vòng lặp cho đến khi tất cả các luồng khác trong sợi dọc của nó đã hoàn thành vòng lặp. Thật không may, các chủ đề khác sẽ không bao giờ hoàn thành bởi vì nó đang chờ người đầu tiên mở khóa.

Dưới đây là một hạt nhân mà sẽ làm các trick mà không có lỗi:

__global__ void k_testLocking(unsigned int* locks, int n) { 
    int id = threadIdx.x % n; 
    bool leaveLoop = false; 
    while (!leaveLoop) { 
     if (atomicExch(&(locks[id]), 1u) == 0u) { 
      //critical section 
      leaveLoop = true; 
      atomicExch(&(locks[id]),0u); 
     } 
    } 
} 

Answer 2

bằng cách u phải nhớ rằng bộ nhớ toàn cầu viết và! đọc không được hoàn thành nơi u viết chúng trong mã ... do đó, để thực hành này bạn cần phải thêm một memfence toàn cầu tức là __threadfence()

Answer 3

Áp phích đã tìm thấy câu trả lời cho vấn đề của riêng mình. Tuy nhiên, trong đoạn mã bên dưới, tôi đang cung cấp một khung chung để triển khai một phần quan trọng trong CUDA. Cụ thể hơn, mã thực hiện đếm khối, nhưng nó có thể dễ dàng sửa đổi để lưu trữ các hoạt động khác được thực hiện trong một phần quan trọng . Dưới đây, tôi cũng báo cáo một số giải thích về mã, với một số lỗi "điển hình" trong việc thực hiện các phần quan trọng trong CUDA.

BỘ LUẬT

#include <stdio.h> 

#include "Utilities.cuh" 

#define NUMBLOCKS 512 
#define NUMTHREADS 512 * 2 

/***************/ 
/* LOCK STRUCT */ 
/***************/ 
struct Lock { 

    int *d_state; 

    // --- Constructor 
    Lock(void) { 
     int h_state = 0;          // --- Host side lock state initializer 
     gpuErrchk(cudaMalloc((void **)&d_state, sizeof(int))); // --- Allocate device side lock state 
     gpuErrchk(cudaMemcpy(d_state, &h_state, sizeof(int), cudaMemcpyHostToDevice)); // --- Initialize device side lock state 
    } 

    // --- Destructor 
    __host__ __device__ ~Lock(void) { 
#if !defined(__CUDACC__) 
     gpuErrchk(cudaFree(d_state)); 
#else 

#endif 
    } 

    // --- Lock function 
    __device__ void lock(void) { while (atomicCAS(d_state, 0, 1) != 0); } 

    // --- Unlock function 
    __device__ void unlock(void) { atomicExch(d_state, 0); } 
}; 

/*************************************/ 
/* BLOCK COUNTER KERNEL WITHOUT LOCK */ 
/*************************************/ 
__global__ void blockCountingKernelNoLock(int *numBlocks) { 

    if (threadIdx.x == 0) { numBlocks[0] = numBlocks[0] + 1; } 
} 

/**********************************/ 
/* BLOCK COUNTER KERNEL WITH LOCK */ 
/**********************************/ 
__global__ void blockCountingKernelLock(Lock lock, int *numBlocks) { 

    if (threadIdx.x == 0) { 
     lock.lock(); 
     numBlocks[0] = numBlocks[0] + 1; 
     lock.unlock(); 
    } 
} 

/****************************************/ 
/* BLOCK COUNTER KERNEL WITH WRONG LOCK */ 
/****************************************/ 
__global__ void blockCountingKernelDeadlock(Lock lock, int *numBlocks) { 

    lock.lock(); 
    if (threadIdx.x == 0) { numBlocks[0] = numBlocks[0] + 1; } 
    lock.unlock(); 
} 

/********/ 
/* MAIN */ 
/********/ 
int main(){ 

    int h_counting, *d_counting; 
    Lock lock; 

    gpuErrchk(cudaMalloc(&d_counting, sizeof(int))); 

    // --- Unlocked case 
    h_counting = 0; 
    gpuErrchk(cudaMemcpy(d_counting, &h_counting, sizeof(int), cudaMemcpyHostToDevice)); 

    blockCountingKernelNoLock << <NUMBLOCKS, NUMTHREADS >> >(d_counting); 
    gpuErrchk(cudaPeekAtLastError()); 
    gpuErrchk(cudaDeviceSynchronize()); 

    gpuErrchk(cudaMemcpy(&h_counting, d_counting, sizeof(int), cudaMemcpyDeviceToHost)); 
    printf("Counting in the unlocked case: %i\n", h_counting); 

    // --- Locked case 
    h_counting = 0; 
    gpuErrchk(cudaMemcpy(d_counting, &h_counting, sizeof(int), cudaMemcpyHostToDevice)); 

    blockCountingKernelLock << <NUMBLOCKS, NUMTHREADS >> >(lock, d_counting); 
    gpuErrchk(cudaPeekAtLastError()); 
    gpuErrchk(cudaDeviceSynchronize()); 

    gpuErrchk(cudaMemcpy(&h_counting, d_counting, sizeof(int), cudaMemcpyDeviceToHost)); 
    printf("Counting in the locked case: %i\n", h_counting); 

    gpuErrchk(cudaFree(d_counting)); 
} 

MÃ GIẢI THÍCH

Phần quan trọng là trình tự của các hoạt động đó phải được thực hiện tuần tự theo chủ đề CUDA.

Giả sử để tạo hạt nhân có nhiệm vụ tính toán số lượng khối chuỗi của lưới ren. Một ý tưởng có thể là để cho mỗi luồng trong mỗi khối có threadIdx.x == 0 tăng bộ đếm toàn cục. Để ngăn chặn các điều kiện chủng tộc, tất cả các gia tăng phải xảy ra tuần tự, vì vậy chúng phải được kết hợp trong một phần quan trọng.

Mã trên có hai hàm hạt nhân: blockCountingKernelNoLock và blockCountingKernelLock. Trước đây không sử dụng một phần quan trọng để tăng bộ đếm và, như người ta có thể thấy, trả về kết quả sai. Sau này đóng gói sự gia tăng truy cập trong một phần quan trọng và do đó tạo ra kết quả chính xác. Nhưng phần quan trọng hoạt động như thế nào?

Phần quan trọng được điều chỉnh bởi trạng thái toàn cầu d_state. Ban đầu, tiểu bang là 0. Hơn nữa, hai phương thức __device__, lock và unlock, có thể thay đổi trạng thái này. Các phương thức lock và unlock chỉ có thể được gọi bởi một chuỗi đơn trong mỗi khối và đặc biệt là bởi chuỗi có chỉ mục chuỗi địa phương threadIdx.x == 0.

Ngẫu nhiên trong khi thực hiện, một trong các chuỗi có chỉ mục chuỗi địa phương threadIdx.x == 0 và chỉ mục chuỗi tổng thể, giả sử, t sẽ là phương thức đầu tiên gọi phương thức lock. Đặc biệt, nó sẽ khởi chạy atomicCAS(d_state, 0, 1). Kể từ lúc đầu tiên d_state == 0, sau đó d_state sẽ được cập nhật thành 1, atomicCAS sẽ trả lại 0 và chuỗi sẽ thoát khỏi chức năng lock, chuyển đến hướng dẫn cập nhật. Trong khi đó một luồng như vậy thực hiện các hoạt động đã đề cập, tất cả các luồng khác của tất cả các khối khác có threadIdx.x == 0 sẽ thực hiện phương thức lock. Tuy nhiên, họ sẽ tìm thấy giá trị là d_state bằng 1, để atomicCAS(d_state, 0, 1) sẽ không thực hiện cập nhật và sẽ trả về 1, vì vậy, để lại các chuỗi này chạy vòng lặp while. Sau đó thread t hoàn thành bản cập nhật, sau đó nó thực thi chức năng unlock, cụ thể là atomicExch(d_state, 0), do đó khôi phục d_state thành 0. Tại thời điểm này, một cách ngẫu nhiên, một chủ đề khác với threadIdx.x == 0 sẽ khóa lại trạng thái.

Mã trên cũng chứa hàm hạt nhân thứ ba, cụ thể là blockCountingKernelDeadlock. Tuy nhiên, đây là một thực hiện sai khác của phần quan trọng, dẫn đến deadlocks. Thật vậy, chúng tôi nhớ lại rằng warps hoạt động ở chế độ khóa và chúng đồng bộ hóa sau mỗi lệnh. Vì vậy, khi chúng tôi thực hiện blockCountingKernelDeadlock, có khả năng là một trong các chủ đề trong một sợi dọc, nói một sợi chỉ với chỉ số thread địa phương t≠0, sẽ khóa trạng thái. Trong trường hợp này, các chủ đề khác trong cùng một sợi dọc của t, bao gồm cả với threadIdx.x == 0, sẽ thực hiện cùng một câu lệnh vòng lặp while như luồng t, là việc thực hiện các luồng trong cùng một đường dọc được thực hiện trong khóa. Theo đó, tất cả các chủ đề sẽ chờ đợi một người nào đó để mở khóa nhà nước, nhưng không có chủ đề khác sẽ có thể làm như vậy, và mã sẽ bị mắc kẹt trong một bế tắc.