我試圖在 CUDA C 中實作 alpha 影像混合演算法。我的代碼沒有錯誤。它編譯得很好。根據執行緒邏輯，如果我以增加的執行緒數運行代碼，則應該減少運行時。在我的代碼中，我得到了一種奇怪的運行時模式。當我用 1 個執行緒運行代碼時，運行時為 8.060539 e-01 秒，當我用 4 個執行緒運行代碼時，我得到運行時 7.579031 e-01 秒，當它運行 8 個執行緒時，運行時為 7.810102e-01，并且對于 256 個執行緒，運行時為 7.875319e-01。
這是我的代碼：
#include <stdio.h>
#include <stdlib.h>
#include "timer.h"
#define STB_IMAGE_IMPLEMENTATION
#include "stb_image.h"
#define STB_IMAGE_WRITE_IMPLEMENTATION
#include "stb_image_write.h"
__global__ void image_blend(unsigned char *Pout, unsigned char *pin1, unsigned char *pin2, int width, int height, int channels, float alpha){
int col = threadIdx.x blockIdx.x*blockDim.x;
int row = threadIdx.y blockIdx.y*blockDim.y;
if(col<width && row<height){
size_t img_size = width * height * channels;
if (Pout != NULL)
{
for (size_t i = 0; i < img_size; i )
{
Pout[i] = ((1.0 - alpha) * pin1[i] alpha * pin2[i]);
}
}
}
}
int main(int argc, char* argv[]){
int thread_count;
double start, finish;
float alpha;
int width, height, channels;
unsigned char *new_img;
thread_count = strtol(argv[1], NULL, 10);
printf("Enter the value for alpha:");
scanf("%f", &alpha);
unsigned char *apple = stbi_load("apple.jpg", &width, &height, &channels, 0);
unsigned char *orange = stbi_load("orange.jpg", &width, &height, &channels, 0);
size_t img_size = width * height * channels;
//unsigned char *new_img = malloc(img_size);
cudaMallocManaged(&new_img,img_size*sizeof(unsigned char));
cudaMallocManaged(&apple,img_size* sizeof(unsigned char));
cudaMallocManaged(&orange, img_size*sizeof(unsigned char));
GET_TIME(start);
image_blend<<<1,16,thread_count>>>(new_img,apple, orange, width, height, channels,alpha);
cudaDeviceSynchronize();
GET_TIME(finish);
stbi_write_jpg("new_image.jpg", width, height, channels, new_img, 100);
cudaFree(new_img);
cudaFree(apple);
cudaFree(orange);
printf("\n Elapsed time for cuda = %e seconds\n", finish-start);
}
在運行時得到一個奇怪的模式后，我對代碼的實作有點懷疑。即使我的代碼沒有錯誤，任何人都可以讓我知道為什么我得到那些運行時。
uj5u.com熱心網友回復：
讓我們從這里開始：
image_blend<<<1,16,thread_count>>>(new_img,apple, orange, width, height, channels,alpha);
顯然您不了解內核啟動語法：
<<<1,16,thread_count>>>
第一個數字 (1) 是要啟動的塊數。第二個數字 (16) 是每個塊的執行緒數。第三個數字 ( thread_count) 是動態分配的共享記憶體的大小（以位元組為單位）。
因此，我們的第一個觀察結果是，盡管您聲稱更改了執行緒數，但您沒有。您正在更改動態分配的共享記憶體的位元組數。由于您的內核代碼不使用共享記憶體，因此這是一個完全沒有意義的變數。
讓我們也觀察一下你的內核代碼：
for (size_t i = 0; i < img_size; i )
{
Pout[i] = ((1.0 - alpha) * pin1[i] alpha * pin2[i]);
}
對于通過 if 測驗的每個執行緒，這些執行緒中的每一個都將執行整個 for 回圈并處理整個影像。這不是撰寫 CUDA 內核的一般想法。一般的想法是分解作業，以便每個執行緒完成一部分作業，而不是整個活動。
這些是非常基本的觀察結果。如果您利用對 CUDA 的有序介紹（例如此處），您可以超越其中一些基本概念。
我們還可以指出，您的內核名義上期望 2D 啟動，而您沒有提供一個，也許還有許多其他觀察結果。您缺少的另一個重要概念是您不能這樣做：
unsigned char *apple = stbi_load("apple.jpg", &width, &height, &channels, 0);
...
cudaMallocManaged(&apple,img_size* sizeof(unsigned char));
并期望從中獲得任何明智的結果。如果您想了解資料如何從主機分配移動到設備，請研究幾乎所有 CUDA 示例代碼，例如vectorAdd. 使用托管分配不允許您像正在做的那樣覆寫指標并從中獲得任何有用的東西。
我將提供一個示例，說明如何去做我認為您建議的事情，而不提供有關 CUDA 的完整教程。舉個例子，我將跳過 STB 影像加載例程。要了解您在此處嘗試執行的作業，實際的影像內容并不重要。
這是一個影像處理內核 (1D) 的示例，它將：
- 處理整個影像，只處理一次
- 粗略地說，隨著執行緒數的增加，使用更少的時間。
你還沒有提供你的計時器例程/代碼，所以我會提供我自己的：
$ cat t2130.cu
#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <sys/time.h>
#define USECPSEC 1000000ULL
unsigned long long dtime_usec(unsigned long long start=0){
timeval tv;
gettimeofday(&tv, 0);
return ((tv.tv_sec*USECPSEC) tv.tv_usec)-start;
}
unsigned char *i_load(int w, int h, int c, int init){
unsigned char *res = new unsigned char[w*h*c];
for (int i = 0; i < w*h*c; i ) res[i] = init;
return res;
}
__global__ void image_blend(unsigned char *Pout, unsigned char *pin1, unsigned char *pin2, int width, int height, int channels, float alpha){
if (Pout != NULL)
{
size_t img_size = width * height * channels;
for (size_t i = blockIdx.x*blockDim.x threadIdx.x; i < img_size; i =gridDim.x*blockDim.x) // grid-stride loop
{
Pout[i] = ((1.0 - alpha) * pin1[i] alpha * pin2[i]);
}
}
}
int main(int argc, char* argv[]){
int threads_per_block = 64;
unsigned long long dt;
float alpha;
int width = 1920;
int height = 1080;
int channels = 3;
size_t img_size = width * height * channels;
int thread_count = img_size;
if (argc > 1) thread_count = atoi(argv[1]);
unsigned char *new_img, *m_apple, *m_orange;
printf("Enter the value for alpha:");
scanf("%f", &alpha);
unsigned char *apple = i_load(width, height, channels, 10);
unsigned char *orange = i_load(width, height, channels, 70);
//unsigned char *new_img = malloc(img_size);
cudaMallocManaged(&new_img,img_size*sizeof(unsigned char));
cudaMallocManaged(&m_apple,img_size* sizeof(unsigned char));
cudaMallocManaged(&m_orange, img_size*sizeof(unsigned char));
memcpy(m_apple, apple, img_size);
memcpy(m_orange, orange, img_size);
int blocks;
if (thread_count < threads_per_block) {threads_per_block = thread_count; blocks = 1;}
else {blocks = thread_count/threads_per_block;}
printf("running with %d blocks of %d threads\n", blocks, threads_per_block);
dt = dtime_usec(0);
image_blend<<<blocks, threads_per_block>>>(new_img,m_apple, m_orange, width, height, channels,alpha);
cudaDeviceSynchronize();
dt = dtime_usec(dt);
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) printf("CUDA Error: %s\n", cudaGetErrorString(err));
else printf("\n Elapsed time for cuda = %e seconds\n", dt/(float)USECPSEC);
cudaFree(new_img);
cudaFree(m_apple);
cudaFree(m_orange);
}
$ nvcc -o t2130 t2130.cu
$ ./t2130 1
Enter the value for alpha:0.2
running with 1 blocks of 1 threads
Elapsed time for cuda = 5.737880e-01 seconds
$ ./t2130 2
Enter the value for alpha:0.2
running with 1 blocks of 2 threads
Elapsed time for cuda = 3.230150e-01 seconds
$ ./t2130 32
Enter the value for alpha:0.2
running with 1 blocks of 32 threads
Elapsed time for cuda = 4.865200e-02 seconds
$ ./t2130 64
Enter the value for alpha:0.2
running with 1 blocks of 64 threads
Elapsed time for cuda = 2.623300e-02 seconds
$ ./t2130 128
Enter the value for alpha:0.2
running with 2 blocks of 64 threads
Elapsed time for cuda = 1.546000e-02 seconds
$ ./t2130
Enter the value for alpha:0.2
running with 97200 blocks of 64 threads
Elapsed time for cuda = 5.809000e-03 seconds
$
（CentOS 7、CUDA 11.4、V100）
允許內核完成所有作業（僅一次）同時有效地使用“任意”數量的執行緒的關鍵方法是grid-stride loop。
轉載請註明出處，本文鏈接：https://www.uj5u.com/qiye/517663.html
標籤：图像处理库达字母混合
上一篇：JS執行機制