7,540
社区成员
发帖
与我相关
我的任务
分享
#pragma once
#include <cudatemplates/array.hpp>
class TileGenerator {
public:
typedef Cuda::Array3D<short> ArrayType;
TileGenerator(Cuda::Size<3> tileSize) : _tile(tileSize) {}
virtual ~TileGenerator() {}
ArrayType& getTile()
{
return _tile;
}
virtual float getMaxDerivative() = 0;
virtual float getStandardDeviation() = 0;
virtual float getFrequency() = 0;
protected:
Cuda::Array3D<short> _tile;
};
class NoiseTileGenerator : public TileGenerator {
public:
NoiseTileGenerator(Cuda::Size<3> tileSize, float targetFrequency, int nGaborNoisePulses = 8)
: TileGenerator(tileSize)
, _frequency(targetFrequency)
{
_generateNoiseTile(targetFrequency, nGaborNoisePulses);
}
virtual float getMaxDerivative() { return _maxDerivative; }
virtual float getStandardDeviation() { return _stdDev; }
virtual float getFrequency() { return _frequency; }
private:
void _generateNoiseTile(float targetFrequency, int nGaborNoisePulses);
float _maxDerivative;
float _stdDev;
float _frequency;
};
class SinusoidTileGenerator : public TileGenerator {
public:
SinusoidTileGenerator(Cuda::Size<3> tileSize, float targetFrequency)
: TileGenerator(tileSize)
, _frequency(targetFrequency)
{
_generateSinusoidTile(targetFrequency);
}
virtual float getMaxDerivative() { return _maxDerivative; }
virtual float getStandardDeviation() { return _stdDev; }
virtual float getFrequency() { return _frequency; }
private:
void _generateSinusoidTile(float targetFrequency);
float _maxDerivative;
float _stdDev;
float _frequency;
};
tilegenerator.cu
#include <tileGenerator.cuh>
#include <cudatemplates/devicememorylinear.hpp>
#include <cudatemplates/copy.hpp>
#include <gaborNoise.cuh>
namespace detail {
class TileGenerateFunction {
public:
__device__ virtual short getValue(const float3& pos) = 0;
__device__ void d_generateTile(Cuda::DeviceMemoryLinear3D<short>::KernelData out)
{
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x >= out.size[0] || y >= out.size[1]) {
return;
}
for (int z = 0; z < out.size[2]; ++z) {
float3 pos = make_float3(x / (float)out.size[0], y / (float)out.size[1], z / (float)out.size[2]);
out.data[z * out.stride[1] + y * out.stride[0] + x] = getValue(pos);
}
}
};
class NoiseTileGenerateFunction : public TileGenerateFunction {
public:
__device__ NoiseTileGenerateFunction(float r, float omegaLen, float a, int nPulses)
: _noise(r, omegaLen, a, nPulses) {}
__device__ virtual short getValue(const float3& pos)
{
return (short)((_noise.noise_evaluate(pos) / 20.f) * 32767);
}
private:
GaborNoise::KernelData _noise;
};
class SinusoidTileGenerateFunction : public TileGenerateFunction {
public:
__device__ SinusoidTileGenerateFunction(float a)
: _a(a) {}
__device__ virtual short getValue(const float3& pos)
{
float x = sinf(_a * pos.x) * sinf(_a* pos.y) * sinf(_a* pos.z);
x = x < 0.f ? -pow(fabsf(x), 0.7f) : pow(fabsf(x), 0.7f);
return (short)((x / 20.f) * 32767);
}
private:
float _a;
};
__global__ void d_generateNoiseTile(Cuda::DeviceMemoryLinear3D<short>::KernelData out, float r, float omegaLen, float a, int nPulses)
{
NoiseTileGenerateFunction(r, omegaLen, a, nPulses).d_generateTile(out);
}
__global__ void d_generateSinusoidTile(Cuda::DeviceMemoryLinear3D<short>::KernelData out, float a)
{
SinusoidTileGenerateFunction(a).d_generateTile(out);
}
}
void NoiseTileGenerator::_generateNoiseTile(float targetFrequency, int nGaborNoisePulses)
{
float h_gridSize = 1.f / targetFrequency;
float h_omegaLen = targetFrequency;
// sqrt(-logf(0.05f) / pi) = 0.97650970247
float h_a = 0.97650970247f / h_gridSize;
const float derivFactor = nGaborNoisePulses; // Experimental
float x = -h_omegaLen/(2*h_a) + sqrtf(h_omegaLen*h_omegaLen/(4*h_a*h_a) + 1/(2*pi*h_a));
_maxDerivative = nGaborNoisePulses * 2 * pi * exp(-pi*h_a*x*x)*(h_omegaLen + h_a*x) / derivFactor;
_stdDev = nGaborNoisePulses / (4 * sqrtf(2.f) * powf(sqrtf(-logf(0.05f) / pi) , 3));
dim3 blockSize(16, 16);
dim3 gridSize((int)ceilf((float)_tile.size[0] / blockSize.x), (int)ceilf((float)_tile.size[1] / blockSize.y));
Cuda::DeviceMemoryLinear3D<short> tileMemory(_tile.size);
detail::d_generateNoiseTile<<<gridSize, blockSize>>>(tileMemory, h_gridSize, h_omegaLen, h_a, nGaborNoisePulses);
//d_generateNoiseTile<<<gridSize, blockSize>>>(tileMemory, h_gridSize, h_omegaLen, h_a, nGaborNoisePulses);
Cuda::copy(_tile, tileMemory);
}
void SinusoidTileGenerator::_generateSinusoidTile(float targetFrequency)
{
// TODO
float a = targetFrequency * pi;
_maxDerivative = a;
_stdDev = 0.2;
// ENDTODO
dim3 blockSize(16, 16);
dim3 gridSize((int)ceilf((float)_tile.size[0] / blockSize.x), (int)ceilf((float)_tile.size[1] / blockSize.y));
Cuda::DeviceMemoryLinear3D<short> tileMemory(_tile.size);
detail::d_generateSinusoidTile<<<gridSize, blockSize>>>(tileMemory, a);
//d_generateSinusoidTile<<<gridSize, blockSize>>>(tileMemory, a);
Cuda::copy(_tile, tileMemory);
}