Почему это ядро ​​CUDA работает медленно?

Мне нужна помощь, чтобы моя программа cuda работала быстрее. Визуальный профилировщик NVIDIA показывает низкую производительность, говоря «Низкое использование вычислений 1,4%»:

Код ниже. Подготовка первых ядер:

void laskeSyvyydet(int& tiilet0, int& tiilet1, int& tiilet2, int& tiilet3) {

cudaArray *tekstuuriSisaan, *tekstuuriUlos;

//take care of synchronazion
cudaEvent_t cEvent;
cudaEventCreate(&cEvent);

//let's take control of OpenGL textures
cudaGraphicsMapResources(1, &cuda.cMaxSyvyys);
cudaEventRecord(cEvent, 0);
cudaGraphicsMapResources(1, &cuda.cDepthTex);
cudaEventRecord(cEvent, 0);

//need to create CUDA pointers
cudaGraphicsSubResourceGetMappedArray(&tekstuuriSisaan, cuda.cDepthTex, 0, 0);
cudaGraphicsSubResourceGetMappedArray(&tekstuuriUlos, cuda.cMaxSyvyys, 0, 0);

cudaProfilerStart();

//launch kernel
cLaskeSyvyydet(tiilet0, tiilet1, tiilet2, tiilet3, tekstuuriSisaan, tekstuuriUlos);
cudaEventRecord(cEvent, 0);

cudaProfilerStop();

//release textures back to OpenGL
cudaGraphicsUnmapResources(1, &cuda.cMaxSyvyys, 0);
cudaEventRecord(cEvent, 0);
cudaGraphicsUnmapResources(1, &cuda.cDepthTex, 0);
cudaEventRecord(cEvent, 0);

//final synchronazion
cudaEventSynchronize(cEvent);
cudaEventDestroy(cEvent);
}

Запуск ядра:

void cLaskeSyvyydet(int& tiilet0, int& tiilet1, int& tiilet2, int& tiilet3, cudaArray* tekstuuriSisaan, cudaArray* tekstuuriUlos) {

cudaBindTextureToArray(surfRefSisaan, tekstuuriSisaan);
cudaBindSurfaceToArray(surfRefUlos, tekstuuriUlos);

    int blocksW = (int)ceilf( tiilet0 / 32.0f );
    int blocksH = (int)ceilf( tiilet1 / 32.0f );
    dim3 gridDim( blocksW, blocksH, 1 );
    dim3 blockDim(32, 32, 1 );

kLaskeSyvyydet<<<gridDim, blockDim>>>(tiilet0, tiilet1, tiilet2, tiilet3);

}

И ядро:

__global__ void kLaskeSyvyydet(const int tiilet0, const int tiilet1, const int tiilet2, const int tiilet3) {

//first define indexes
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int j = blockIdx.y * blockDim.y + threadIdx.y;
    if (i >= tiilet0 || j >= tiilet1) return;

//if we are inside boundaries, let's find the greatest depth value
    unsigned int takana=0;
    unsigned int ddd;
    uchar4 syvyys;
    uchar4 dd;

//there's possibly four different tile sizes to choose between
    if (j!=tiilet1-1 && i!=tiilet0-1) {

    for (int y=j*BLOCK_SIZE; y<(j+1)*BLOCK_SIZE; y++) {
        for (int x=i*BLOCK_SIZE; x<(i+1)*BLOCK_SIZE; x++) {
            dd=tex2D(surfRefSisaan, x, y);
            ddd=(dd.x << 24) | (dd.y << 16) | (dd.z << 8) | (dd.w);
            takana=max(takana, ddd);
        }
    }

    } else if (j==tiilet1-1 && i!=tiilet0-1) {

    for (int y=j*BLOCK_SIZE; y<j*BLOCK_SIZE+tiilet3; y++) {
        for (int x=i*BLOCK_SIZE; x<(i+1)*BLOCK_SIZE; x++) {
            dd=tex2D(surfRefSisaan, x, y);
            ddd=(dd.x << 24) | (dd.y << 16) | (dd.z << 8) | (dd.w);
            takana=max(takana, ddd);
            }
        }

    } else if (j!=tiilet1-1 && i==tiilet0-1) {

    for (int y=j*BLOCK_SIZE; y<(j+1)*BLOCK_SIZE; y++) {
        for (int x=i*BLOCK_SIZE; x<i*BLOCK_SIZE+tiilet2; x++) {
            dd=tex2D(surfRefSisaan, x, y);
            ddd=(dd.x << 24) | (dd.y << 16) | (dd.z << 8) | (dd.w);
            takana=max(takana, ddd);
        }
    }

    } else if (j==tiilet1-1 && i==tiilet0-1) {

    for (int y=j*BLOCK_SIZE; y<j*BLOCK_SIZE+tiilet3; y++) {
        for (int x=i*BLOCK_SIZE; x<i*BLOCK_SIZE+tiilet2; x++) {
            dd=tex2D(surfRefSisaan, x, y);
            ddd=(dd.x << 24) | (dd.y << 16) | (dd.z << 8) | (dd.w);
            takana=max(takana, ddd);
        }
    }
    }

//if there's empty texture, then we choose the maximum possible value
    if (takana==0) {
    takana=1000000000;
    }

//after slicing the greatest 32bit depth value into four 8bit pieces we write the value into another texture
    syvyys.x=(takana & 0xFF000000) >> 24;
    syvyys.y=(takana & 0x00FF0000) >> 16;
    syvyys.z=(takana & 0x0000FF00) >> 8;
    syvyys.w=(takana & 0x000000FF) >> 0;

    surf2Dwrite(syvyys, surfRefUlos, i*sizeof(syvyys), j, cudaBoundaryModeZero);

}

Пожалуйста, помогите мне заставить это работать быстрее, у меня нет идей...