AI/rocm-systems
2
0
Garpu 0
Kode Masalah Tarik Permintaan Actions Packages Projects Rilis Wiki Kegiatan

SWDEV-331248 - Add more image tests in sample (#2709)

Telusuri Sumber 
In samples/2_Cookbook/11_texture_driver, add
Vector data types(char4, short4, int4, float4);
More arithmetic data types(char, short, int);

Change-Id: I54aa482213d340d32cf912601adead0812c2323a

[ROCm/hip-tests commit: 1a422d9798]
This commit is contained in:
ROCm CI Service Account
2022-06-02 17:01:14 +05:30
committed by GitHub
orang tua c91bfdb8c6
melakukan b48478288b
2 mengubah file dengan 233 tambahan dan 87 penghapusan
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip-tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp
+53 -6
Melihat File
@@ -19,15 +19,62 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include "hip/hip_runtime.h"
texture<float, 2, hipReadModeElementType> tex;
texture<char, hipTextureType2D, hipReadModeElementType> texChar;
texture<short, hipTextureType2D, hipReadModeElementType> texShort;
texture<int, hipTextureType2D, hipReadModeElementType> texInt;
texture<float, hipTextureType2D, hipReadModeElementType> texFloat;
extern "C" __global__ void tex2dKernel(float* outputData, int width, int height) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
texture<char4, hipTextureType2D, hipReadModeElementType> texChar4;
texture<short4, hipTextureType2D, hipReadModeElementType> texShort4;
texture<int4, hipTextureType2D, hipReadModeElementType> texInt4;
texture<float4, hipTextureType2D, hipReadModeElementType> texFloat4;
extern "C" __global__ void tex2dKernelChar(char* outputData, int width, int height) {
int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
outputData[y * width + x] = tex2D(tex, x, y);
#endif
outputData[y * width + x] = tex2D(texChar, x, y);
}
extern "C" __global__ void tex2dKernelShort(short* outputData, int width, int height) {
int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
outputData[y * width + x] = tex2D(texShort, x, y);
}
extern "C" __global__ void tex2dKernelInt(int* outputData, int width, int height) {
int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
outputData[y * width + x] = tex2D(texInt, x, y);
}
extern "C" __global__ void tex2dKernelFloat(float* outputData, int width, int height) {
int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
outputData[y * width + x] = tex2D(texFloat, x, y);
}
extern "C" __global__ void tex2dKernelChar4(char4* outputData, int width, int height) {
int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
outputData[y * width + x] = tex2D(texChar4, x, y);
}
extern "C" __global__ void tex2dKernelShort4(short4* outputData, int width, int height) {
int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
outputData[y * width + x] = tex2D(texShort4, x, y);
}
extern "C" __global__ void tex2dKernelInt4(int4* outputData, int width, int height) {
int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
outputData[y * width + x] = tex2D(texInt4, x, y);
}
extern "C" __global__ void tex2dKernelFloat4(float4* outputData, int width, int height) {
int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
outputData[y * width + x] = tex2D(texFloat4, x, y);
}