CUDAでウェーブレットノイズ

  1
  2
  3
  4
  5
  6
  7
  8
  9
 10
 11
 12
 13
 14
 15
 16
 17
 18
 19
 20
 21
 22
 23
 24
 25
 26
 27
 28
 29
 30
 31
 32
 33
 34
 35
 36
 37
 38
 39
 40
 41
 42
 43
 44
 45
 46
 47
 48
 49
 50

float *g_dNoiseTile = 0;
uint g_uNoiseTileSize = 0;
 
void CuSetWaveletTile(float *tile, int n)
{
    if(!g_dNoiseTile && g_uNoiseTileSize != n){
                int size = n*n*sizeof(float);
        cutilSafeCall(cudaMalloc((void**)&g_dNoiseTile, size));
        cutilSafeCall(cudaMemcpy((void*)g_dNoiseTile, (void*)tile, size, cudaMemcpyHostToDevice));
 
        g_uNoiseTileSize = n;
    }
}
 
void CuWNoise2D(float *hW, int nx, int ny, int level)
{
    printf("\n[CuWNoise2D]\n");
 
    if(!g_dNoiseTile) return;
 
        float *dW;
    int size;
    float w = powf(2.0f, (float)level);
 
        size = nx*ny*sizeof(float);
    cutilSafeCall(cudaMalloc((void**)&dW, size));
    cutilSafeCall(cudaMemset((void*)dW, 0, size));
 
    if(!hW) hW = new float[nx*ny];
 
    dim3 block(BLOCK_SIZE, BLOCK_SIZE);
    dim3 grid((nx+block.x-1)/block.x, (ny+block.y-1)/block.y);
 
    wavelet_noise_2d<<< grid, block >>>(g_dNoiseTile, g_uNoiseTileSize, dW, nx, ny, w);
 
        cutilCheckMsg("Kernel execution failed");
    cutilSafeCall(cudaThreadSynchronize());
 
        size = nx*ny*sizeof(float);
    cutilSafeCall(cudaMemcpy(hW, dW, size, cudaMemcpyDeviceToHost));
 
        cutilSafeCall(cudaFree(dW));
}

  1
  2
  3
  4
  5
  6
  7
  8
  9
 10
 11
 12
 13
 14
 15
 16
 17
 18
 19
 20
 21
 22
 23
 24
 25
 26
 27
 28
 29
 30
 31
 32

__device__
float wnoise2d(float *tile, int tile_n, float p[2])
{
    int f[3], c[3];        // filter, noise coef. indices
    int mid[3];
 
    float w[2][3], t, result = 0;
 
        for(int k = 0; k < 2; ++k){
        mid[k] = ceil(p[k]-0.5);
        t = mid[k]-(p[k]-0.5);
 
        w[k][0] = t*t/2; 
        w[k][2] = (1-t)*(1-t)/2; 
        w[k][1] = 1-w[k][0]-w[k][2];
    }
 
        for(f[1] = -1; f[1] <= 1; ++f[1]){
        for(f[0] = -1; f[0] <= 1; ++f[0]){
            float weight = 1;
            for(int k = 0; k < 2; ++k){
                c[k] = fmodf(mid[k]+f[k], tile_n);
                weight *= w[k][f[k]+1];
            }
            result += weight*tile[c[1]*tile_n+c[0]];
        }
    }
 
    return result;
}

  1
  2
  3
  4
  5
  6
  7
  8
  9
 10

        if(k == d){
            w[k][0] = -t; 
            w[k][1] = 2*t-1;
            w[k][2] = 1-t;
        }
        else{
            w[k][0] = t*t/2; 
            w[k][2] = (1-t)*(1-t)/2; 
            w[k][1] = 1-w[k][0]-w[k][2];
        }

  1
  2
  3
  4
  5
  6
  7
  8
  9
 10
 11
 12
 13
 14

__global__
void wavelet_noise_2d(float *tile, int tile_n, float *wt, int nx, int ny, float wpos)
{
    int i = blockIdx.x*blockDim.x+threadIdx.x;
    int j = blockIdx.y*blockDim.y+threadIdx.y;
 
    if(i < nx && j < ny){
        float p[2];
        p[0] = (float)(i+0.5)*wpos/nx;
        p[1] = (float)(j+0.5)*wpos/ny;
 
        wt[i+j*nx] = wnoise2d(tile, tile_n, p);
    }
}

  1
  2
  3
  4
  5
  6
  7
  8
  9
 10
 11
 12
 13
 14
 15
 16
 17
 18
 19
 20
 21
 22
 23
 24
 25
 26
 27
 28
 29
 30
 31
 32
 33
 34
 35

__global__
void wavelet_multiband_noise_2d(float *tile, int tile_n, float *wt, int nx, int ny, int first, int nbands)
{
    int i = blockIdx.x*blockDim.x+threadIdx.x;
    int j = blockIdx.y*blockDim.y+threadIdx.y;
 
    if(i < nx && j < ny){
        float p[2];
        p[0] = (float)(i+0.5)/nx;
        p[1] = (float)(j+0.5)/ny;
 
        float result = 0;
        float w = powf(2.0, (float)first);
        float q[2];
        float sigma_m = 0;
 
        for(int b = 0; b < nbands; ++b){
            q[0] = p[0]*w;
            q[1] = p[1]*w;
            result += wnoise2d(tile, tile_n, q)/w;
 
            sigma_m += 1.0/w*w;
            w *= 2.0;
        }
 
                float sigma_n = 0.265;
        sigma_m = sqrtf(sigma_n*sigma_m);
 
                if(sigma_m) result /= sigma_m;
 
        wt[i+j*nx] = result;
    }
}