Giảm số lượng thanh ghi được sử dụng trong hạt nhân CUDA

Question

Tôi có một hạt nhân sử dụng 17 thanh ghi, giảm xuống còn 16 sẽ mang lại cho tôi 100% dung lượng. Câu hỏi của tôi là: có những phương pháp có thể được sử dụng để giảm số lượng hoặc sổ đăng ký được sử dụng, không bao gồm viết lại hoàn toàn thuật toán của tôi theo cách khác. Tôi đã luôn luôn loại giả định trình biên dịch là thông minh hơn rất nhiều so với tôi, vì vậy ví dụ tôi thường sử dụng các biến thêm cho lợi ích của rõ ràng một mình. Tôi có sai trong suy nghĩ này không?

Xin lưu ý: Tôi biết về --max_registers (hoặc bất kỳ cú pháp là) cờ, nhưng việc sử dụng bộ nhớ địa phương sẽ có nhiều bất lợi hơn so với một công suất thấp hơn 25% (tôi nên kiểm tra này)

Answer 1

Nó thực sự khó để nói, trình biên dịch nvcc không phải là rất thông minh theo ý kiến ​​của tôi.
Bạn có thể thử những điều hiển nhiên, ví dụ: sử dụng ngắn thay vì int, chuyển và sử dụng biến theo tham chiếu (ví dụ: & biến), bỏ vòng lặp, sử dụng mẫu (như trong C++). Nếu bạn có các bộ phận, các hàm siêu việt, được áp dụng theo thứ tự, hãy cố gắng làm cho chúng thành một vòng lặp. Cố gắng loại bỏ các điều kiện, có thể thay thế chúng bằng các tính toán dư thừa.

Nếu bạn đăng một số mã, có thể bạn sẽ nhận được câu trả lời cụ thể.

Answer 2

Công việc có thể gây nhầm lẫn một chút và 100% số người dùng không phải là mục tiêu chính của bạn. Nếu bạn có thể truy cập hoàn toàn vào bộ nhớ toàn cục thì khi đó dung lượng chiếm 50% của GPU sẽ đủ để che giấu độ trễ cho bộ nhớ toàn cầu (đối với phao nổi, thậm chí còn thấp hơn để tăng gấp đôi). Hãy xem bài trình bày Advanced CUDA C từ GTC năm ngoái để biết thêm thông tin về chủ đề này.

Trong trường hợp của bạn, bạn nên đo hiệu suất cả có và không có maxrregcount được đặt thành 16. Độ trễ cho bộ nhớ cục bộ phải được ẩn do có đủ luồng, giả sử bạn không truy cập ngẫu nhiên vào mảng cục bộ (dẫn đến các truy cập không được kết hợp).

Để trả lời cho bạn câu hỏi cụ thể về việc giảm đăng ký, hãy đăng mã để có câu trả lời chi tiết hơn! Hiểu cách trình biên dịch làm việc nói chung có thể hữu ích, nhưng hãy nhớ rằng nvcc là một trình biên dịch tối ưu hóa với một không gian tham số lớn, vì vậy việc giảm thiểu số lượng đăng ký phải được cân bằng với hiệu năng tổng thể.

Answer 3

Tăng số lượng lệnh khi giảm mức sử dụng đăng ký có giải thích đơn giản. Trình biên dịch có thể sử dụng thanh ghi để lưu trữ kết quả của một số thao tác được sử dụng nhiều lần thông qua mã của bạn để tránh tính lại các giá trị đó, khi buộc phải sử dụng ít thanh ghi hơn, trình biên dịch quyết định tính toán lại các giá trị được lưu trữ trong sổ đăng ký nếu không thì.

Answer 4

Đây không phải là cách tiếp cận tốt để giảm thiểu áp lực đăng ký. Trình biên dịch thực hiện tốt công việc tối ưu hóa hiệu suất hạt nhân dự kiến ​​tổng thể, và nó sẽ tính đến rất nhiều yếu tố, bao gồm đăng ký.

Làm thế nào nó hoạt động khi đăng ký giảm do tốc độ chậm hơn

lẽ Hầu hết các trình biên dịch đã phải đổ đủ dữ liệu đăng ký vào bộ nhớ "địa phương", mà chủ yếu là giống như bộ nhớ toàn cầu, và do đó rất chậm

Vì mục đích tối ưu hóa, tôi khuyên bạn nên sử dụng các từ khóa như const, dễ bay hơi, v.v ... khi cần thiết, để giúp trình biên dịch trong giai đoạn tối ưu hóa.

Dù sao, nó không phải là những vấn đề nhỏ như đăng ký mà thường làm cho hạt nhân CUDA chạy chậm.Tôi khuyên bạn nên tối ưu hóa công việc với bộ nhớ toàn cầu, mẫu truy cập, bộ nhớ đệm trong bộ nhớ kết cấu nếu có thể, giao dịch qua PCIe.

Answer 5

Bằng cách sử dụng bộ nhớ chia sẻ như bộ nhớ cache có thể dẫn ít đăng ký sử dụng và ngăn chặn đăng ký tràn vào bộ nhớ địa phương ...

Nghĩ rằng kernel tính toán một số giá trị và các giá trị tính toán được sử dụng bởi tất cả các chủ đề,

__global__ void kernel(...) { 
    int idx = threadIdx.x + blockDim.x * blockIdx.x; 
    int id0 = blockDim.x * blockIdx.x; 

    int reg = id0 * ...; 
    int reg0 = reg * a/x + y; 


    ... 

    int val = reg + reg0 + 2 * idx; 

    output[idx] = val > 10; 
} 

Vì vậy, thay vì giữ reg và reg0 làm sổ đăng ký và làm cho chúng có khả năng tràn ra bộ nhớ cục bộ (bộ nhớ toàn cục), chúng tôi có thể sử dụng bộ nhớ dùng chung.

__global__ void kernel(...) { 
    __shared__ int cache[10]; 

    int idx = threadIdx.x + blockDim.x * blockIdx.x; 

    if (threadIdx.x == 0) { 
     int id0 = blockDim.x * blockIdx.x; 

     cache[0] = id0 * ...; 
     cache[1] = cache[0] * a/x + y; 
    } 
    __syncthreads(); 


    ... 

    int val = cache[0] + cache[1] + 2 * idx; 

    output[idx] = val > 10; 
} 

Hãy xem paper này để biết thêm thông tin ..